nir_def *y = nir_channel(b, id, 1);
nir_def *z = nir_channel(b, id, 2);
- nir_def *grid_size = nir_load_num_workgroups(b, 32);
+ nir_def *grid_size = nir_load_num_workgroups(b);
nir_def *grid_size_x = nir_channel(b, grid_size, 0);
nir_def *grid_size_y = nir_channel(b, grid_size, 1);
system_value("workgroup_index", 1)
system_value("base_workgroup_id", 3, bit_sizes=[32, 64])
system_value("user_clip_plane", 4, indices=[UCP_ID])
-system_value("num_workgroups", 3, bit_sizes=[32, 64])
+system_value("num_workgroups", 3)
system_value("num_vertices", 1)
system_value("helper_invocation", 1, bit_sizes=[1, 32])
system_value("layer_id", 1)
build_global_group_size(nir_builder *b, unsigned bit_size)
{
nir_def *group_size = nir_load_workgroup_size(b);
- nir_def *num_workgroups = nir_load_num_workgroups(b, bit_size);
+ nir_def *num_workgroups = nir_load_num_workgroups(b);
return nir_imul(b, nir_u2uN(b, group_size, bit_size),
- num_workgroups);
+ nir_u2uN(b, num_workgroups, bit_size));
}
static bool
case nir_intrinsic_load_local_invocation_id:
case nir_intrinsic_load_local_invocation_index:
+ case nir_intrinsic_load_num_workgroups:
case nir_intrinsic_load_workgroup_id:
case nir_intrinsic_load_workgroup_size:
return sanitize_32bit_sysval(b, intrin);
if (val)
return val;
+ nir_def *num_workgroups = nir_load_num_workgroups(b);
return lower_id_to_index_no_umod(b, wg_idx,
- nir_load_num_workgroups(b, bit_size),
+ nir_u2uN(b, num_workgroups, bit_size),
bit_size,
options->num_workgroups,
options->shortcut_1d_workgroup_id);