1/* 2 * Copyright (c) 2016 Intel Corporation 3 * 4 * Permission is hereby granted, free of charge, to any person obtaining a 5 * copy of this software and associated documentation files (the "Software"), 6 * to deal in the Software without restriction, including without limitation 7 * the rights to use, copy, modify, merge, publish, distribute, sublicense, 8 * and/or sell copies of the Software, and to permit persons to whom the 9 * Software is furnished to do so, subject to the following conditions: 10 * 11 * The above copyright notice and this permission notice (including the next 12 * paragraph) shall be included in all copies or substantial portions of the 13 * Software. 14 * 15 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR 16 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, 17 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL 18 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER 19 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING 20 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS 21 * IN THE SOFTWARE. 22 */ 23 24#include "brw_nir.h" 25#include "compiler/nir/nir_builder.h" 26 27struct lower_intrinsics_state { 28 nir_shader *nir; 29 nir_function_impl *impl; 30 bool progress; 31 nir_builder builder; 32}; 33 34static bool 35lower_cs_intrinsics_convert_block(struct lower_intrinsics_state *state, 36 nir_block *block) 37{ 38 bool progress = false; 39 nir_builder *b = &state->builder; 40 nir_shader *nir = state->nir; 41 42 /* Reuse calculated values inside the block. */ 43 nir_ssa_def *local_index = NULL; 44 nir_ssa_def *local_id = NULL; 45 46 nir_foreach_instr_safe(instr, block) { 47 if (instr->type != nir_instr_type_intrinsic) 48 continue; 49 50 nir_intrinsic_instr *intrinsic = nir_instr_as_intrinsic(instr); 51 52 b->cursor = nir_after_instr(&intrinsic->instr); 53 54 nir_ssa_def *sysval; 55 switch (intrinsic->intrinsic) { 56 case nir_intrinsic_load_workgroup_size: 57 case nir_intrinsic_load_workgroup_id: 58 case nir_intrinsic_load_num_workgroups: 59 /* Convert this to 32-bit if it's not */ 60 if (intrinsic->dest.ssa.bit_size == 64) { 61 intrinsic->dest.ssa.bit_size = 32; 62 sysval = nir_u2u64(b, &intrinsic->dest.ssa); 63 nir_ssa_def_rewrite_uses_after(&intrinsic->dest.ssa, 64 sysval, 65 sysval->parent_instr); 66 } 67 continue; 68 69 case nir_intrinsic_load_local_invocation_index: 70 case nir_intrinsic_load_local_invocation_id: { 71 /* First time we are using those, so let's calculate them. */ 72 if (!local_index) { 73 assert(!local_id); 74 75 nir_ssa_def *subgroup_id = nir_load_subgroup_id(b); 76 77 nir_ssa_def *thread_local_id = 78 nir_imul(b, subgroup_id, nir_load_simd_width_intel(b)); 79 nir_ssa_def *channel = nir_load_subgroup_invocation(b); 80 nir_ssa_def *linear = nir_iadd(b, channel, thread_local_id); 81 82 nir_ssa_def *size_x; 83 nir_ssa_def *size_y; 84 if (state->nir->info.workgroup_size_variable) { 85 nir_ssa_def *size_xyz = nir_load_workgroup_size(b); 86 size_x = nir_channel(b, size_xyz, 0); 87 size_y = nir_channel(b, size_xyz, 1); 88 } else { 89 size_x = nir_imm_int(b, nir->info.workgroup_size[0]); 90 size_y = nir_imm_int(b, nir->info.workgroup_size[1]); 91 } 92 nir_ssa_def *size_xy = nir_imul(b, size_x, size_y); 93 94 /* The local invocation index and ID must respect the following 95 * 96 * gl_LocalInvocationID.x = 97 * gl_LocalInvocationIndex % gl_WorkGroupSize.x; 98 * gl_LocalInvocationID.y = 99 * (gl_LocalInvocationIndex / gl_WorkGroupSize.x) % 100 * gl_WorkGroupSize.y; 101 * gl_LocalInvocationID.z = 102 * (gl_LocalInvocationIndex / 103 * (gl_WorkGroupSize.x * gl_WorkGroupSize.y)) % 104 * gl_WorkGroupSize.z; 105 * 106 * However, the final % gl_WorkGroupSize.z does nothing unless we 107 * accidentally end up with a gl_LocalInvocationIndex that is too 108 * large so it can safely be omitted. 109 */ 110 111 nir_ssa_def *id_x, *id_y, *id_z; 112 switch (state->nir->info.cs.derivative_group) { 113 case DERIVATIVE_GROUP_NONE: 114 if (nir->info.num_images == 0 && 115 nir->info.num_textures == 0) { 116 /* X-major lid order. Optimal for linear accesses only, 117 * which are usually buffers. X,Y ordering will look like: 118 * (0,0) (1,0) (2,0) ... (size_x-1,0) (0,1) (1,1) ... 119 */ 120 id_x = nir_umod(b, linear, size_x); 121 id_y = nir_umod(b, nir_udiv(b, linear, size_x), size_y); 122 local_index = linear; 123 } else if (!nir->info.workgroup_size_variable && 124 nir->info.workgroup_size[1] % 4 == 0) { 125 /* 1x4 block X-major lid order. Same as X-major except increments in 126 * blocks of width=1 height=4. Always optimal for tileY and usually 127 * optimal for linear accesses. 128 * x = (linear / 4) % size_x 129 * y = ((linear % 4) + (linear / 4 / size_x) * 4) % size_y 130 * X,Y ordering will look like: (0,0) (0,1) (0,2) (0,3) (1,0) (1,1) 131 * (1,2) (1,3) (2,0) ... (size_x-1,3) (0,4) (0,5) (0,6) (0,7) (1,4) ... 132 */ 133 const unsigned height = 4; 134 nir_ssa_def *block = nir_udiv_imm(b, linear, height); 135 id_x = nir_umod(b, block, size_x); 136 id_y = nir_umod(b, 137 nir_iadd(b, 138 nir_umod(b, linear, nir_imm_int(b, height)), 139 nir_imul_imm(b, 140 nir_udiv(b, block, size_x), 141 height)), 142 size_y); 143 } else { 144 /* Y-major lid order. Optimal for tileY accesses only, 145 * which are usually images. X,Y ordering will look like: 146 * (0,0) (0,1) (0,2) ... (0,size_y-1) (1,0) (1,1) ... 147 */ 148 id_y = nir_umod(b, linear, size_y); 149 id_x = nir_umod(b, nir_udiv(b, linear, size_y), size_x); 150 } 151 152 id_z = nir_udiv(b, linear, size_xy); 153 local_id = nir_vec3(b, id_x, id_y, id_z); 154 if (!local_index) { 155 local_index = nir_iadd(b, nir_iadd(b, id_x, 156 nir_imul(b, id_y, size_x)), 157 nir_imul(b, id_z, size_xy)); 158 } 159 break; 160 case DERIVATIVE_GROUP_LINEAR: 161 /* For linear, just set the local invocation index linearly, 162 * and calculate local invocation ID from that. 163 */ 164 id_x = nir_umod(b, linear, size_x); 165 id_y = nir_umod(b, nir_udiv(b, linear, size_x), size_y); 166 id_z = nir_udiv(b, linear, size_xy); 167 local_id = nir_vec3(b, id_x, id_y, id_z); 168 local_index = linear; 169 break; 170 case DERIVATIVE_GROUP_QUADS: { 171 /* For quads, first we figure out the 2x2 grid the invocation 172 * belongs to -- treating extra Z layers as just more rows. 173 * Then map that into local invocation ID (trivial) and local 174 * invocation index. Skipping Z simplify index calculation. 175 */ 176 177 nir_ssa_def *one = nir_imm_int(b, 1); 178 nir_ssa_def *double_size_x = nir_ishl(b, size_x, one); 179 180 /* ID within a pair of rows, where each group of 4 is 2x2 quad. */ 181 nir_ssa_def *row_pair_id = nir_umod(b, linear, double_size_x); 182 nir_ssa_def *y_row_pairs = nir_udiv(b, linear, double_size_x); 183 184 nir_ssa_def *x = 185 nir_ior(b, 186 nir_iand(b, row_pair_id, one), 187 nir_iand(b, nir_ishr(b, row_pair_id, one), 188 nir_imm_int(b, 0xfffffffe))); 189 nir_ssa_def *y = 190 nir_ior(b, 191 nir_ishl(b, y_row_pairs, one), 192 nir_iand(b, nir_ishr(b, row_pair_id, one), one)); 193 194 local_id = nir_vec3(b, x, 195 nir_umod(b, y, size_y), 196 nir_udiv(b, y, size_y)); 197 local_index = nir_iadd(b, x, nir_imul(b, y, size_x)); 198 break; 199 } 200 default: 201 unreachable("invalid derivative group"); 202 } 203 } 204 205 assert(local_id); 206 assert(local_index); 207 if (intrinsic->intrinsic == nir_intrinsic_load_local_invocation_id) 208 sysval = local_id; 209 else 210 sysval = local_index; 211 break; 212 } 213 214 case nir_intrinsic_load_num_subgroups: { 215 nir_ssa_def *size; 216 if (state->nir->info.workgroup_size_variable) { 217 nir_ssa_def *size_xyz = nir_load_workgroup_size(b); 218 nir_ssa_def *size_x = nir_channel(b, size_xyz, 0); 219 nir_ssa_def *size_y = nir_channel(b, size_xyz, 1); 220 nir_ssa_def *size_z = nir_channel(b, size_xyz, 2); 221 size = nir_imul(b, nir_imul(b, size_x, size_y), size_z); 222 } else { 223 size = nir_imm_int(b, nir->info.workgroup_size[0] * 224 nir->info.workgroup_size[1] * 225 nir->info.workgroup_size[2]); 226 } 227 228 /* Calculate the equivalent of DIV_ROUND_UP. */ 229 nir_ssa_def *simd_width = nir_load_simd_width_intel(b); 230 sysval = 231 nir_udiv(b, nir_iadd_imm(b, nir_iadd(b, size, simd_width), -1), 232 simd_width); 233 break; 234 } 235 236 default: 237 continue; 238 } 239 240 if (intrinsic->dest.ssa.bit_size == 64) 241 sysval = nir_u2u64(b, sysval); 242 243 nir_ssa_def_rewrite_uses(&intrinsic->dest.ssa, sysval); 244 nir_instr_remove(&intrinsic->instr); 245 246 state->progress = true; 247 } 248 249 return progress; 250} 251 252static void 253lower_cs_intrinsics_convert_impl(struct lower_intrinsics_state *state) 254{ 255 nir_builder_init(&state->builder, state->impl); 256 257 nir_foreach_block(block, state->impl) { 258 lower_cs_intrinsics_convert_block(state, block); 259 } 260 261 nir_metadata_preserve(state->impl, 262 nir_metadata_block_index | nir_metadata_dominance); 263} 264 265bool 266brw_nir_lower_cs_intrinsics(nir_shader *nir) 267{ 268 assert(nir->info.stage == MESA_SHADER_COMPUTE || 269 nir->info.stage == MESA_SHADER_KERNEL); 270 271 struct lower_intrinsics_state state = { 272 .nir = nir, 273 }; 274 275 /* Constraints from NV_compute_shader_derivatives. */ 276 if (!nir->info.workgroup_size_variable) { 277 if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_QUADS) { 278 assert(nir->info.workgroup_size[0] % 2 == 0); 279 assert(nir->info.workgroup_size[1] % 2 == 0); 280 } else if (nir->info.cs.derivative_group == DERIVATIVE_GROUP_LINEAR) { 281 ASSERTED unsigned workgroup_size = 282 nir->info.workgroup_size[0] * 283 nir->info.workgroup_size[1] * 284 nir->info.workgroup_size[2]; 285 assert(workgroup_size % 4 == 0); 286 } 287 } 288 289 nir_foreach_function(function, nir) { 290 if (function->impl) { 291 state.impl = function->impl; 292 lower_cs_intrinsics_convert_impl(&state); 293 } 294 } 295 296 return state.progress; 297} 298