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 <kher...@redhat.com>
---
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) \
{ \
- 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)),
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
mesa-dev@lists.freedesktop.org
https://lists.freedesktop.org/mailman/listinfo/mesa-dev
