case nir_intrinsic_load_num_workgroups:
info->uses_grid_size = true;
break;
- case nir_intrinsic_load_workgroup_size:
- /* The block size is translated to IMM with a fixed block size. */
- if (info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0)
- info->uses_block_size = true;
- break;
- case nir_intrinsic_load_local_invocation_id:
- case nir_intrinsic_load_workgroup_id: {
- unsigned mask = nir_ssa_def_components_read(&intr->dest.ssa);
- while (mask) {
- unsigned i = u_bit_scan(&mask);
-
- if (intr->intrinsic == nir_intrinsic_load_workgroup_id)
- info->uses_block_id[i] = true;
- else
- info->uses_thread_id[i] = true;
- }
- break;
- }
case nir_intrinsic_load_vertex_id:
info->uses_vertexid = 1;
break;
if (info->processor == PIPE_SHADER_COMPUTE &&
src->Register.File == TGSI_FILE_SYSTEM_VALUE) {
- unsigned name, mask;
+ unsigned name;
name = info->system_value_semantic_name[src->Register.Index];
switch (name) {
- case TGSI_SEMANTIC_THREAD_ID:
- case TGSI_SEMANTIC_BLOCK_ID:
- mask = usage_mask_after_swizzle & TGSI_WRITEMASK_XYZ;
- while (mask) {
- unsigned i = u_bit_scan(&mask);
-
- if (name == TGSI_SEMANTIC_THREAD_ID)
- info->uses_thread_id[i] = true;
- else
- info->uses_block_id[i] = true;
- }
- break;
- case TGSI_SEMANTIC_BLOCK_SIZE:
- /* The block size is translated to IMM with a fixed block size. */
- if (info->properties[TGSI_PROPERTY_CS_FIXED_BLOCK_WIDTH] == 0)
- info->uses_block_size = true;
- break;
case TGSI_SEMANTIC_GRID_SIZE:
info->uses_grid_size = true;
break;