<div dir="ltr"><div class="gmail_quote"><div dir="ltr">On Mon, Jul 16, 2018 at 7:29 AM Karol Herbst <<a href="mailto:kherbst@redhat.com">kherbst@redhat.com</a>> wrote: </div><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. We could probably be a lot smarter and specify which system values might be valid as 32 and/or 64 bit, but I get the feeling it isn't really worth the effort and we can simply depend on the dest type choosen by the API. </blockquote><div> </div><div>I'm not sure what I think about this. Most system values (other than the few used by OpenCL) are always 32 bits all the time and back-ends are likely to depend on this. At the very least, it might be nice to have some sort of validation that the bit sizes are correct before we get there.</div><div> </div><div>One solution to this would be to add a dest_bit_size field to nir_intrinsic_info and use the convention of dest_bit_size == 0 means it can be anything. Then nir_builder_opcodes.py can use that to create functions which either require the bit size or don't. Also, we could have the validator properly validate intrinsic destination bit sizes. It's also something that we could do somewhat incramentally because defaulting everything to dest_bit_size = 0 gives the current behavior.</div><div> </div><div>--Jason </div><div> </div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex"> Signed-off-by: Karol Herbst <<a href="mailto:kherbst@redhat.com" target="_blank">kherbst@redhat.com</a>> --- src/compiler/nir/nir_builder_opcodes_h.py | 9 ++-- 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 | 49 +++++++++++-------- 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 +- .../compiler/brw_nir_lower_cs_intrinsics.c | 6 +-- src/mesa/drivers/dri/i965/brw_tcs.c | 2 +- 13 files changed, 52 insertions(+), 40 deletions(-) diff --git a/src/compiler/nir/nir_builder_opcodes_h.py b/src/compiler/nir/nir_builder_opcodes_h.py index 72cf5b4549d..d16dac6b16e 100644 --- a/src/compiler/nir/nir_builder_opcodes_h.py +++ b/src/compiler/nir/nir_builder_opcodes_h.py @@ -44,22 +44,23 @@ nir_${name}(nir_builder *build, ${src_decl_list(opcode.num_inputs)}) /* 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; } % for name, opcode in filter(lambda v: v[1].sysval, sorted(INTR_OPCODES.iteritems())): static inline nir_ssa_def * -nir_${name}(nir_builder *build) +nir_${name}(nir_builder *build, unsigned bit_size) { - return nir_load_system_value(build, nir_intrinsic_${name}, 0); + return nir_load_system_value(build, nir_intrinsic_${name}, 0, bit_size); } % endfor diff --git a/src/compiler/nir/nir_lower_alpha_test.c b/src/compiler/nir/nir_lower_alpha_test.c index ddd815765bd..8341a0246d2 100644 --- a/src/compiler/nir/nir_lower_alpha_test.c +++ b/src/compiler/nir/nir_lower_alpha_test.c @@ -95,7 +95,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 ee5e8bd644b..c474b9fd27a 100644 --- a/src/compiler/nir/nir_lower_subgroups.c +++ b/src/compiler/nir/nir_lower_subgroups.c @@ -226,7 +226,7 @@ static nir_ssa_def * lower_shuffle(nir_builder *b, nir_intrinsic_instr *intrin, bool lower_to_scalar, bool lower_to_32bit) { - 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); @@ -338,7 +338,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: @@ -411,7 +411,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; @@ -434,7 +434,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 da04895d66c..41f939dd935 100644 --- a/src/compiler/nir/nir_lower_system_values.c +++ b/src/compiler/nir/nir_lower_system_values.c @@ -29,14 +29,22 @@ #include "nir_builder.h" static nir_ssa_def* -build_local_group_size(nir_builder *b) +build_local_group_size(nir_builder *b, unsigned bit_size) { 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]; - return nir_build_imm(b, 3, 32, local_size); + if (bit_size == 64) { + 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]; + } else if (bit_size == 32) { + 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]; + } else { + assert(!"local group size can't be smaller than 32 bits"); + } + return nir_build_imm(b, 3, bit_size, local_size); } static bool @@ -67,6 +75,7 @@ convert_block(nir_block *block, nir_builder *b) } nir_variable *var = deref->var; + unsigned bit_size = load_deref->dest.ssa.bit_size; b->cursor = nir_after_instr(&load_deref->instr); nir_ssa_def *sysval = NULL; @@ -77,9 +86,9 @@ convert_block(nir_block *block, nir_builder *b) * "The value of gl_GlobalInvocationID is equal to * gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInvocationID" */ - nir_ssa_def *group_size = build_local_group_size(b); - 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_size = build_local_group_size(b, bit_size); + 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, group_size), local_id); break; @@ -99,7 +108,7 @@ 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]); @@ -115,17 +124,17 @@ convert_block(nir_block *block, nir_builder *b) } case SYSTEM_VALUE_LOCAL_GROUP_SIZE: { - sysval = build_local_group_size(b); + sysval = build_local_group_size(b, bit_size); break; } 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_first_vertex(b)); + nir_load_vertex_id_zero_base(b, bit_size), + nir_load_first_vertex(b, bit_size)); } else { - sysval = nir_load_vertex_id(b); + sysval = nir_load_vertex_id(b, bit_size); } break; @@ -140,14 +149,14 @@ convert_block(nir_block *block, nir_builder *b) */ if (b->shader->options->lower_base_vertex) sysval = nir_iand(b, - nir_load_is_indexed_draw(b), - nir_load_first_vertex(b)); + nir_load_is_indexed_draw(b, bit_size), + nir_load_first_vertex(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: @@ -172,8 +181,8 @@ convert_block(nir_block *block, nir_builder *b) break; case SYSTEM_VALUE_GLOBAL_GROUP_SIZE: { - nir_ssa_def *group_size = build_local_group_size(b); - nir_ssa_def *num_work_groups = nir_load_num_work_groups(b); + nir_ssa_def *group_size = nir_load_local_group_size(b, bit_size); + nir_ssa_def *num_work_groups = nir_load_num_work_groups(b, bit_size); sysval = nir_imul(b, group_size, num_work_groups); break; } @@ -185,7 +194,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_deref->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 b6f3529c766..f2151244e17 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 ecec3aa62d0..d71ae5284e4 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 1b31b564246..a4e46c54b4e 100644 --- a/src/gallium/auxiliary/nir/tgsi_to_nir.c +++ b/src/gallium/auxiliary/nir/tgsi_to_nir.c @@ -591,7 +591,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 f719aac1b86..1dc0f5f25e8 100644 --- a/src/intel/blorp/blorp_blit.c +++ b/src/intel/blorp/blorp_blit.c @@ -116,7 +116,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 b4c744020d9..a377fb5a212 100644 --- a/src/intel/blorp/blorp_clear.c +++ b/src/intel/blorp/blorp_clear.c @@ -967,7 +967,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 3b4642033fe..84a1d162607 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.17.1 _______________________________________________ mesa-dev mailing list <a href="mailto:mesa-dev@lists.freedesktop.org" target="_blank">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>