<div dir="ltr"><div class="gmail_extra"><div class="gmail_quote">On Fri, Mar 23, 2018 at 12:33 PM, Karol Herbst <<a href="mailto:kherbst@redhat.com" target="_blank">kherbst@redhat.com</a>> wrote: <blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">With OpenCL the size of some system value depends on the Physical model choosen, so we need a way to load any system value as 32 or 64 bit. Signed-off-by: Karol Herbst <<a href="mailto:kherbst@redhat.com">kherbst@redhat.com</a>> --- src/compiler/nir/nir_builder.h | 10 +++++--- src/compiler/nir/nir_lower_alpha_test.c | 2 +- src/compiler/nir/nir_lower_clip.c | 3 ++- src/compiler/nir/nir_lower_subgroups.c | 8 +++--- src/compiler/nir/nir_lower_system_values.c | 31 ++++++++++++------------ src/compiler/nir/nir_lower_two_sided_color.c | 2 +- src/compiler/nir/nir_lower_wpos_center.c | 2 +- src/compiler/spirv/vtn_subgroup.c | 2 +- src/gallium/auxiliary/nir/tgsi_to_nir.c | 3 ++- src/intel/blorp/blorp_blit.c | 2 +- src/intel/blorp/blorp_clear.c | 2 +- src/intel/compiler/brw_nir_lower_cs_intrinsics.c | 6 ++--- src/mesa/drivers/dri/i965/brw_tcs.c | 2 +- 13 files changed, 40 insertions(+), 35 deletions(-) diff --git a/src/compiler/nir/nir_builder.h b/src/compiler/nir/nir_builder.h index 36e0ae3ac63..4e93cd08169 100644 --- a/src/compiler/nir/nir_builder.h +++ b/src/compiler/nir/nir_builder.h @@ -612,13 +612,14 @@ nir_copy_var(nir_builder *build, nir_variable *dest, nir_variable *src) /* Generic builder for system values. */ static inline nir_ssa_def * -nir_load_system_value(nir_builder *build, nir_intrinsic_op op, int index) +nir_load_system_value(nir_builder *build, nir_intrinsic_op op, int index, + unsigned bit_size) { nir_intrinsic_instr *load = nir_intrinsic_instr_create(build->shader, op); load->num_components = nir_intrinsic_infos[op].dest_components; load->const_index[0] = index; nir_ssa_dest_init(&load->instr, &load->dest, - nir_intrinsic_infos[op].dest_components, 32, NULL); + nir_intrinsic_infos[op].dest_components, bit_size, NULL); nir_builder_instr_insert(build, &load->instr); return &load->dest.ssa; } @@ -630,9 +631,10 @@ nir_load_system_value(nir_builder *build, nir_intrinsic_op op, int index) #define DEFINE_SYSTEM_VALUE(name) \ static inline nir_ssa_def * \ - nir_load_##name(nir_builder *build) \ + nir_load_##name(nir_builder *build, unsigned bit_size) \ </blockquote><div> </div><div>I was really hoping that this change wouldn't touch every single intrinsic helper. Maybe with Rob's python-based intrinsics table we can do something better. </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> { \ - return nir_load_system_value(build, nir_intrinsic_load_##name, 0); \ + return nir_load_system_value(build, nir_intrinsic_load_##name, 0, \ + bit_size); \ } #include "nir_intrinsics.h" diff --git a/src/compiler/nir/nir_lower_alpha_test.c b/src/compiler/nir/nir_lower_alpha_test.c index 6bf9ff142df..29f91ab9428 100644 --- a/src/compiler/nir/nir_lower_alpha_test.c +++ b/src/compiler/nir/nir_lower_alpha_test.c @@ -92,7 +92,7 @@ nir_lower_alpha_test(nir_shader *shader, enum compare_func func, nir_ssa_def *condition = nir_compare_func(&b, func, - alpha, nir_load_alpha_ref_float(&b)); + alpha, nir_load_alpha_ref_float(&b, 32)); nir_intrinsic_instr *discard = nir_intrinsic_instr_create(b.shader, diff --git a/src/compiler/nir/nir_lower_clip.c b/src/compiler/nir/nir_lower_clip.c index ea12f51a7bb..b9a91f7d40b 100644 --- a/src/compiler/nir/nir_lower_clip.c +++ b/src/compiler/nir/nir_lower_clip.c @@ -174,7 +174,8 @@ lower_clip_vs(nir_function_impl *impl, unsigned ucp_enables, for (int plane = 0; plane < MAX_CLIP_PLANES; plane++) { if (ucp_enables & (1 << plane)) { nir_ssa_def *ucp = - nir_load_system_value(&b, nir_intrinsic_load_user_clip_plane, plane); + nir_load_system_value(&b, nir_intrinsic_load_user_clip_plane, + plane, 32); /* calculate clipdist[plane] - dot(ucp, cv): */ clipdist[plane] = nir_fdot4(&b, ucp, cv); diff --git a/src/compiler/nir/nir_lower_subgroups.c b/src/compiler/nir/nir_lower_subgroups.c index 0d3c83b7951..7e910c013a9 100644 --- a/src/compiler/nir/nir_lower_subgroups.c +++ b/src/compiler/nir/nir_lower_subgroups.c @@ -190,7 +190,7 @@ static nir_ssa_def * lower_shuffle(nir_builder *b, nir_intrinsic_instr *intrin, bool lower_to_scalar) { - nir_ssa_def *index = nir_load_subgroup_invocation(b); + nir_ssa_def *index = nir_load_subgroup_invocation(b, 32); switch (intrin->intrinsic) { case nir_intrinsic_shuffle_xor: assert(intrin->src[1].is_ssa); @@ -300,7 +300,7 @@ lower_subgroups_intrin(nir_builder *b, nir_intrinsic_instr *intrin, assert(options->subgroup_size <= 64); uint64_t group_mask = ~0ull >> (64 - options->subgroup_size); - nir_ssa_def *count = nir_load_subgroup_invocation(b); + nir_ssa_def *count = nir_load_subgroup_invocation(b, 32); nir_ssa_def *val; switch (intrin->intrinsic) { case nir_intrinsic_load_subgroup_eq_mask: @@ -373,7 +373,7 @@ lower_subgroups_intrin(nir_builder *b, nir_intrinsic_instr *intrin, case nir_intrinsic_ballot_bit_count_exclusive: case nir_intrinsic_ballot_bit_count_inclusive: { - nir_ssa_def *count = nir_load_subgroup_invocation(b); + nir_ssa_def *count = nir_load_subgroup_invocation(b, 32); nir_ssa_def *mask = nir_imm_intN_t(b, ~0ull, options->ballot_bit_size); if (intrin->intrinsic == nir_intrinsic_ballot_bit_count_inclusive) { const unsigned bits = options->ballot_bit_size; @@ -396,7 +396,7 @@ lower_subgroups_intrin(nir_builder *b, nir_intrinsic_instr *intrin, nir_ssa_dest_init(&first->instr, &first->dest, 1, 32, NULL); nir_builder_instr_insert(b, &first->instr); - return nir_ieq(b, nir_load_subgroup_invocation(b), &first->dest.ssa); + return nir_ieq(b, nir_load_subgroup_invocation(b, 32), &first->dest.ssa); } case nir_intrinsic_shuffle: diff --git a/src/compiler/nir/nir_lower_system_values.c b/src/compiler/nir/nir_lower_system_values.c index fb560ee21bb..d507c28f421 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -46,6 +46,7 @@ convert_block(nir_block *block, nir_builder *b) if (var->data.mode != nir_var_system_value) continue; + unsigned bit_size = load_var->dest.ssa.bit_size; b->cursor = nir_after_instr(&load_var->instr); nir_ssa_def *sysval = NULL; @@ -59,15 +60,15 @@ convert_block(nir_block *block, nir_builder *b) nir_const_value local_size; memset(&local_size, 0, sizeof(local_size)); - local_size.u32[0] = b->shader->info.cs.local_size[0]; - local_size.u32[1] = b->shader->info.cs.local_size[1]; - local_size.u32[2] = b->shader->info.cs.local_size[2]; + local_size.u64[0] = b->shader->info.cs.local_size[0]; + local_size.u64[1] = b->shader->info.cs.local_size[1]; + local_size.u64[2] = b->shader->info.cs.local_size[2]; - nir_ssa_def *group_id = nir_load_work_group_id(b); - nir_ssa_def *local_id = nir_load_local_invocation_id(b); + nir_ssa_def *group_id = nir_load_work_group_id(b, bit_size); + nir_ssa_def *local_id = nir_load_local_invocation_id(b, bit_size); sysval = nir_iadd(b, nir_imul(b, group_id, - nir_build_imm(b, 3, 32, local_size)), + nir_build_imm(b, 3, bit_size, local_size)), </blockquote><div> </div><div>This doesn't do what you think it does. Due to the way that the different arrays in nir_const_value alias, you can't put 64-bit values in the nir_const_value and then use 32 for nir_build_imm and expect it to work. We can either make a smarter immediate builder or just insert a u2u64 instruction which will get properly constant folded. </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> local_id); break; } @@ -86,12 +87,12 @@ convert_block(nir_block *block, nir_builder *b) * gl_WorkGroupSize.y + gl_LocalInvocationID.y * * gl_WorkGroupSize.x + gl_LocalInvocationID.x" */ - nir_ssa_def *local_id = nir_load_local_invocation_id(b); + nir_ssa_def *local_id = nir_load_local_invocation_id(b, bit_size); nir_ssa_def *size_x = - nir_imm_int(b, b->shader->info.cs.local_size[0]); + nir_imm_intN_t(b, b->shader->info.cs.local_size[0], bit_size); nir_ssa_def *size_y = - nir_imm_int(b, b->shader->info.cs.local_size[1]); + nir_imm_intN_t(b, b->shader->info.cs.local_size[1], bit_size); sysval = nir_imul(b, nir_channel(b, local_id, 2), nir_imul(b, size_x, size_y)); @@ -104,17 +105,17 @@ convert_block(nir_block *block, nir_builder *b) case SYSTEM_VALUE_VERTEX_ID: if (b->shader->options->vertex_id_zero_based) { sysval = nir_iadd(b, - nir_load_vertex_id_zero_base(b), - nir_load_base_vertex(b)); + nir_load_vertex_id_zero_base(b, bit_size), + nir_load_base_vertex(b, bit_size)); } else { - sysval = nir_load_vertex_id(b); + sysval = nir_load_vertex_id(b, bit_size); } break; case SYSTEM_VALUE_INSTANCE_INDEX: sysval = nir_iadd(b, - nir_load_instance_id(b), - nir_load_base_instance(b)); + nir_load_instance_id(b, bit_size), + nir_load_base_instance(b, bit_size)); break; case SYSTEM_VALUE_SUBGROUP_EQ_MASK: @@ -145,7 +146,7 @@ convert_block(nir_block *block, nir_builder *b) if (sysval == NULL) { nir_intrinsic_op sysval_op = nir_intrinsic_from_system_value(var->data.location); - sysval = nir_load_system_value(b, sysval_op, 0); + sysval = nir_load_system_value(b, sysval_op, 0, bit_size); } nir_ssa_def_rewrite_uses(&load_var->dest.ssa, nir_src_for_ssa(sysval)); diff --git a/src/compiler/nir/nir_lower_two_sided_color.c b/src/compiler/nir/nir_lower_two_sided_color.c index b6742ab2462..20af88b6aec 100644 --- a/src/compiler/nir/nir_lower_two_sided_color.c +++ b/src/compiler/nir/nir_lower_two_sided_color.c @@ -158,7 +158,7 @@ nir_lower_two_sided_color_block(nir_block *block, * bcsel(load_system_value(FACE), load_input(COLn), load_input(BFCn)) */ b->cursor = nir_before_instr(&intr->instr); - nir_ssa_def *face = nir_load_front_face(b); + nir_ssa_def *face = nir_load_front_face(b, 32); nir_ssa_def *front = load_input(b, state->colors[idx].front); nir_ssa_def *back = load_input(b, state->colors[idx].back); nir_ssa_def *color = nir_bcsel(b, face, front, back); diff --git a/src/compiler/nir/nir_lower_wpos_center.c b/src/compiler/nir/nir_lower_wpos_center.c index dca810d735e..a0d9719e270 100644 --- a/src/compiler/nir/nir_lower_wpos_center.c +++ b/src/compiler/nir/nir_lower_wpos_center.c @@ -58,7 +58,7 @@ update_fragcoord(nir_builder *b, nir_intrinsic_instr *intr, wpos = nir_fadd(b, wpos, nir_imm_vec4(b, 0.5f, 0.5f, 0.0f, 0.0f)); } else { nir_ssa_def *spos = - nir_load_system_value(b, nir_intrinsic_load_sample_pos, 0); + nir_load_system_value(b, nir_intrinsic_load_sample_pos, 0, 32); wpos = nir_fadd(b, wpos, nir_vec4(b, diff --git a/src/compiler/spirv/vtn_subgroup.c b/src/compiler/spirv/vtn_subgroup.c index bd3143962be..50a4ecc2dcc 100644 --- a/src/compiler/spirv/vtn_subgroup.c +++ b/src/compiler/spirv/vtn_subgroup.c @@ -110,7 +110,7 @@ vtn_handle_subgroup(struct vtn_builder *b, SpvOp opcode, nir_intrinsic_ballot_bitfield_extract); intrin->src[0] = nir_src_for_ssa(vtn_ssa_value(b, w[4])->def); - intrin->src[1] = nir_src_for_ssa(nir_load_subgroup_invocation(&b->nb)); + intrin->src[1] = nir_src_for_ssa(nir_load_subgroup_invocation(&b->nb, 32)); nir_ssa_dest_init(&intrin->instr, &intrin->dest, 1, 32, NULL); nir_builder_instr_insert(&b->nb, &intrin->instr); diff --git a/src/gallium/auxiliary/nir/tgsi_to_nir.c b/src/gallium/auxiliary/nir/tgsi_to_nir.c index f8df4c10137..852b24eaaf1 100644 --- a/src/gallium/auxiliary/nir/tgsi_to_nir.c +++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c @@ -610,7 +610,8 @@ ttn_src_for_file_and_index(struct ttn_compile *c, unsigned file, unsigned index, nir_ssa_def *tgsi_frontface[4] = { nir_bcsel(&c->build, nir_load_system_value(&c->build, - nir_intrinsic_load_front_face, 0), + nir_intrinsic_load_front_face, + 0, 32), nir_imm_float(&c->build, 1.0), nir_imm_float(&c->build, -1.0)), nir_imm_float(&c->build, 0.0), diff --git a/src/intel/blorp/blorp_blit.c b/src/intel/blorp/blorp_blit.c index 0757db0d04b..ca70734981a 100644 --- a/src/intel/blorp/blorp_blit.c +++ b/src/intel/blorp/blorp_blit.c @@ -114,7 +114,7 @@ blorp_blit_get_frag_coords(nir_builder *b, if (key->persample_msaa_dispatch) { return nir_vec3(b, nir_channel(b, coord, 0), nir_channel(b, coord, 1), - nir_load_sample_id(b)); + nir_load_sample_id(b, 32)); } else { return nir_vec2(b, nir_channel(b, coord, 0), nir_channel(b, coord, 1)); } diff --git a/src/intel/blorp/blorp_clear.c b/src/intel/blorp/blorp_clear.c index 832e8ee26f9..c0207d8fa0c 100644 --- a/src/intel/blorp/blorp_clear.c +++ b/src/intel/blorp/blorp_clear.c @@ -880,7 +880,7 @@ blorp_params_get_mcs_partial_resolve_kernel(struct blorp_context *blorp, /* Do an MCS fetch and check if it is equal to the magic clear value */ nir_ssa_def *mcs = blorp_nir_txf_ms_mcs(&b, nir_f2i32(&b, blorp_nir_frag_coord(&b)), - nir_load_layer_id(&b)); + nir_load_layer_id(&b, 32)); nir_ssa_def *is_clear = blorp_nir_mcs_is_clear_color(&b, mcs, blorp_key.num_samples); diff --git a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c index bfbdea0e8fa..846e82ffdf9 100644 --- a/src/intel/compiler/brw_nir_lower_cs_intrinsics.c +++ b/src/intel/compiler/brw_nir_lower_cs_intrinsics.c @@ -61,11 +61,11 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, if (state->local_workgroup_size <= state->dispatch_width) subgroup_id = nir_imm_int(b, 0); else - subgroup_id = nir_load_subgroup_id(b); + subgroup_id = nir_load_subgroup_id(b, 32); nir_ssa_def *thread_local_id = nir_imul(b, subgroup_id, nir_imm_int(b, state->dispatch_width)); - nir_ssa_def *channel = nir_load_subgroup_invocation(b); + nir_ssa_def *channel = nir_load_subgroup_invocation(b, 32); sysval = nir_iadd(b, channel, thread_local_id); break; } @@ -86,7 +86,7 @@ lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, */ unsigned *size = nir->info.cs.local_size; - nir_ssa_def *local_index = nir_load_local_invocation_index(b); + nir_ssa_def *local_index = nir_load_local_invocation_index(b, 32); nir_const_value uvec3; memset(&uvec3, 0, sizeof(uvec3)); diff --git a/src/mesa/drivers/dri/i965/brw_tcs.c b/src/mesa/drivers/dri/i965/brw_tcs.c index 931ef64166c..dda6431108d 100644 --- a/src/mesa/drivers/dri/i965/brw_tcs.c +++ b/src/mesa/drivers/dri/i965/brw_tcs.c @@ -48,7 +48,7 @@ create_passthrough_tcs(void *mem_ctx, const struct brw_compiler *compiler, nir_intrinsic_instr *store; nir_ssa_def *zero = nir_imm_int(&b, 0); nir_ssa_def *invoc_id = - nir_load_system_value(&b, nir_intrinsic_load_invocation_id, 0); + nir_load_system_value(&b, nir_intrinsic_load_invocation_id, 0, 32); nir->info.inputs_read = key->outputs_written & ~(VARYING_BIT_TESS_LEVEL_INNER | VARYING_BIT_TESS_LEVEL_OUTER); -- 2.14.3 _______________________________________________ mesa-dev mailing list <a href="mailto:mesa-dev@lists.freedesktop.org">mesa-dev@lists.freedesktop.org</a> <a href="https://lists.freedesktop.org/mailman/listinfo/mesa-dev" rel="noreferrer" target="_blank">https://lists.freedesktop.org/mailman/listinfo/mesa-dev</a> </blockquote></div> </div></div>