fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-05 16:08:04 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

vtn: more CL subgroups

Browse source 
v2: handle ExecutionModes

Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Nora Allen <blackcatgames@protonmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/22893>
This commit is contained in:
Karol Herbst 2023-05-06 20:58:22 +02:00 • committed by Marge Bot
parent 11cb1a7bd7
commit 17e749dc00
3 changed files with 16 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/compiler/shader_info.h
View file

@ -247,6 +247,7 @@ typedef struct shader_info {
uint16_t workgroup_size[3];
enum gl_subgroup_size subgroup_size;
uint8_t num_subgroups;
/**
* Uses subgroup intrinsics which can communicate across a quad.

13
src/compiler/spirv/spirv_to_nir.c
View file

@ -4762,8 +4762,6 @@ vtn_handle_preamble_instruction(struct vtn_builder *b, SpvOp opcode,
/* Missing :
* - SpvOpGetKernelLocalSizeForSubgroupCount
* - SpvOpGetKernelMaxNumSubgroups
* - SpvExecutionModeSubgroupsPerWorkgroup
* - SpvExecutionModeSubgroupsPerWorkgroupId
*/
vtn_warn("Not fully supported capability: %s",
spirv_capability_to_string(cap));
@ -5442,6 +5440,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
case SpvExecutionModeLocalSizeId:
case SpvExecutionModeLocalSizeHintId:
case SpvExecutionModeSubgroupsPerWorkgroupId:
/* Handled later by vtn_handle_execution_mode_id(). */
break;
@ -5451,6 +5450,11 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
b->shader->info.subgroup_size = mode->operands[0];
break;
case SpvExecutionModeSubgroupsPerWorkgroup:
vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
b->shader->info.num_subgroups = mode->operands[0];
break;
case SpvExecutionModeSubgroupUniformControlFlowKHR:
/* There's no corresponding SPIR-V capability, so check here. */
vtn_fail_if(!b->options->caps.subgroup_uniform_control_flow,
@ -5525,6 +5529,11 @@ vtn_handle_execution_mode_id(struct vtn_builder *b, struct vtn_value *entry_poin
b->shader->info.cs.workgroup_size_hint[2] = vtn_constant_uint(b, mode->operands[2]);
break;
case SpvExecutionModeSubgroupsPerWorkgroupId:
vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
b->shader->info.num_subgroups = vtn_constant_uint(b, mode->operands[0]);
break;
default:
/* Nothing to do. Literal execution modes already handled by
* vtn_handle_execution_mode(). */

4
src/compiler/spirv/vtn_variables.c
View file

@ -982,6 +982,8 @@ vtn_get_builtin_location(struct vtn_builder *b,
set_mode_system_value(b, mode);
break;
case SpvBuiltInSubgroupSize:
/* TODO once we support non uniform work groups we have to fix this */
case SpvBuiltInSubgroupMaxSize:
*location = SYSTEM_VALUE_SUBGROUP_SIZE;
set_mode_system_value(b, mode);
break;
@ -994,6 +996,8 @@ vtn_get_builtin_location(struct vtn_builder *b,
set_mode_system_value(b, mode);
break;
case SpvBuiltInNumSubgroups:
/* TODO once we support non uniform work groups we have to fix this */
case SpvBuiltInNumEnqueuedSubgroups:
*location = SYSTEM_VALUE_NUM_SUBGROUPS;
set_mode_system_value(b, mode);
break;