fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-02-24 03:00:30 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

zink: implement compiler handling for subgroup ballot builtins/intrinsics

Browse source 
these are all lowered and unremarkable

Reviewed-by: Dave Airlie <airlied@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11990>
This commit is contained in:
Mike Blumenkrantz 2021-04-05 12:02:08 -04:00 • committed by Marge Bot
parent 252165d925
commit a05693d332
1 changed files with 55 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

55
src/gallium/drivers/zink/nir_to_spirv/nir_to_spirv.c
View file

@ -91,6 +91,15 @@ struct ntv_context {
local_group_size_var,
shared_block_var,
base_vertex_var, base_instance_var, draw_id_var;
SpvId subgroup_eq_mask_var,
subgroup_ge_mask_var,
subgroup_gt_mask_var,
subgroup_id_var,
subgroup_invocation_var,
subgroup_le_mask_var,
subgroup_lt_mask_var,
subgroup_size_var;
};
static SpvId
@ -2796,6 +2805,48 @@ emit_intrinsic(struct ntv_context *ctx, nir_intrinsic_instr *intr)
emit_load_uint_input(ctx, intr, &ctx->local_invocation_index_var, "gl_LocalInvocationIndex", SpvBuiltInLocalInvocationIndex);
break;
#define LOAD_SHADER_BALLOT(lowercase, camelcase) \
case nir_intrinsic_load_##lowercase: \
emit_load_uint_input(ctx, intr, &ctx->lowercase##_var, "gl_"#camelcase, SpvBuiltIn##camelcase); \
break
LOAD_SHADER_BALLOT(subgroup_id, SubgroupId);
LOAD_SHADER_BALLOT(subgroup_eq_mask, SubgroupEqMask);
LOAD_SHADER_BALLOT(subgroup_ge_mask, SubgroupGeMask);
LOAD_SHADER_BALLOT(subgroup_invocation, SubgroupLocalInvocationId);
LOAD_SHADER_BALLOT(subgroup_le_mask, SubgroupLeMask);
LOAD_SHADER_BALLOT(subgroup_lt_mask, SubgroupLtMask);
LOAD_SHADER_BALLOT(subgroup_size, SubgroupSize);
case nir_intrinsic_ballot: {
spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
SpvId type = get_dest_uvec_type(ctx, &intr->dest);
SpvId result = emit_unop(ctx, SpvOpSubgroupBallotKHR, type, get_src(ctx, &intr->src[0]));
store_dest(ctx, &intr->dest, result, nir_type_uint);
break;
}
case nir_intrinsic_read_first_invocation: {
spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
SpvId type = get_dest_type(ctx, &intr->dest, nir_type_uint);
SpvId result = emit_unop(ctx, SpvOpSubgroupFirstInvocationKHR, type, get_src(ctx, &intr->src[0]));
store_dest(ctx, &intr->dest, result, nir_type_uint);
break;
}
case nir_intrinsic_read_invocation: {
spirv_builder_emit_cap(&ctx->builder, SpvCapabilitySubgroupBallotKHR);
spirv_builder_emit_extension(&ctx->builder, "SPV_KHR_shader_ballot");
SpvId type = get_dest_type(ctx, &intr->dest, nir_type_uint);
SpvId result = emit_binop(ctx, SpvOpSubgroupReadInvocationKHR, type,
get_src(ctx, &intr->src[0]),
get_src(ctx, &intr->src[1]));
store_dest(ctx, &intr->dest, result, nir_type_uint);
break;
}
case nir_intrinsic_load_workgroup_size: {
assert(ctx->local_group_size_var);
store_dest(ctx, &intr->dest, ctx->local_group_size_var, nir_type_uint);
@ -3842,6 +3893,10 @@ nir_to_spirv(struct nir_shader *s, const struct zink_so_info *so_info, uint32_t
default:
break;
}
if (BITSET_TEST_RANGE(s->info.system_values_read, SYSTEM_VALUE_SUBGROUP_SIZE, SYSTEM_VALUE_SUBGROUP_LT_MASK)) {
spirv_builder_emit_cap(&ctx.builder, SpvCapabilitySubgroupBallotKHR);
spirv_builder_emit_extension(&ctx.builder, "SPV_KHR_shader_ballot");
}
if (s->info.has_transform_feedback_varyings) {
spirv_builder_emit_cap(&ctx.builder, SpvCapabilityTransformFeedback);
spirv_builder_emit_exec_mode(&ctx.builder, entry_point,