1 /* 2 * Copyright 2012 Advanced Micro Devices, Inc. 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 "ac_shader_util.h" 25 #include "ac_gpu_info.h" 26 27 #include "sid.h" 28 #include "u_math.h" 29 30 #include <assert.h> 31 #include <stdlib.h> 32 #include <string.h> 33 34 unsigned ac_get_spi_shader_z_format(bool writes_z, bool writes_stencil, bool writes_samplemask) 35 { 36 if (writes_z) { 37 /* Z needs 32 bits. */ 38 if (writes_samplemask) 39 return V_028710_SPI_SHADER_32_ABGR; 40 else if (writes_stencil) 41 return V_028710_SPI_SHADER_32_GR; 42 else 43 return V_028710_SPI_SHADER_32_R; 44 } else if (writes_stencil || writes_samplemask) { 45 /* Both stencil and sample mask need only 16 bits. */ 46 return V_028710_SPI_SHADER_UINT16_ABGR; 47 } else { 48 return V_028710_SPI_SHADER_ZERO; 49 } 50 } 51 52 unsigned ac_get_cb_shader_mask(unsigned spi_shader_col_format) 53 { 54 unsigned i, cb_shader_mask = 0; 55 56 for (i = 0; i < 8; i++) { 57 switch ((spi_shader_col_format >> (i * 4)) & 0xf) { 58 case V_028714_SPI_SHADER_ZERO: 59 break; 60 case V_028714_SPI_SHADER_32_R: 61 cb_shader_mask |= 0x1 << (i * 4); 62 break; 63 case V_028714_SPI_SHADER_32_GR: 64 cb_shader_mask |= 0x3 << (i * 4); 65 break; 66 case V_028714_SPI_SHADER_32_AR: 67 cb_shader_mask |= 0x9u << (i * 4); 68 break; 69 case V_028714_SPI_SHADER_FP16_ABGR: 70 case V_028714_SPI_SHADER_UNORM16_ABGR: 71 case V_028714_SPI_SHADER_SNORM16_ABGR: 72 case V_028714_SPI_SHADER_UINT16_ABGR: 73 case V_028714_SPI_SHADER_SINT16_ABGR: 74 case V_028714_SPI_SHADER_32_ABGR: 75 cb_shader_mask |= 0xfu << (i * 4); 76 break; 77 default: 78 assert(0); 79 } 80 } 81 return cb_shader_mask; 82 } 83 84 /** 85 * Calculate the appropriate setting of VGT_GS_MODE when \p shader is a 86 * geometry shader. 87 */ 88 uint32_t ac_vgt_gs_mode(unsigned gs_max_vert_out, enum chip_class chip_class) 89 { 90 unsigned cut_mode; 91 92 if (gs_max_vert_out <= 128) { 93 cut_mode = V_028A40_GS_CUT_128; 94 } else if (gs_max_vert_out <= 256) { 95 cut_mode = V_028A40_GS_CUT_256; 96 } else if (gs_max_vert_out <= 512) { 97 cut_mode = V_028A40_GS_CUT_512; 98 } else { 99 assert(gs_max_vert_out <= 1024); 100 cut_mode = V_028A40_GS_CUT_1024; 101 } 102 103 return S_028A40_MODE(V_028A40_GS_SCENARIO_G) | S_028A40_CUT_MODE(cut_mode) | 104 S_028A40_ES_WRITE_OPTIMIZE(chip_class <= GFX8) | S_028A40_GS_WRITE_OPTIMIZE(1) | 105 S_028A40_ONCHIP(chip_class >= GFX9 ? 1 : 0); 106 } 107 108 /// Translate a (dfmt, nfmt) pair into a chip-appropriate combined format 109 /// value for LLVM8+ tbuffer intrinsics. 110 unsigned ac_get_tbuffer_format(enum chip_class chip_class, unsigned dfmt, unsigned nfmt) 111 { 112 // Some games try to access vertex buffers without a valid format. 113 // This is a game bug, but we should still handle it gracefully. 114 if (dfmt == V_008F0C_GFX10_FORMAT_INVALID) 115 return V_008F0C_GFX10_FORMAT_INVALID; 116 117 if (chip_class >= GFX10) { 118 unsigned format; 119 switch (dfmt) { 120 default: 121 unreachable("bad dfmt"); 122 case V_008F0C_BUF_DATA_FORMAT_INVALID: 123 format = V_008F0C_GFX10_FORMAT_INVALID; 124 break; 125 case V_008F0C_BUF_DATA_FORMAT_8: 126 format = V_008F0C_GFX10_FORMAT_8_UINT; 127 break; 128 case V_008F0C_BUF_DATA_FORMAT_8_8: 129 format = V_008F0C_GFX10_FORMAT_8_8_UINT; 130 break; 131 case V_008F0C_BUF_DATA_FORMAT_8_8_8_8: 132 format = V_008F0C_GFX10_FORMAT_8_8_8_8_UINT; 133 break; 134 case V_008F0C_BUF_DATA_FORMAT_16: 135 format = V_008F0C_GFX10_FORMAT_16_UINT; 136 break; 137 case V_008F0C_BUF_DATA_FORMAT_16_16: 138 format = V_008F0C_GFX10_FORMAT_16_16_UINT; 139 break; 140 case V_008F0C_BUF_DATA_FORMAT_16_16_16_16: 141 format = V_008F0C_GFX10_FORMAT_16_16_16_16_UINT; 142 break; 143 case V_008F0C_BUF_DATA_FORMAT_32: 144 format = V_008F0C_GFX10_FORMAT_32_UINT; 145 break; 146 case V_008F0C_BUF_DATA_FORMAT_32_32: 147 format = V_008F0C_GFX10_FORMAT_32_32_UINT; 148 break; 149 case V_008F0C_BUF_DATA_FORMAT_32_32_32: 150 format = V_008F0C_GFX10_FORMAT_32_32_32_UINT; 151 break; 152 case V_008F0C_BUF_DATA_FORMAT_32_32_32_32: 153 format = V_008F0C_GFX10_FORMAT_32_32_32_32_UINT; 154 break; 155 case V_008F0C_BUF_DATA_FORMAT_2_10_10_10: 156 format = V_008F0C_GFX10_FORMAT_2_10_10_10_UINT; 157 break; 158 case V_008F0C_BUF_DATA_FORMAT_10_11_11: 159 format = V_008F0C_GFX10_FORMAT_10_11_11_UINT; 160 break; 161 } 162 163 // Use the regularity properties of the combined format enum. 164 // 165 // Note: float is incompatible with 8-bit data formats, 166 // [us]{norm,scaled} are incomparible with 32-bit data formats. 167 // [us]scaled are not writable. 168 switch (nfmt) { 169 case V_008F0C_BUF_NUM_FORMAT_UNORM: 170 format -= 4; 171 break; 172 case V_008F0C_BUF_NUM_FORMAT_SNORM: 173 format -= 3; 174 break; 175 case V_008F0C_BUF_NUM_FORMAT_USCALED: 176 format -= 2; 177 break; 178 case V_008F0C_BUF_NUM_FORMAT_SSCALED: 179 format -= 1; 180 break; 181 default: 182 unreachable("bad nfmt"); 183 case V_008F0C_BUF_NUM_FORMAT_UINT: 184 break; 185 case V_008F0C_BUF_NUM_FORMAT_SINT: 186 format += 1; 187 break; 188 case V_008F0C_BUF_NUM_FORMAT_FLOAT: 189 format += 2; 190 break; 191 } 192 193 return format; 194 } else { 195 return dfmt | (nfmt << 4); 196 } 197 } 198 199 static const struct ac_data_format_info data_format_table[] = { 200 [V_008F0C_BUF_DATA_FORMAT_INVALID] = {0, 4, 0, V_008F0C_BUF_DATA_FORMAT_INVALID}, 201 [V_008F0C_BUF_DATA_FORMAT_8] = {1, 1, 1, V_008F0C_BUF_DATA_FORMAT_8}, 202 [V_008F0C_BUF_DATA_FORMAT_16] = {2, 1, 2, V_008F0C_BUF_DATA_FORMAT_16}, 203 [V_008F0C_BUF_DATA_FORMAT_8_8] = {2, 2, 1, V_008F0C_BUF_DATA_FORMAT_8}, 204 [V_008F0C_BUF_DATA_FORMAT_32] = {4, 1, 4, V_008F0C_BUF_DATA_FORMAT_32}, 205 [V_008F0C_BUF_DATA_FORMAT_16_16] = {4, 2, 2, V_008F0C_BUF_DATA_FORMAT_16}, 206 [V_008F0C_BUF_DATA_FORMAT_10_11_11] = {4, 3, 0, V_008F0C_BUF_DATA_FORMAT_10_11_11}, 207 [V_008F0C_BUF_DATA_FORMAT_11_11_10] = {4, 3, 0, V_008F0C_BUF_DATA_FORMAT_11_11_10}, 208 [V_008F0C_BUF_DATA_FORMAT_10_10_10_2] = {4, 4, 0, V_008F0C_BUF_DATA_FORMAT_10_10_10_2}, 209 [V_008F0C_BUF_DATA_FORMAT_2_10_10_10] = {4, 4, 0, V_008F0C_BUF_DATA_FORMAT_2_10_10_10}, 210 [V_008F0C_BUF_DATA_FORMAT_8_8_8_8] = {4, 4, 1, V_008F0C_BUF_DATA_FORMAT_8}, 211 [V_008F0C_BUF_DATA_FORMAT_32_32] = {8, 2, 4, V_008F0C_BUF_DATA_FORMAT_32}, 212 [V_008F0C_BUF_DATA_FORMAT_16_16_16_16] = {8, 4, 2, V_008F0C_BUF_DATA_FORMAT_16}, 213 [V_008F0C_BUF_DATA_FORMAT_32_32_32] = {12, 3, 4, V_008F0C_BUF_DATA_FORMAT_32}, 214 [V_008F0C_BUF_DATA_FORMAT_32_32_32_32] = {16, 4, 4, V_008F0C_BUF_DATA_FORMAT_32}, 215 }; 216 217 const struct ac_data_format_info *ac_get_data_format_info(unsigned dfmt) 218 { 219 assert(dfmt < ARRAY_SIZE(data_format_table)); 220 return &data_format_table[dfmt]; 221 } 222 223 enum ac_image_dim ac_get_sampler_dim(enum chip_class chip_class, enum glsl_sampler_dim dim, 224 bool is_array) 225 { 226 switch (dim) { 227 case GLSL_SAMPLER_DIM_1D: 228 if (chip_class == GFX9) 229 return is_array ? ac_image_2darray : ac_image_2d; 230 return is_array ? ac_image_1darray : ac_image_1d; 231 case GLSL_SAMPLER_DIM_2D: 232 case GLSL_SAMPLER_DIM_RECT: 233 case GLSL_SAMPLER_DIM_EXTERNAL: 234 return is_array ? ac_image_2darray : ac_image_2d; 235 case GLSL_SAMPLER_DIM_3D: 236 return ac_image_3d; 237 case GLSL_SAMPLER_DIM_CUBE: 238 return ac_image_cube; 239 case GLSL_SAMPLER_DIM_MS: 240 return is_array ? ac_image_2darraymsaa : ac_image_2dmsaa; 241 case GLSL_SAMPLER_DIM_SUBPASS: 242 return ac_image_2darray; 243 case GLSL_SAMPLER_DIM_SUBPASS_MS: 244 return ac_image_2darraymsaa; 245 default: 246 unreachable("bad sampler dim"); 247 } 248 } 249 250 enum ac_image_dim ac_get_image_dim(enum chip_class chip_class, enum glsl_sampler_dim sdim, 251 bool is_array) 252 { 253 enum ac_image_dim dim = ac_get_sampler_dim(chip_class, sdim, is_array); 254 255 /* Match the resource type set in the descriptor. */ 256 if (dim == ac_image_cube || (chip_class <= GFX8 && dim == ac_image_3d)) 257 dim = ac_image_2darray; 258 else if (sdim == GLSL_SAMPLER_DIM_2D && !is_array && chip_class == GFX9) { 259 /* When a single layer of a 3D texture is bound, the shader 260 * will refer to a 2D target, but the descriptor has a 3D type. 261 * Since the HW ignores BASE_ARRAY in this case, we need to 262 * send 3 coordinates. This doesn't hurt when the underlying 263 * texture is non-3D. 264 */ 265 dim = ac_image_3d; 266 } 267 268 return dim; 269 } 270 271 unsigned ac_get_fs_input_vgpr_cnt(const struct ac_shader_config *config, 272 signed char *face_vgpr_index_ptr, 273 signed char *ancillary_vgpr_index_ptr) 274 { 275 unsigned num_input_vgprs = 0; 276 signed char face_vgpr_index = -1; 277 signed char ancillary_vgpr_index = -1; 278 279 if (G_0286CC_PERSP_SAMPLE_ENA(config->spi_ps_input_addr)) 280 num_input_vgprs += 2; 281 if (G_0286CC_PERSP_CENTER_ENA(config->spi_ps_input_addr)) 282 num_input_vgprs += 2; 283 if (G_0286CC_PERSP_CENTROID_ENA(config->spi_ps_input_addr)) 284 num_input_vgprs += 2; 285 if (G_0286CC_PERSP_PULL_MODEL_ENA(config->spi_ps_input_addr)) 286 num_input_vgprs += 3; 287 if (G_0286CC_LINEAR_SAMPLE_ENA(config->spi_ps_input_addr)) 288 num_input_vgprs += 2; 289 if (G_0286CC_LINEAR_CENTER_ENA(config->spi_ps_input_addr)) 290 num_input_vgprs += 2; 291 if (G_0286CC_LINEAR_CENTROID_ENA(config->spi_ps_input_addr)) 292 num_input_vgprs += 2; 293 if (G_0286CC_LINE_STIPPLE_TEX_ENA(config->spi_ps_input_addr)) 294 num_input_vgprs += 1; 295 if (G_0286CC_POS_X_FLOAT_ENA(config->spi_ps_input_addr)) 296 num_input_vgprs += 1; 297 if (G_0286CC_POS_Y_FLOAT_ENA(config->spi_ps_input_addr)) 298 num_input_vgprs += 1; 299 if (G_0286CC_POS_Z_FLOAT_ENA(config->spi_ps_input_addr)) 300 num_input_vgprs += 1; 301 if (G_0286CC_POS_W_FLOAT_ENA(config->spi_ps_input_addr)) 302 num_input_vgprs += 1; 303 if (G_0286CC_FRONT_FACE_ENA(config->spi_ps_input_addr)) { 304 face_vgpr_index = num_input_vgprs; 305 num_input_vgprs += 1; 306 } 307 if (G_0286CC_ANCILLARY_ENA(config->spi_ps_input_addr)) { 308 ancillary_vgpr_index = num_input_vgprs; 309 num_input_vgprs += 1; 310 } 311 if (G_0286CC_SAMPLE_COVERAGE_ENA(config->spi_ps_input_addr)) 312 num_input_vgprs += 1; 313 if (G_0286CC_POS_FIXED_PT_ENA(config->spi_ps_input_addr)) 314 num_input_vgprs += 1; 315 316 if (face_vgpr_index_ptr) 317 *face_vgpr_index_ptr = face_vgpr_index; 318 if (ancillary_vgpr_index_ptr) 319 *ancillary_vgpr_index_ptr = ancillary_vgpr_index; 320 321 return num_input_vgprs; 322 } 323 324 void ac_choose_spi_color_formats(unsigned format, unsigned swap, unsigned ntype, 325 bool is_depth, bool use_rbplus, 326 struct ac_spi_color_formats *formats) 327 { 328 /* Alpha is needed for alpha-to-coverage. 329 * Blending may be with or without alpha. 330 */ 331 unsigned normal = 0; /* most optimal, may not support blending or export alpha */ 332 unsigned alpha = 0; /* exports alpha, but may not support blending */ 333 unsigned blend = 0; /* supports blending, but may not export alpha */ 334 unsigned blend_alpha = 0; /* least optimal, supports blending and exports alpha */ 335 336 /* Choose the SPI color formats. These are required values for RB+. 337 * Other chips have multiple choices, though they are not necessarily better. 338 */ 339 switch (format) { 340 case V_028C70_COLOR_5_6_5: 341 case V_028C70_COLOR_1_5_5_5: 342 case V_028C70_COLOR_5_5_5_1: 343 case V_028C70_COLOR_4_4_4_4: 344 case V_028C70_COLOR_10_11_11: 345 case V_028C70_COLOR_11_11_10: 346 case V_028C70_COLOR_5_9_9_9: 347 case V_028C70_COLOR_8: 348 case V_028C70_COLOR_8_8: 349 case V_028C70_COLOR_8_8_8_8: 350 case V_028C70_COLOR_10_10_10_2: 351 case V_028C70_COLOR_2_10_10_10: 352 if (ntype == V_028C70_NUMBER_UINT) 353 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_UINT16_ABGR; 354 else if (ntype == V_028C70_NUMBER_SINT) 355 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_SINT16_ABGR; 356 else 357 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_FP16_ABGR; 358 359 if (!use_rbplus && format == V_028C70_COLOR_8 && 360 ntype != V_028C70_NUMBER_SRGB && swap == V_028C70_SWAP_STD) /* R */ { 361 /* When RB+ is enabled, R8_UNORM should use FP16_ABGR for 2x 362 * exporting performance. Otherwise, use 32_R to remove useless 363 * instructions needed for 16-bit compressed exports. 364 */ 365 blend = normal = V_028714_SPI_SHADER_32_R; 366 } 367 break; 368 369 case V_028C70_COLOR_16: 370 case V_028C70_COLOR_16_16: 371 case V_028C70_COLOR_16_16_16_16: 372 if (ntype == V_028C70_NUMBER_UNORM || ntype == V_028C70_NUMBER_SNORM) { 373 /* UNORM16 and SNORM16 don't support blending */ 374 if (ntype == V_028C70_NUMBER_UNORM) 375 normal = alpha = V_028714_SPI_SHADER_UNORM16_ABGR; 376 else 377 normal = alpha = V_028714_SPI_SHADER_SNORM16_ABGR; 378 379 /* Use 32 bits per channel for blending. */ 380 if (format == V_028C70_COLOR_16) { 381 if (swap == V_028C70_SWAP_STD) { /* R */ 382 blend = V_028714_SPI_SHADER_32_R; 383 blend_alpha = V_028714_SPI_SHADER_32_AR; 384 } else if (swap == V_028C70_SWAP_ALT_REV) /* A */ 385 blend = blend_alpha = V_028714_SPI_SHADER_32_AR; 386 else 387 assert(0); 388 } else if (format == V_028C70_COLOR_16_16) { 389 if (swap == V_028C70_SWAP_STD) { /* RG */ 390 blend = V_028714_SPI_SHADER_32_GR; 391 blend_alpha = V_028714_SPI_SHADER_32_ABGR; 392 } else if (swap == V_028C70_SWAP_ALT) /* RA */ 393 blend = blend_alpha = V_028714_SPI_SHADER_32_AR; 394 else 395 assert(0); 396 } else /* 16_16_16_16 */ 397 blend = blend_alpha = V_028714_SPI_SHADER_32_ABGR; 398 } else if (ntype == V_028C70_NUMBER_UINT) 399 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_UINT16_ABGR; 400 else if (ntype == V_028C70_NUMBER_SINT) 401 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_SINT16_ABGR; 402 else if (ntype == V_028C70_NUMBER_FLOAT) 403 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_FP16_ABGR; 404 else 405 assert(0); 406 break; 407 408 case V_028C70_COLOR_32: 409 if (swap == V_028C70_SWAP_STD) { /* R */ 410 blend = normal = V_028714_SPI_SHADER_32_R; 411 alpha = blend_alpha = V_028714_SPI_SHADER_32_AR; 412 } else if (swap == V_028C70_SWAP_ALT_REV) /* A */ 413 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_32_AR; 414 else 415 assert(0); 416 break; 417 418 case V_028C70_COLOR_32_32: 419 if (swap == V_028C70_SWAP_STD) { /* RG */ 420 blend = normal = V_028714_SPI_SHADER_32_GR; 421 alpha = blend_alpha = V_028714_SPI_SHADER_32_ABGR; 422 } else if (swap == V_028C70_SWAP_ALT) /* RA */ 423 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_32_AR; 424 else 425 assert(0); 426 break; 427 428 case V_028C70_COLOR_32_32_32_32: 429 case V_028C70_COLOR_8_24: 430 case V_028C70_COLOR_24_8: 431 case V_028C70_COLOR_X24_8_32_FLOAT: 432 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_32_ABGR; 433 break; 434 435 default: 436 assert(0); 437 return; 438 } 439 440 /* The DB->CB copy needs 32_ABGR. */ 441 if (is_depth) 442 alpha = blend = blend_alpha = normal = V_028714_SPI_SHADER_32_ABGR; 443 444 formats->normal = normal; 445 formats->alpha = alpha; 446 formats->blend = blend; 447 formats->blend_alpha = blend_alpha; 448 } 449 450 void ac_compute_late_alloc(const struct radeon_info *info, bool ngg, bool ngg_culling, 451 bool uses_scratch, unsigned *late_alloc_wave64, unsigned *cu_mask) 452 { 453 *late_alloc_wave64 = 0; /* The limit is per SA. */ 454 *cu_mask = 0xffff; 455 456 /* CU masking can decrease performance and cause a hang with <= 2 CUs per SA. */ 457 if (info->min_good_cu_per_sa <= 2) 458 return; 459 460 /* If scratch is used with late alloc, the GPU could deadlock if PS uses scratch too. A more 461 * complicated computation is needed to enable late alloc with scratch (see PAL). 462 */ 463 if (uses_scratch) 464 return; 465 466 /* Late alloc is not used for NGG on Navi14 due to a hw bug. */ 467 if (ngg && info->family == CHIP_NAVI14) 468 return; 469 470 if (info->chip_class >= GFX10) { 471 /* For Wave32, the hw will launch twice the number of late alloc waves, so 1 == 2x wave32. 472 * These limits are estimated because they are all safe but they vary in performance. 473 */ 474 if (ngg_culling) 475 *late_alloc_wave64 = info->min_good_cu_per_sa * 10; 476 else 477 *late_alloc_wave64 = info->min_good_cu_per_sa * 4; 478 479 /* Limit LATE_ALLOC_GS to prevent a hang (hw bug) on gfx10. */ 480 if (info->chip_class == GFX10 && ngg) 481 *late_alloc_wave64 = MIN2(*late_alloc_wave64, 64); 482 483 /* Gfx10: CU2 & CU3 must be disabled to prevent a hw deadlock. 484 * Others: CU1 must be disabled to prevent a hw deadlock. 485 * 486 * The deadlock is caused by late alloc, which usually increases performance. 487 */ 488 *cu_mask &= info->chip_class == GFX10 ? ~BITFIELD_RANGE(2, 2) : 489 ~BITFIELD_RANGE(1, 1); 490 } else { 491 if (info->min_good_cu_per_sa <= 4) { 492 /* Too few available compute units per SA. Disallowing VS to run on one CU could hurt us 493 * more than late VS allocation would help. 494 * 495 * 2 is the highest safe number that allows us to keep all CUs enabled. 496 */ 497 *late_alloc_wave64 = 2; 498 } else { 499 /* This is a good initial value, allowing 1 late_alloc wave per SIMD on num_cu - 2. 500 */ 501 *late_alloc_wave64 = (info->min_good_cu_per_sa - 2) * 4; 502 } 503 504 /* VS can't execute on one CU if the limit is > 2. */ 505 if (*late_alloc_wave64 > 2) 506 *cu_mask = 0xfffe; /* 1 CU disabled */ 507 } 508 509 /* Max number that fits into the register field. */ 510 if (ngg) /* GS */ 511 *late_alloc_wave64 = MIN2(*late_alloc_wave64, G_00B204_SPI_SHADER_LATE_ALLOC_GS_GFX10(~0u)); 512 else /* VS */ 513 *late_alloc_wave64 = MIN2(*late_alloc_wave64, G_00B11C_LIMIT(~0u)); 514 } 515 516 unsigned ac_compute_cs_workgroup_size(uint16_t sizes[3], bool variable, unsigned max) 517 { 518 if (variable) 519 return max; 520 521 return sizes[0] * sizes[1] * sizes[2]; 522 } 523 524 unsigned ac_compute_lshs_workgroup_size(enum chip_class chip_class, gl_shader_stage stage, 525 unsigned tess_num_patches, 526 unsigned tess_patch_in_vtx, 527 unsigned tess_patch_out_vtx) 528 { 529 /* When tessellation is used, API VS runs on HW LS, API TCS runs on HW HS. 530 * These two HW stages are merged on GFX9+. 531 */ 532 533 bool merged_shaders = chip_class >= GFX9; 534 unsigned ls_workgroup_size = tess_num_patches * tess_patch_in_vtx; 535 unsigned hs_workgroup_size = tess_num_patches * tess_patch_out_vtx; 536 537 if (merged_shaders) 538 return MAX2(ls_workgroup_size, hs_workgroup_size); 539 else if (stage == MESA_SHADER_VERTEX) 540 return ls_workgroup_size; 541 else if (stage == MESA_SHADER_TESS_CTRL) 542 return hs_workgroup_size; 543 else 544 unreachable("invalid LSHS shader stage"); 545 } 546 547 unsigned ac_compute_esgs_workgroup_size(enum chip_class chip_class, unsigned wave_size, 548 unsigned es_verts, unsigned gs_inst_prims) 549 { 550 /* ESGS may operate in workgroups if on-chip GS (LDS rings) are enabled. 551 * 552 * GFX6: Not possible in the HW. 553 * GFX7-8 (unmerged): possible in the HW, but not implemented in Mesa. 554 * GFX9+ (merged): implemented in Mesa. 555 */ 556 557 if (chip_class <= GFX8) 558 return wave_size; 559 560 unsigned workgroup_size = MAX2(es_verts, gs_inst_prims); 561 return CLAMP(workgroup_size, 1, 256); 562 } 563 564 unsigned ac_compute_ngg_workgroup_size(unsigned es_verts, unsigned gs_inst_prims, 565 unsigned max_vtx_out, unsigned prim_amp_factor) 566 { 567 /* NGG always operates in workgroups. 568 * 569 * For API VS/TES/GS: 570 * - 1 invocation per input vertex 571 * - 1 invocation per input primitive 572 * 573 * The same invocation can process both an input vertex and primitive, 574 * however 1 invocation can only output up to 1 vertex and 1 primitive. 575 */ 576 577 unsigned max_vtx_in = es_verts < 256 ? es_verts : 3 * gs_inst_prims; 578 unsigned max_prim_in = gs_inst_prims; 579 unsigned max_prim_out = gs_inst_prims * prim_amp_factor; 580 unsigned workgroup_size = MAX4(max_vtx_in, max_vtx_out, max_prim_in, max_prim_out); 581 582 return CLAMP(workgroup_size, 1, 256); 583 } 584
Indexes created Sat Aug 01 00:24:56 UTC 2026