Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1 | /* |
| 2 | * Copyright © 2016 Red Hat. |
| 3 | * Copyright © 2016 Bas Nieuwenhuizen |
| 4 | * |
| 5 | * based in part on anv driver which is: |
| 6 | * Copyright © 2015 Intel Corporation |
| 7 | * |
| 8 | * Permission is hereby granted, free of charge, to any person obtaining a |
| 9 | * copy of this software and associated documentation files (the "Software"), |
| 10 | * to deal in the Software without restriction, including without limitation |
| 11 | * the rights to use, copy, modify, merge, publish, distribute, sublicense, |
| 12 | * and/or sell copies of the Software, and to permit persons to whom the |
| 13 | * Software is furnished to do so, subject to the following conditions: |
| 14 | * |
| 15 | * The above copyright notice and this permission notice (including the next |
| 16 | * paragraph) shall be included in all copies or substantial portions of the |
| 17 | * Software. |
| 18 | * |
| 19 | * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR |
| 20 | * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, |
| 21 | * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL |
| 22 | * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER |
| 23 | * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING |
| 24 | * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS |
| 25 | * IN THE SOFTWARE. |
| 26 | */ |
| 27 | |
| 28 | #include "util/mesa-sha1.h" |
| 29 | #include "util/u_atomic.h" |
| 30 | #include "radv_debug.h" |
| 31 | #include "radv_private.h" |
| 32 | #include "radv_shader.h" |
Dave Airlie | 6f3aee4 | 2018-06-27 11:34:25 +1000 | [diff] [blame] | 33 | #include "radv_shader_helper.h" |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 34 | #include "nir/nir.h" |
| 35 | #include "nir/nir_builder.h" |
| 36 | #include "spirv/nir_spirv.h" |
| 37 | |
| 38 | #include <llvm-c/Core.h> |
| 39 | #include <llvm-c/TargetMachine.h> |
Samuel Pitoiset | 135e4d4 | 2018-06-08 11:38:01 +0200 | [diff] [blame] | 40 | #include <llvm-c/Support.h> |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 41 | |
| 42 | #include "sid.h" |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 43 | #include "ac_binary.h" |
| 44 | #include "ac_llvm_util.h" |
| 45 | #include "ac_nir_to_llvm.h" |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 46 | #include "ac_rtld.h" |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 47 | #include "vk_format.h" |
| 48 | #include "util/debug.h" |
| 49 | #include "ac_exp_param.h" |
| 50 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 51 | #include "aco_interface.h" |
| 52 | |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 53 | #include "util/string_buffer.h" |
| 54 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 55 | static const struct nir_shader_compiler_options nir_options_llvm = { |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 56 | .vertex_id_zero_based = true, |
| 57 | .lower_scmp = true, |
Rhys Perry | 0af95f0 | 2018-12-06 14:01:15 +0000 | [diff] [blame] | 58 | .lower_flrp16 = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 59 | .lower_flrp32 = true, |
Timothy Arceri | f0d74ec | 2018-01-12 11:12:09 +1100 | [diff] [blame] | 60 | .lower_flrp64 = true, |
Bas Nieuwenhuizen | 5240fdd | 2018-01-21 17:13:26 +0100 | [diff] [blame] | 61 | .lower_device_index_to_zero = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 62 | .lower_fsat = true, |
| 63 | .lower_fdiv = true, |
Samuel Pitoiset | 5ebe1a1 | 2019-10-03 16:20:40 +0200 | [diff] [blame] | 64 | .lower_fmod = true, |
Daniel Schürmann | 48a75e7 | 2019-01-25 16:08:38 +0100 | [diff] [blame] | 65 | .lower_bitfield_insert_to_bitfield_select = true, |
Daniel Schürmann | 0daeb1d | 2019-01-25 16:24:55 +0100 | [diff] [blame] | 66 | .lower_bitfield_extract = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 67 | .lower_sub = true, |
| 68 | .lower_pack_snorm_2x16 = true, |
| 69 | .lower_pack_snorm_4x8 = true, |
| 70 | .lower_pack_unorm_2x16 = true, |
| 71 | .lower_pack_unorm_4x8 = true, |
| 72 | .lower_unpack_snorm_2x16 = true, |
| 73 | .lower_unpack_snorm_4x8 = true, |
| 74 | .lower_unpack_unorm_2x16 = true, |
| 75 | .lower_unpack_unorm_4x8 = true, |
| 76 | .lower_extract_byte = true, |
| 77 | .lower_extract_word = true, |
Dave Airlie | 2c61594 | 2017-10-04 06:33:02 +1000 | [diff] [blame] | 78 | .lower_ffma = true, |
Samuel Pitoiset | 7aa008d | 2018-02-02 19:04:57 +0100 | [diff] [blame] | 79 | .lower_fpow = true, |
Samuel Pitoiset | 71ffa00 | 2019-03-06 22:35:31 +0100 | [diff] [blame] | 80 | .lower_mul_2x32_64 = true, |
Sagar Ghuge | 456557a | 2019-06-03 17:11:57 -0700 | [diff] [blame] | 81 | .lower_rotate = true, |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 82 | .max_unroll_iterations = 32, |
| 83 | .use_interpolated_input_intrinsics = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 84 | }; |
| 85 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 86 | static const struct nir_shader_compiler_options nir_options_aco = { |
| 87 | .vertex_id_zero_based = true, |
| 88 | .lower_scmp = true, |
| 89 | .lower_flrp16 = true, |
| 90 | .lower_flrp32 = true, |
| 91 | .lower_flrp64 = true, |
| 92 | .lower_device_index_to_zero = true, |
| 93 | .lower_fdiv = true, |
Rhys Perry | a87b0f5 | 2019-10-03 15:32:19 +0100 | [diff] [blame] | 94 | .lower_fmod = true, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 95 | .lower_bitfield_insert_to_bitfield_select = true, |
| 96 | .lower_bitfield_extract = true, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 97 | .lower_pack_snorm_2x16 = true, |
| 98 | .lower_pack_snorm_4x8 = true, |
| 99 | .lower_pack_unorm_2x16 = true, |
| 100 | .lower_pack_unorm_4x8 = true, |
| 101 | .lower_unpack_snorm_2x16 = true, |
| 102 | .lower_unpack_snorm_4x8 = true, |
| 103 | .lower_unpack_unorm_2x16 = true, |
| 104 | .lower_unpack_unorm_4x8 = true, |
| 105 | .lower_unpack_half_2x16 = true, |
| 106 | .lower_extract_byte = true, |
| 107 | .lower_extract_word = true, |
| 108 | .lower_ffma = true, |
| 109 | .lower_fpow = true, |
| 110 | .lower_mul_2x32_64 = true, |
| 111 | .lower_rotate = true, |
| 112 | .max_unroll_iterations = 32, |
| 113 | .use_interpolated_input_intrinsics = true, |
| 114 | }; |
| 115 | |
Daniel Schürmann | 45638e1 | 2019-07-29 17:51:01 +0200 | [diff] [blame] | 116 | bool |
| 117 | radv_can_dump_shader(struct radv_device *device, |
| 118 | struct radv_shader_module *module, |
| 119 | bool is_gs_copy_shader) |
| 120 | { |
| 121 | if (!(device->instance->debug_flags & RADV_DEBUG_DUMP_SHADERS)) |
| 122 | return false; |
Timur Kristóf | 30f0c0e | 2019-09-18 14:39:10 +0200 | [diff] [blame] | 123 | if (module) |
| 124 | return !module->nir || |
| 125 | (device->instance->debug_flags & RADV_DEBUG_DUMP_META_SHADERS); |
Daniel Schürmann | 45638e1 | 2019-07-29 17:51:01 +0200 | [diff] [blame] | 126 | |
Timur Kristóf | 30f0c0e | 2019-09-18 14:39:10 +0200 | [diff] [blame] | 127 | return is_gs_copy_shader; |
Daniel Schürmann | 45638e1 | 2019-07-29 17:51:01 +0200 | [diff] [blame] | 128 | } |
| 129 | |
| 130 | bool |
| 131 | radv_can_dump_shader_stats(struct radv_device *device, |
| 132 | struct radv_shader_module *module) |
| 133 | { |
| 134 | /* Only dump non-meta shader stats. */ |
| 135 | return device->instance->debug_flags & RADV_DEBUG_DUMP_SHADER_STATS && |
| 136 | module && !module->nir; |
| 137 | } |
| 138 | |
| 139 | unsigned shader_io_get_unique_index(gl_varying_slot slot) |
| 140 | { |
| 141 | /* handle patch indices separate */ |
| 142 | if (slot == VARYING_SLOT_TESS_LEVEL_OUTER) |
| 143 | return 0; |
| 144 | if (slot == VARYING_SLOT_TESS_LEVEL_INNER) |
| 145 | return 1; |
| 146 | if (slot >= VARYING_SLOT_PATCH0 && slot <= VARYING_SLOT_TESS_MAX) |
| 147 | return 2 + (slot - VARYING_SLOT_PATCH0); |
| 148 | if (slot == VARYING_SLOT_POS) |
| 149 | return 0; |
| 150 | if (slot == VARYING_SLOT_PSIZ) |
| 151 | return 1; |
| 152 | if (slot == VARYING_SLOT_CLIP_DIST0) |
| 153 | return 2; |
| 154 | if (slot == VARYING_SLOT_CLIP_DIST1) |
| 155 | return 3; |
| 156 | /* 3 is reserved for clip dist as well */ |
| 157 | if (slot >= VARYING_SLOT_VAR0 && slot <= VARYING_SLOT_VAR31) |
| 158 | return 4 + (slot - VARYING_SLOT_VAR0); |
| 159 | unreachable("illegal slot in get unique index\n"); |
| 160 | } |
| 161 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 162 | VkResult radv_CreateShaderModule( |
| 163 | VkDevice _device, |
| 164 | const VkShaderModuleCreateInfo* pCreateInfo, |
| 165 | const VkAllocationCallbacks* pAllocator, |
| 166 | VkShaderModule* pShaderModule) |
| 167 | { |
| 168 | RADV_FROM_HANDLE(radv_device, device, _device); |
| 169 | struct radv_shader_module *module; |
| 170 | |
| 171 | assert(pCreateInfo->sType == VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO); |
| 172 | assert(pCreateInfo->flags == 0); |
| 173 | |
| 174 | module = vk_alloc2(&device->alloc, pAllocator, |
| 175 | sizeof(*module) + pCreateInfo->codeSize, 8, |
| 176 | VK_SYSTEM_ALLOCATION_SCOPE_OBJECT); |
| 177 | if (module == NULL) |
Bas Nieuwenhuizen | 38933c1 | 2018-05-31 01:06:41 +0200 | [diff] [blame] | 178 | return vk_error(device->instance, VK_ERROR_OUT_OF_HOST_MEMORY); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 179 | |
| 180 | module->nir = NULL; |
| 181 | module->size = pCreateInfo->codeSize; |
| 182 | memcpy(module->data, pCreateInfo->pCode, module->size); |
| 183 | |
| 184 | _mesa_sha1_compute(module->data, module->size, module->sha1); |
| 185 | |
| 186 | *pShaderModule = radv_shader_module_to_handle(module); |
| 187 | |
| 188 | return VK_SUCCESS; |
| 189 | } |
| 190 | |
| 191 | void radv_DestroyShaderModule( |
| 192 | VkDevice _device, |
| 193 | VkShaderModule _module, |
| 194 | const VkAllocationCallbacks* pAllocator) |
| 195 | { |
| 196 | RADV_FROM_HANDLE(radv_device, device, _device); |
| 197 | RADV_FROM_HANDLE(radv_shader_module, module, _module); |
| 198 | |
| 199 | if (!module) |
| 200 | return; |
| 201 | |
| 202 | vk_free2(&device->alloc, pAllocator, module); |
| 203 | } |
| 204 | |
Bas Nieuwenhuizen | 06f0504 | 2017-02-09 00:12:10 +0100 | [diff] [blame] | 205 | void |
Timothy Arceri | 0667571 | 2018-10-18 09:42:17 +1100 | [diff] [blame] | 206 | radv_optimize_nir(struct nir_shader *shader, bool optimize_conservatively, |
| 207 | bool allow_copies) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 208 | { |
| 209 | bool progress; |
Ian Romanick | d41cdef | 2018-08-18 16:42:04 -0700 | [diff] [blame] | 210 | unsigned lower_flrp = |
| 211 | (shader->options->lower_flrp16 ? 16 : 0) | |
| 212 | (shader->options->lower_flrp32 ? 32 : 0) | |
| 213 | (shader->options->lower_flrp64 ? 64 : 0); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 214 | |
| 215 | do { |
| 216 | progress = false; |
| 217 | |
Karol Herbst | 9b24028 | 2019-01-16 00:05:04 +0100 | [diff] [blame] | 218 | NIR_PASS(progress, shader, nir_split_array_vars, nir_var_function_temp); |
| 219 | NIR_PASS(progress, shader, nir_shrink_vec_array_vars, nir_var_function_temp); |
Timothy Arceri | 8086fa1 | 2018-10-18 10:19:16 +1100 | [diff] [blame] | 220 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 221 | NIR_PASS_V(shader, nir_lower_vars_to_ssa); |
Iago Toral Quiroga | 2d648e5 | 2018-04-27 09:28:48 +0200 | [diff] [blame] | 222 | NIR_PASS_V(shader, nir_lower_pack); |
Timothy Arceri | 9d5b106 | 2018-10-18 08:55:46 +1100 | [diff] [blame] | 223 | |
Timothy Arceri | 0667571 | 2018-10-18 09:42:17 +1100 | [diff] [blame] | 224 | if (allow_copies) { |
| 225 | /* Only run this pass in the first call to |
| 226 | * radv_optimize_nir. Later calls assume that we've |
| 227 | * lowered away any copy_deref instructions and we |
| 228 | * don't want to introduce any more. |
| 229 | */ |
| 230 | NIR_PASS(progress, shader, nir_opt_find_array_copies); |
| 231 | } |
| 232 | |
Timothy Arceri | 9d5b106 | 2018-10-18 08:55:46 +1100 | [diff] [blame] | 233 | NIR_PASS(progress, shader, nir_opt_copy_prop_vars); |
| 234 | NIR_PASS(progress, shader, nir_opt_dead_write_vars); |
Connor Abbott | a69ab1b | 2019-06-26 14:03:31 +0200 | [diff] [blame] | 235 | NIR_PASS(progress, shader, nir_remove_dead_variables, |
| 236 | nir_var_function_temp); |
Timothy Arceri | 9d5b106 | 2018-10-18 08:55:46 +1100 | [diff] [blame] | 237 | |
Vasily Khoruzhick | 9367d2c | 2019-08-29 21:14:54 -0700 | [diff] [blame] | 238 | NIR_PASS_V(shader, nir_lower_alu_to_scalar, NULL, NULL); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 239 | NIR_PASS_V(shader, nir_lower_phis_to_scalar); |
| 240 | |
| 241 | NIR_PASS(progress, shader, nir_copy_prop); |
| 242 | NIR_PASS(progress, shader, nir_opt_remove_phis); |
| 243 | NIR_PASS(progress, shader, nir_opt_dce); |
| 244 | if (nir_opt_trivial_continues(shader)) { |
| 245 | progress = true; |
| 246 | NIR_PASS(progress, shader, nir_copy_prop); |
Dave Airlie | 64d9bd1 | 2017-09-13 03:49:31 +0100 | [diff] [blame] | 247 | NIR_PASS(progress, shader, nir_opt_remove_phis); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 248 | NIR_PASS(progress, shader, nir_opt_dce); |
| 249 | } |
Timothy Arceri | e30804c | 2019-04-08 20:13:49 +1000 | [diff] [blame] | 250 | NIR_PASS(progress, shader, nir_opt_if, true); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 251 | NIR_PASS(progress, shader, nir_opt_dead_cf); |
| 252 | NIR_PASS(progress, shader, nir_opt_cse); |
Ian Romanick | 378f996 | 2018-06-18 16:11:55 -0700 | [diff] [blame] | 253 | NIR_PASS(progress, shader, nir_opt_peephole_select, 8, true, true); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 254 | NIR_PASS(progress, shader, nir_opt_constant_folding); |
Timothy Arceri | e19a8fe | 2019-05-02 13:38:52 +1000 | [diff] [blame] | 255 | NIR_PASS(progress, shader, nir_opt_algebraic); |
Ian Romanick | d41cdef | 2018-08-18 16:42:04 -0700 | [diff] [blame] | 256 | |
| 257 | if (lower_flrp != 0) { |
Ian Romanick | 1f1007a | 2019-05-08 07:32:43 -0700 | [diff] [blame] | 258 | bool lower_flrp_progress = false; |
Ian Romanick | d41cdef | 2018-08-18 16:42:04 -0700 | [diff] [blame] | 259 | NIR_PASS(lower_flrp_progress, |
| 260 | shader, |
| 261 | nir_lower_flrp, |
| 262 | lower_flrp, |
| 263 | false /* always_precise */, |
| 264 | shader->options->lower_ffma); |
| 265 | if (lower_flrp_progress) { |
| 266 | NIR_PASS(progress, shader, |
| 267 | nir_opt_constant_folding); |
| 268 | progress = true; |
| 269 | } |
| 270 | |
| 271 | /* Nothing should rematerialize any flrps, so we only |
| 272 | * need to do this lowering once. |
| 273 | */ |
| 274 | lower_flrp = 0; |
| 275 | } |
| 276 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 277 | NIR_PASS(progress, shader, nir_opt_undef); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 278 | if (shader->options->max_unroll_iterations) { |
| 279 | NIR_PASS(progress, shader, nir_opt_loop_unroll, 0); |
| 280 | } |
Timothy Arceri | ce18881 | 2018-05-08 14:57:55 +1000 | [diff] [blame] | 281 | } while (progress && !optimize_conservatively); |
Samuel Pitoiset | 3488a3f | 2018-01-29 17:19:18 +0100 | [diff] [blame] | 282 | |
Daniel Schürmann | 64b7386 | 2019-07-20 19:21:14 +0200 | [diff] [blame] | 283 | NIR_PASS(progress, shader, nir_opt_conditional_discard); |
Samuel Pitoiset | 3488a3f | 2018-01-29 17:19:18 +0100 | [diff] [blame] | 284 | NIR_PASS(progress, shader, nir_opt_shrink_load); |
Rhys Perry | 7740149 | 2019-07-24 19:23:21 +0100 | [diff] [blame] | 285 | NIR_PASS(progress, shader, nir_opt_move, nir_move_load_ubo); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 286 | } |
| 287 | |
| 288 | nir_shader * |
| 289 | radv_shader_compile_to_nir(struct radv_device *device, |
| 290 | struct radv_shader_module *module, |
| 291 | const char *entrypoint_name, |
| 292 | gl_shader_stage stage, |
Timothy Arceri | ce18881 | 2018-05-08 14:57:55 +1000 | [diff] [blame] | 293 | const VkSpecializationInfo *spec_info, |
Bas Nieuwenhuizen | 5c3467e | 2019-03-30 14:28:06 +0100 | [diff] [blame] | 294 | const VkPipelineCreateFlags flags, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 295 | const struct radv_pipeline_layout *layout, |
| 296 | bool use_aco) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 297 | { |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 298 | nir_shader *nir; |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 299 | const nir_shader_compiler_options *nir_options = use_aco ? &nir_options_aco : |
| 300 | &nir_options_llvm; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 301 | if (module->nir) { |
| 302 | /* Some things such as our meta clear/blit code will give us a NIR |
| 303 | * shader directly. In that case, we just ignore the SPIR-V entirely |
| 304 | * and just use the NIR shader */ |
| 305 | nir = module->nir; |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 306 | nir->options = nir_options; |
Jason Ekstrand | 28bb6ab | 2018-10-18 15:18:30 -0500 | [diff] [blame] | 307 | nir_validate_shader(nir, "in internal shader"); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 308 | |
| 309 | assert(exec_list_length(&nir->functions) == 1); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 310 | } else { |
| 311 | uint32_t *spirv = (uint32_t *) module->data; |
| 312 | assert(module->size % 4 == 0); |
| 313 | |
Timothy Arceri | 7664aaf | 2017-10-11 11:59:20 +1100 | [diff] [blame] | 314 | if (device->instance->debug_flags & RADV_DEBUG_DUMP_SPIRV) |
Samuel Pitoiset | 844ae72 | 2017-09-22 16:56:40 +0200 | [diff] [blame] | 315 | radv_print_spirv(spirv, module->size, stderr); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 316 | |
| 317 | uint32_t num_spec_entries = 0; |
| 318 | struct nir_spirv_specialization *spec_entries = NULL; |
| 319 | if (spec_info && spec_info->mapEntryCount > 0) { |
| 320 | num_spec_entries = spec_info->mapEntryCount; |
| 321 | spec_entries = malloc(num_spec_entries * sizeof(*spec_entries)); |
| 322 | for (uint32_t i = 0; i < num_spec_entries; i++) { |
| 323 | VkSpecializationMapEntry entry = spec_info->pMapEntries[i]; |
| 324 | const void *data = spec_info->pData + entry.offset; |
| 325 | assert(data + entry.size <= spec_info->pData + spec_info->dataSize); |
| 326 | |
| 327 | spec_entries[i].id = spec_info->pMapEntries[i].constantID; |
| 328 | if (spec_info->dataSize == 8) |
| 329 | spec_entries[i].data64 = *(const uint64_t *)data; |
| 330 | else |
| 331 | spec_entries[i].data32 = *(const uint32_t *)data; |
| 332 | } |
| 333 | } |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 334 | const struct spirv_to_nir_options spirv_options = { |
Jason Ekstrand | 63b9aa2 | 2018-12-14 18:36:01 -0600 | [diff] [blame] | 335 | .lower_ubo_ssbo_access_to_offsets = true, |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 336 | .caps = { |
Daniel Schürmann | 7a858f2 | 2018-05-09 20:41:23 +0200 | [diff] [blame] | 337 | .amd_gcn_shader = true, |
Samuel Pitoiset | e73d863 | 2019-08-21 08:38:24 +0200 | [diff] [blame] | 338 | .amd_shader_ballot = device->physical_device->use_shader_ballot, |
Daniel Schürmann | 7a858f2 | 2018-05-09 20:41:23 +0200 | [diff] [blame] | 339 | .amd_trinary_minmax = true, |
Daniel Schürmann | 2812622 | 2019-09-17 17:09:52 +0200 | [diff] [blame] | 340 | .demote_to_helper_invocation = device->physical_device->use_aco, |
Samuel Pitoiset | b3e3440 | 2019-04-19 12:40:37 +0200 | [diff] [blame] | 341 | .derivative_group = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 342 | .descriptor_array_dynamic_indexing = true, |
Juan A. Suarez Romero | 06c9d7f | 2019-04-29 17:05:13 +0200 | [diff] [blame] | 343 | .descriptor_array_non_uniform_indexing = true, |
| 344 | .descriptor_indexing = true, |
Bas Nieuwenhuizen | 5240fdd | 2018-01-21 17:13:26 +0100 | [diff] [blame] | 345 | .device_group = true, |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 346 | .draw_parameters = true, |
Samuel Pitoiset | 7c50214 | 2019-10-14 11:27:32 +0200 | [diff] [blame] | 347 | .float_controls = true, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 348 | .float16 = !device->physical_device->use_aco, |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 349 | .float64 = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 350 | .geometry_streams = true, |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 351 | .image_read_without_format = true, |
| 352 | .image_write_without_format = true, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 353 | .int8 = !device->physical_device->use_aco, |
| 354 | .int16 = !device->physical_device->use_aco, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 355 | .int64 = true, |
Samuel Pitoiset | 9cf55b0 | 2019-04-16 10:38:24 +0200 | [diff] [blame] | 356 | .int64_atomics = true, |
Jason Ekstrand | e19c623 | 2017-10-18 17:28:19 -0700 | [diff] [blame] | 357 | .multiview = true, |
Bas Nieuwenhuizen | 13ab63b | 2019-01-24 02:06:27 +0100 | [diff] [blame] | 358 | .physical_storage_buffer_address = true, |
Samuel Pitoiset | 07ff367 | 2019-07-16 17:11:50 +0200 | [diff] [blame] | 359 | .post_depth_coverage = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 360 | .runtime_descriptor_array = true, |
Samuel Pitoiset | cbd6f0a | 2019-10-07 10:26:22 +0200 | [diff] [blame] | 361 | .shader_clock = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 362 | .shader_viewport_index_layer = true, |
| 363 | .stencil_export = true, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 364 | .storage_8bit = !device->physical_device->use_aco, |
| 365 | .storage_16bit = !device->physical_device->use_aco, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 366 | .storage_image_ms = true, |
Samuel Pitoiset | 3565682 | 2018-09-18 15:27:52 +0200 | [diff] [blame] | 367 | .subgroup_arithmetic = true, |
Daniel Schürmann | f2c6a55 | 2018-03-06 15:05:13 +0100 | [diff] [blame] | 368 | .subgroup_ballot = true, |
Bas Nieuwenhuizen | 8f9af58 | 2018-01-21 15:06:10 +0100 | [diff] [blame] | 369 | .subgroup_basic = true, |
Daniel Schürmann | f2c6a55 | 2018-03-06 15:05:13 +0100 | [diff] [blame] | 370 | .subgroup_quad = true, |
| 371 | .subgroup_shuffle = true, |
| 372 | .subgroup_vote = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 373 | .tessellation = true, |
Samuel Pitoiset | b4eb029 | 2018-10-05 18:04:56 +0200 | [diff] [blame] | 374 | .transform_feedback = true, |
Jason Ekstrand | 05d72d6 | 2019-01-07 10:28:23 -0600 | [diff] [blame] | 375 | .variable_pointers = true, |
Daniel Schürmann | ffbf75c | 2018-02-23 13:55:01 +0100 | [diff] [blame] | 376 | }, |
Caio Marcelo de Oliveira Filho | 31a7476 | 2019-05-01 14:15:32 -0700 | [diff] [blame] | 377 | .ubo_addr_format = nir_address_format_32bit_index_offset, |
| 378 | .ssbo_addr_format = nir_address_format_32bit_index_offset, |
| 379 | .phys_ssbo_addr_format = nir_address_format_64bit_global, |
| 380 | .push_const_addr_format = nir_address_format_logical, |
| 381 | .shared_addr_format = nir_address_format_32bit_offset, |
Connor Abbott | 27f0c3c | 2019-05-13 15:39:54 +0200 | [diff] [blame] | 382 | .frag_coord_is_sysval = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 383 | }; |
Caio Marcelo de Oliveira Filho | e45bf01 | 2019-05-19 00:22:17 -0700 | [diff] [blame] | 384 | nir = spirv_to_nir(spirv, module->size / 4, |
| 385 | spec_entries, num_spec_entries, |
| 386 | stage, entrypoint_name, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 387 | &spirv_options, nir_options); |
Jason Ekstrand | 59fb59a | 2017-09-14 19:52:38 -0700 | [diff] [blame] | 388 | assert(nir->info.stage == stage); |
Jason Ekstrand | 28bb6ab | 2018-10-18 15:18:30 -0500 | [diff] [blame] | 389 | nir_validate_shader(nir, "after spirv_to_nir"); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 390 | |
| 391 | free(spec_entries); |
| 392 | |
| 393 | /* We have to lower away local constant initializers right before we |
| 394 | * inline functions. That way they get properly initialized at the top |
| 395 | * of the function and not at the top of its caller. |
| 396 | */ |
Karol Herbst | 9b24028 | 2019-01-16 00:05:04 +0100 | [diff] [blame] | 397 | NIR_PASS_V(nir, nir_lower_constant_initializers, nir_var_function_temp); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 398 | NIR_PASS_V(nir, nir_lower_returns); |
| 399 | NIR_PASS_V(nir, nir_inline_functions); |
Jason Ekstrand | fc9c4f8 | 2018-12-13 11:08:13 -0600 | [diff] [blame] | 400 | NIR_PASS_V(nir, nir_opt_deref); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 401 | |
| 402 | /* Pick off the single entrypoint that we want */ |
| 403 | foreach_list_typed_safe(nir_function, func, node, &nir->functions) { |
Caio Marcelo de Oliveira Filho | a3bfdac | 2019-05-19 00:11:37 -0700 | [diff] [blame] | 404 | if (func->is_entrypoint) |
| 405 | func->name = ralloc_strdup(func, "main"); |
| 406 | else |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 407 | exec_node_remove(&func->node); |
| 408 | } |
| 409 | assert(exec_list_length(&nir->functions) == 1); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 410 | |
Dave Airlie | e8d9b7a | 2018-03-19 04:27:49 +0000 | [diff] [blame] | 411 | /* Make sure we lower constant initializers on output variables so that |
| 412 | * nir_remove_dead_variables below sees the corresponding stores |
| 413 | */ |
| 414 | NIR_PASS_V(nir, nir_lower_constant_initializers, nir_var_shader_out); |
| 415 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 416 | /* Now that we've deleted all but the main function, we can go ahead and |
| 417 | * lower the rest of the constant initializers. |
| 418 | */ |
| 419 | NIR_PASS_V(nir, nir_lower_constant_initializers, ~0); |
Jason Ekstrand | b0c643d | 2018-03-21 17:30:22 -0700 | [diff] [blame] | 420 | |
| 421 | /* Split member structs. We do this before lower_io_to_temporaries so that |
| 422 | * it doesn't lower system values to temporaries by accident. |
| 423 | */ |
| 424 | NIR_PASS_V(nir, nir_split_var_copies); |
| 425 | NIR_PASS_V(nir, nir_split_per_member_structs); |
| 426 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 427 | if (nir->info.stage == MESA_SHADER_FRAGMENT && use_aco) |
| 428 | NIR_PASS_V(nir, nir_lower_io_to_vector, nir_var_shader_out); |
Daniel Schürmann | e41e932 | 2019-04-05 11:01:39 +0200 | [diff] [blame] | 429 | if (nir->info.stage == MESA_SHADER_FRAGMENT) |
Connor Abbott | 27f0c3c | 2019-05-13 15:39:54 +0200 | [diff] [blame] | 430 | NIR_PASS_V(nir, nir_lower_input_attachments, true); |
Daniel Schürmann | e41e932 | 2019-04-05 11:01:39 +0200 | [diff] [blame] | 431 | |
Samuel Pitoiset | 24ee532 | 2018-08-22 12:34:13 +0200 | [diff] [blame] | 432 | NIR_PASS_V(nir, nir_remove_dead_variables, |
Daniel Schürmann | 8b78cce | 2019-09-17 18:24:06 +0200 | [diff] [blame] | 433 | nir_var_shader_in | nir_var_shader_out | nir_var_system_value | nir_var_mem_shared); |
Samuel Pitoiset | 24ee532 | 2018-08-22 12:34:13 +0200 | [diff] [blame] | 434 | |
Connor Abbott | 3f5b541 | 2019-09-05 13:57:11 +0200 | [diff] [blame] | 435 | NIR_PASS_V(nir, nir_propagate_invariant); |
| 436 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 437 | NIR_PASS_V(nir, nir_lower_system_values); |
| 438 | NIR_PASS_V(nir, nir_lower_clip_cull_distance_arrays); |
Bas Nieuwenhuizen | 5c3467e | 2019-03-30 14:28:06 +0100 | [diff] [blame] | 439 | NIR_PASS_V(nir, radv_nir_lower_ycbcr_textures, layout); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 440 | } |
| 441 | |
| 442 | /* Vulkan uses the separate-shader linking model */ |
| 443 | nir->info.separate_shader = true; |
| 444 | |
Caio Marcelo de Oliveira Filho | a3bfdac | 2019-05-19 00:11:37 -0700 | [diff] [blame] | 445 | nir_shader_gather_info(nir, nir_shader_get_entrypoint(nir)); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 446 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 447 | static const nir_lower_tex_options tex_options = { |
| 448 | .lower_txp = ~0, |
Jason Ekstrand | 08f804e | 2019-03-19 13:55:21 -0500 | [diff] [blame] | 449 | .lower_tg4_offsets = true, |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 450 | }; |
| 451 | |
| 452 | nir_lower_tex(nir, &tex_options); |
| 453 | |
| 454 | nir_lower_vars_to_ssa(nir); |
Samuel Pitoiset | ded1509 | 2018-05-23 14:31:55 +0200 | [diff] [blame] | 455 | |
Samuel Pitoiset | 38a8c59 | 2018-05-23 14:31:56 +0200 | [diff] [blame] | 456 | if (nir->info.stage == MESA_SHADER_VERTEX || |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 457 | nir->info.stage == MESA_SHADER_GEOMETRY || |
| 458 | nir->info.stage == MESA_SHADER_FRAGMENT) { |
Samuel Pitoiset | 38a8c59 | 2018-05-23 14:31:56 +0200 | [diff] [blame] | 459 | NIR_PASS_V(nir, nir_lower_io_to_temporaries, |
| 460 | nir_shader_get_entrypoint(nir), true, true); |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 461 | } else if (nir->info.stage == MESA_SHADER_TESS_EVAL) { |
Samuel Pitoiset | 38a8c59 | 2018-05-23 14:31:56 +0200 | [diff] [blame] | 462 | NIR_PASS_V(nir, nir_lower_io_to_temporaries, |
| 463 | nir_shader_get_entrypoint(nir), true, false); |
| 464 | } |
| 465 | |
Samuel Pitoiset | ded1509 | 2018-05-23 14:31:55 +0200 | [diff] [blame] | 466 | nir_split_var_copies(nir); |
Samuel Pitoiset | ded1509 | 2018-05-23 14:31:55 +0200 | [diff] [blame] | 467 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 468 | nir_lower_global_vars_to_local(nir); |
Karol Herbst | 9b24028 | 2019-01-16 00:05:04 +0100 | [diff] [blame] | 469 | nir_remove_dead_variables(nir, nir_var_function_temp); |
Bas Nieuwenhuizen | 8f9af58 | 2018-01-21 15:06:10 +0100 | [diff] [blame] | 470 | nir_lower_subgroups(nir, &(struct nir_lower_subgroups_options) { |
| 471 | .subgroup_size = 64, |
| 472 | .ballot_bit_size = 64, |
| 473 | .lower_to_scalar = 1, |
| 474 | .lower_subgroup_masks = 1, |
| 475 | .lower_shuffle = 1, |
Daniel Schürmann | f2c6a55 | 2018-03-06 15:05:13 +0100 | [diff] [blame] | 476 | .lower_shuffle_to_32bit = 1, |
| 477 | .lower_vote_eq_to_ballot = 1, |
Bas Nieuwenhuizen | 8f9af58 | 2018-01-21 15:06:10 +0100 | [diff] [blame] | 478 | }); |
| 479 | |
Timothy Arceri | 72e4287 | 2018-09-24 18:18:48 +1000 | [diff] [blame] | 480 | nir_lower_load_const_to_scalar(nir); |
| 481 | |
Timothy Arceri | ce18881 | 2018-05-08 14:57:55 +1000 | [diff] [blame] | 482 | if (!(flags & VK_PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT)) |
Timothy Arceri | 0667571 | 2018-10-18 09:42:17 +1100 | [diff] [blame] | 483 | radv_optimize_nir(nir, false, true); |
| 484 | |
| 485 | /* We call nir_lower_var_copies() after the first radv_optimize_nir() |
| 486 | * to remove any copies introduced by nir_opt_find_array_copies(). |
| 487 | */ |
| 488 | nir_lower_var_copies(nir); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 489 | |
Connor Abbott | 71a6794 | 2019-08-30 16:08:47 +0200 | [diff] [blame] | 490 | /* Lower large variables that are always constant with load_constant |
| 491 | * intrinsics, which get turned into PC-relative loads from a data |
| 492 | * section next to the shader. |
| 493 | */ |
| 494 | NIR_PASS_V(nir, nir_opt_large_constants, |
| 495 | glsl_get_natural_size_align_bytes, 16); |
| 496 | |
Timothy Arceri | 9a243ec | 2018-03-08 16:20:48 +1100 | [diff] [blame] | 497 | /* Indirect lowering must be called after the radv_optimize_nir() loop |
| 498 | * has been called at least once. Otherwise indirect lowering can |
| 499 | * bloat the instruction count of the loop and cause it to be |
| 500 | * considered too large for unrolling. |
| 501 | */ |
| 502 | ac_lower_indirect_derefs(nir, device->physical_device->rad_info.chip_class); |
Timothy Arceri | 0667571 | 2018-10-18 09:42:17 +1100 | [diff] [blame] | 503 | radv_optimize_nir(nir, flags & VK_PIPELINE_CREATE_DISABLE_OPTIMIZATION_BIT, false); |
Timothy Arceri | 9a243ec | 2018-03-08 16:20:48 +1100 | [diff] [blame] | 504 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 505 | return nir; |
| 506 | } |
| 507 | |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 508 | static int |
| 509 | type_size_vec4(const struct glsl_type *type, bool bindless) |
| 510 | { |
| 511 | return glsl_count_attribute_slots(type, false); |
| 512 | } |
| 513 | |
| 514 | static nir_variable * |
| 515 | find_layer_in_var(nir_shader *nir) |
| 516 | { |
| 517 | nir_foreach_variable(var, &nir->inputs) { |
| 518 | if (var->data.location == VARYING_SLOT_LAYER) { |
| 519 | return var; |
| 520 | } |
| 521 | } |
| 522 | |
| 523 | nir_variable *var = |
| 524 | nir_variable_create(nir, nir_var_shader_in, glsl_int_type(), "layer id"); |
| 525 | var->data.location = VARYING_SLOT_LAYER; |
| 526 | var->data.interpolation = INTERP_MODE_FLAT; |
| 527 | return var; |
| 528 | } |
| 529 | |
| 530 | /* We use layered rendering to implement multiview, which means we need to map |
| 531 | * view_index to gl_Layer. The attachment lowering also uses needs to know the |
| 532 | * layer so that it can sample from the correct layer. The code generates a |
| 533 | * load from the layer_id sysval, but since we don't have a way to get at this |
| 534 | * information from the fragment shader, we also need to lower this to the |
| 535 | * gl_Layer varying. This pass lowers both to a varying load from the LAYER |
| 536 | * slot, before lowering io, so that nir_assign_var_locations() will give the |
| 537 | * LAYER varying the correct driver_location. |
| 538 | */ |
| 539 | |
| 540 | static bool |
| 541 | lower_view_index(nir_shader *nir) |
| 542 | { |
| 543 | bool progress = false; |
| 544 | nir_function_impl *entry = nir_shader_get_entrypoint(nir); |
| 545 | nir_builder b; |
| 546 | nir_builder_init(&b, entry); |
| 547 | |
| 548 | nir_variable *layer = NULL; |
| 549 | nir_foreach_block(block, entry) { |
| 550 | nir_foreach_instr_safe(instr, block) { |
| 551 | if (instr->type != nir_instr_type_intrinsic) |
| 552 | continue; |
| 553 | |
| 554 | nir_intrinsic_instr *load = nir_instr_as_intrinsic(instr); |
| 555 | if (load->intrinsic != nir_intrinsic_load_view_index && |
| 556 | load->intrinsic != nir_intrinsic_load_layer_id) |
| 557 | continue; |
| 558 | |
| 559 | if (!layer) |
| 560 | layer = find_layer_in_var(nir); |
| 561 | |
| 562 | b.cursor = nir_before_instr(instr); |
| 563 | nir_ssa_def *def = nir_load_var(&b, layer); |
| 564 | nir_ssa_def_rewrite_uses(&load->dest.ssa, |
| 565 | nir_src_for_ssa(def)); |
| 566 | |
| 567 | nir_instr_remove(instr); |
| 568 | progress = true; |
| 569 | } |
| 570 | } |
| 571 | |
| 572 | return progress; |
| 573 | } |
| 574 | |
Samuel Pitoiset | 8d44f83 | 2019-08-29 11:16:44 +0200 | [diff] [blame] | 575 | void |
| 576 | radv_lower_fs_io(nir_shader *nir) |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 577 | { |
| 578 | NIR_PASS_V(nir, lower_view_index); |
| 579 | nir_assign_io_var_locations(&nir->inputs, &nir->num_inputs, |
| 580 | MESA_SHADER_FRAGMENT); |
| 581 | |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 582 | NIR_PASS_V(nir, nir_lower_io, nir_var_shader_in, type_size_vec4, 0); |
| 583 | |
| 584 | /* This pass needs actual constants */ |
| 585 | nir_opt_constant_folding(nir); |
| 586 | |
| 587 | NIR_PASS_V(nir, nir_io_add_const_offset_to_base, nir_var_shader_in); |
Connor Abbott | 118a66d | 2019-05-10 10:44:20 +0200 | [diff] [blame] | 588 | } |
| 589 | |
| 590 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 591 | void * |
| 592 | radv_alloc_shader_memory(struct radv_device *device, |
| 593 | struct radv_shader_variant *shader) |
| 594 | { |
| 595 | mtx_lock(&device->shader_slab_mutex); |
| 596 | list_for_each_entry(struct radv_shader_slab, slab, &device->shader_slabs, slabs) { |
| 597 | uint64_t offset = 0; |
| 598 | list_for_each_entry(struct radv_shader_variant, s, &slab->shaders, slab_list) { |
| 599 | if (s->bo_offset - offset >= shader->code_size) { |
| 600 | shader->bo = slab->bo; |
| 601 | shader->bo_offset = offset; |
| 602 | list_addtail(&shader->slab_list, &s->slab_list); |
| 603 | mtx_unlock(&device->shader_slab_mutex); |
| 604 | return slab->ptr + offset; |
| 605 | } |
| 606 | offset = align_u64(s->bo_offset + s->code_size, 256); |
| 607 | } |
| 608 | if (slab->size - offset >= shader->code_size) { |
| 609 | shader->bo = slab->bo; |
| 610 | shader->bo_offset = offset; |
| 611 | list_addtail(&shader->slab_list, &slab->shaders); |
| 612 | mtx_unlock(&device->shader_slab_mutex); |
| 613 | return slab->ptr + offset; |
| 614 | } |
| 615 | } |
| 616 | |
| 617 | mtx_unlock(&device->shader_slab_mutex); |
| 618 | struct radv_shader_slab *slab = calloc(1, sizeof(struct radv_shader_slab)); |
| 619 | |
| 620 | slab->size = 256 * 1024; |
| 621 | slab->bo = device->ws->buffer_create(device->ws, slab->size, 256, |
Samuel Pitoiset | a3c2a86 | 2018-01-04 15:19:47 +0100 | [diff] [blame] | 622 | RADEON_DOMAIN_VRAM, |
| 623 | RADEON_FLAG_NO_INTERPROCESS_SHARING | |
Samuel Pitoiset | 2b9c371 | 2019-08-20 17:20:42 +0200 | [diff] [blame] | 624 | (device->physical_device->rad_info.cpdma_prefetch_writes_memory ? |
Bas Nieuwenhuizen | ead54d4 | 2019-01-28 00:28:05 +0100 | [diff] [blame] | 625 | 0 : RADEON_FLAG_READ_ONLY), |
| 626 | RADV_BO_PRIORITY_SHADER); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 627 | slab->ptr = (char*)device->ws->buffer_map(slab->bo); |
| 628 | list_inithead(&slab->shaders); |
| 629 | |
| 630 | mtx_lock(&device->shader_slab_mutex); |
| 631 | list_add(&slab->slabs, &device->shader_slabs); |
| 632 | |
| 633 | shader->bo = slab->bo; |
| 634 | shader->bo_offset = 0; |
| 635 | list_add(&shader->slab_list, &slab->shaders); |
| 636 | mtx_unlock(&device->shader_slab_mutex); |
| 637 | return slab->ptr; |
| 638 | } |
| 639 | |
| 640 | void |
| 641 | radv_destroy_shader_slabs(struct radv_device *device) |
| 642 | { |
| 643 | list_for_each_entry_safe(struct radv_shader_slab, slab, &device->shader_slabs, slabs) { |
| 644 | device->ws->buffer_destroy(slab->bo); |
| 645 | free(slab); |
| 646 | } |
| 647 | mtx_destroy(&device->shader_slab_mutex); |
| 648 | } |
| 649 | |
Samuel Pitoiset | 939e5a3 | 2018-06-27 10:39:51 +0200 | [diff] [blame] | 650 | /* For the UMR disassembler. */ |
| 651 | #define DEBUGGER_END_OF_CODE_MARKER 0xbf9f0000 /* invalid instruction */ |
| 652 | #define DEBUGGER_NUM_MARKERS 5 |
| 653 | |
| 654 | static unsigned |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 655 | radv_get_shader_binary_size(size_t code_size) |
Samuel Pitoiset | 939e5a3 | 2018-06-27 10:39:51 +0200 | [diff] [blame] | 656 | { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 657 | return code_size + DEBUGGER_NUM_MARKERS * 4; |
Samuel Pitoiset | 939e5a3 | 2018-06-27 10:39:51 +0200 | [diff] [blame] | 658 | } |
| 659 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 660 | static void radv_postprocess_config(const struct radv_physical_device *pdevice, |
| 661 | const struct ac_shader_config *config_in, |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 662 | const struct radv_shader_info *info, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 663 | gl_shader_stage stage, |
| 664 | struct ac_shader_config *config_out) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 665 | { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 666 | bool scratch_enabled = config_in->scratch_bytes_per_wave > 0; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 667 | unsigned vgpr_comp_cnt = 0; |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 668 | unsigned num_input_vgprs = info->num_input_vgprs; |
| 669 | |
| 670 | if (stage == MESA_SHADER_FRAGMENT) { |
Timur Kristóf | a4fd8ba | 2019-09-25 16:40:07 +0200 | [diff] [blame] | 671 | num_input_vgprs = ac_get_fs_input_vgpr_cnt(config_in, NULL, NULL); |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 672 | } |
| 673 | |
| 674 | unsigned num_vgprs = MAX2(config_in->num_vgprs, num_input_vgprs); |
| 675 | /* +3 for scratch wave offset and VCC */ |
| 676 | unsigned num_sgprs = MAX2(config_in->num_sgprs, info->num_input_sgprs + 3); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 677 | unsigned num_shared_vgprs = config_in->num_shared_vgprs; |
| 678 | /* shared VGPRs are introduced in Navi and are allocated in blocks of 8 (RDNA ref 3.6.5) */ |
| 679 | assert((pdevice->rad_info.chip_class >= GFX10 && num_shared_vgprs % 8 == 0) |
| 680 | || (pdevice->rad_info.chip_class < GFX10 && num_shared_vgprs == 0)); |
| 681 | unsigned num_shared_vgpr_blocks = num_shared_vgprs / 8; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 682 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 683 | *config_out = *config_in; |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 684 | config_out->num_vgprs = num_vgprs; |
| 685 | config_out->num_sgprs = num_sgprs; |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 686 | config_out->num_shared_vgprs = num_shared_vgprs; |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 687 | |
| 688 | /* Enable 64-bit and 16-bit denormals, because there is no performance |
| 689 | * cost. |
| 690 | * |
| 691 | * If denormals are enabled, all floating-point output modifiers are |
| 692 | * ignored. |
| 693 | * |
| 694 | * Don't enable denormals for 32-bit floats, because: |
| 695 | * - Floating-point output modifiers would be ignored by the hw. |
| 696 | * - Some opcodes don't support denormals, such as v_mad_f32. We would |
| 697 | * have to stop using those. |
| 698 | * - GFX6 & GFX7 would be very slow. |
| 699 | */ |
| 700 | config_out->float_mode |= V_00B028_FP_64_DENORMS; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 701 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 702 | config_out->rsrc2 = S_00B12C_USER_SGPR(info->num_user_sgprs) | |
Samuel Pitoiset | a15b3bc | 2019-09-09 10:23:30 +0200 | [diff] [blame] | 703 | S_00B12C_SCRATCH_EN(scratch_enabled); |
| 704 | |
| 705 | if (!pdevice->use_ngg_streamout) { |
| 706 | config_out->rsrc2 |= S_00B12C_SO_BASE0_EN(!!info->so.strides[0]) | |
| 707 | S_00B12C_SO_BASE1_EN(!!info->so.strides[1]) | |
| 708 | S_00B12C_SO_BASE2_EN(!!info->so.strides[2]) | |
| 709 | S_00B12C_SO_BASE3_EN(!!info->so.strides[3]) | |
| 710 | S_00B12C_SO_EN(!!info->so.num_outputs); |
| 711 | } |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 712 | |
Samuel Pitoiset | ea38565 | 2019-07-30 18:32:42 +0200 | [diff] [blame] | 713 | config_out->rsrc1 = S_00B848_VGPRS((num_vgprs - 1) / |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 714 | (info->wave_size == 32 ? 8 : 4)) | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 715 | S_00B848_DX10_CLAMP(1) | |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 716 | S_00B848_FLOAT_MODE(config_out->float_mode); |
Bas Nieuwenhuizen | 228325f | 2017-10-18 00:59:16 +0200 | [diff] [blame] | 717 | |
Samuel Pitoiset | 4c82094 | 2019-06-25 13:33:03 +0200 | [diff] [blame] | 718 | if (pdevice->rad_info.chip_class >= GFX10) { |
| 719 | config_out->rsrc2 |= S_00B22C_USER_SGPR_MSB_GFX10(info->num_user_sgprs >> 5); |
| 720 | } else { |
| 721 | config_out->rsrc1 |= S_00B228_SGPRS((num_sgprs - 1) / 8); |
Samuel Pitoiset | 09abe57 | 2019-07-23 14:55:16 +0200 | [diff] [blame] | 722 | config_out->rsrc2 |= S_00B22C_USER_SGPR_MSB_GFX9(info->num_user_sgprs >> 5); |
Samuel Pitoiset | 4c82094 | 2019-06-25 13:33:03 +0200 | [diff] [blame] | 723 | } |
| 724 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 725 | switch (stage) { |
| 726 | case MESA_SHADER_TESS_EVAL: |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 727 | if (info->is_ngg) { |
| 728 | config_out->rsrc1 |= S_00B228_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10); |
| 729 | config_out->rsrc2 |= S_00B22C_OC_LDS_EN(1); |
| 730 | } else if (info->tes.as_es) { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 731 | assert(pdevice->rad_info.chip_class <= GFX8); |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 732 | vgpr_comp_cnt = info->uses_prim_id ? 3 : 2; |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 733 | |
| 734 | config_out->rsrc2 |= S_00B12C_OC_LDS_EN(1); |
Samuel Pitoiset | b4477fa | 2019-06-26 15:11:00 +0200 | [diff] [blame] | 735 | } else { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 736 | bool enable_prim_id = info->tes.export_prim_id || info->uses_prim_id; |
Samuel Pitoiset | b4477fa | 2019-06-26 15:11:00 +0200 | [diff] [blame] | 737 | vgpr_comp_cnt = enable_prim_id ? 3 : 2; |
Bas Nieuwenhuizen | aeb5b1a | 2019-07-06 12:31:25 +0200 | [diff] [blame] | 738 | |
| 739 | config_out->rsrc1 |= S_00B128_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10); |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 740 | config_out->rsrc2 |= S_00B12C_OC_LDS_EN(1); |
Samuel Pitoiset | b4477fa | 2019-06-26 15:11:00 +0200 | [diff] [blame] | 741 | } |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 742 | config_out->rsrc2 |= S_00B22C_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
Bas Nieuwenhuizen | 228325f | 2017-10-18 00:59:16 +0200 | [diff] [blame] | 743 | break; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 744 | case MESA_SHADER_TESS_CTRL: |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 745 | if (pdevice->rad_info.chip_class >= GFX9) { |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 746 | /* We need at least 2 components for LS. |
| 747 | * VGPR0-3: (VertexID, RelAutoindex, InstanceID / StepRate0, InstanceID). |
| 748 | * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded. |
| 749 | */ |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 750 | if (pdevice->rad_info.chip_class >= GFX10) { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 751 | vgpr_comp_cnt = info->vs.needs_instance_id ? 3 : 1; |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 752 | } else { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 753 | vgpr_comp_cnt = info->vs.needs_instance_id ? 2 : 1; |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 754 | } |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 755 | } else { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 756 | config_out->rsrc2 |= S_00B12C_OC_LDS_EN(1); |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 757 | } |
Samuel Pitoiset | e68b55f | 2019-07-12 12:17:16 +0200 | [diff] [blame] | 758 | config_out->rsrc1 |= S_00B428_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10) | |
| 759 | S_00B848_WGP_MODE(pdevice->rad_info.chip_class >= GFX10); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 760 | config_out->rsrc2 |= S_00B42C_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 761 | break; |
| 762 | case MESA_SHADER_VERTEX: |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 763 | if (info->is_ngg) { |
| 764 | config_out->rsrc1 |= S_00B228_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10); |
| 765 | } else if (info->vs.as_ls) { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 766 | assert(pdevice->rad_info.chip_class <= GFX8); |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 767 | /* We need at least 2 components for LS. |
| 768 | * VGPR0-3: (VertexID, RelAutoindex, InstanceID / StepRate0, InstanceID). |
| 769 | * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded. |
| 770 | */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 771 | vgpr_comp_cnt = info->vs.needs_instance_id ? 2 : 1; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 772 | } else if (info->vs.as_es) { |
| 773 | assert(pdevice->rad_info.chip_class <= GFX8); |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 774 | /* VGPR0-3: (VertexID, InstanceID / StepRate0, ...) */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 775 | vgpr_comp_cnt = info->vs.needs_instance_id ? 1 : 0; |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 776 | } else { |
| 777 | /* VGPR0-3: (VertexID, InstanceID / StepRate0, PrimID, InstanceID) |
| 778 | * If PrimID is disabled. InstanceID / StepRate1 is loaded instead. |
| 779 | * StepRate0 is set to 1. so that VGPR3 doesn't have to be loaded. |
| 780 | */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 781 | if (info->vs.needs_instance_id && pdevice->rad_info.chip_class >= GFX10) { |
Bas Nieuwenhuizen | 2e763f7 | 2019-08-21 01:50:53 +0200 | [diff] [blame] | 782 | vgpr_comp_cnt = 3; |
| 783 | } else if (info->vs.export_prim_id) { |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 784 | vgpr_comp_cnt = 2; |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 785 | } else if (info->vs.needs_instance_id) { |
Bas Nieuwenhuizen | 2e763f7 | 2019-08-21 01:50:53 +0200 | [diff] [blame] | 786 | vgpr_comp_cnt = 1; |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 787 | } else { |
| 788 | vgpr_comp_cnt = 0; |
| 789 | } |
Bas Nieuwenhuizen | aeb5b1a | 2019-07-06 12:31:25 +0200 | [diff] [blame] | 790 | |
| 791 | config_out->rsrc1 |= S_00B128_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 792 | config_out->rsrc2 |= S_00B12C_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 793 | } |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 794 | break; |
| 795 | case MESA_SHADER_FRAGMENT: |
Bas Nieuwenhuizen | aeb5b1a | 2019-07-06 12:31:25 +0200 | [diff] [blame] | 796 | config_out->rsrc1 |= S_00B028_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 797 | config_out->rsrc2 |= S_00B02C_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
Bas Nieuwenhuizen | aeb5b1a | 2019-07-06 12:31:25 +0200 | [diff] [blame] | 798 | break; |
Samuel Pitoiset | f4d2c47 | 2019-06-26 15:11:01 +0200 | [diff] [blame] | 799 | case MESA_SHADER_GEOMETRY: |
Samuel Pitoiset | e68b55f | 2019-07-12 12:17:16 +0200 | [diff] [blame] | 800 | config_out->rsrc1 |= S_00B228_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10) | |
| 801 | S_00B848_WGP_MODE(pdevice->rad_info.chip_class >= GFX10); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 802 | config_out->rsrc2 |= S_00B22C_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 803 | break; |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 804 | case MESA_SHADER_COMPUTE: |
Samuel Pitoiset | e68b55f | 2019-07-12 12:17:16 +0200 | [diff] [blame] | 805 | config_out->rsrc1 |= S_00B848_MEM_ORDERED(pdevice->rad_info.chip_class >= GFX10) | |
| 806 | S_00B848_WGP_MODE(pdevice->rad_info.chip_class >= GFX10); |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 807 | config_out->rsrc2 |= |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 808 | S_00B84C_TGID_X_EN(info->cs.uses_block_id[0]) | |
| 809 | S_00B84C_TGID_Y_EN(info->cs.uses_block_id[1]) | |
| 810 | S_00B84C_TGID_Z_EN(info->cs.uses_block_id[2]) | |
| 811 | S_00B84C_TIDIG_COMP_CNT(info->cs.uses_thread_id[2] ? 2 : |
| 812 | info->cs.uses_thread_id[1] ? 1 : 0) | |
| 813 | S_00B84C_TG_SIZE_EN(info->cs.uses_local_invocation_idx) | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 814 | S_00B84C_LDS_SIZE(config_in->lds_size); |
Timur Kristóf | 83eebdb | 2019-09-13 15:53:09 +0200 | [diff] [blame] | 815 | config_out->rsrc3 |= S_00B8A0_SHARED_VGPR_CNT(num_shared_vgpr_blocks); |
| 816 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 817 | break; |
| 818 | default: |
| 819 | unreachable("unsupported shader type"); |
| 820 | break; |
| 821 | } |
| 822 | |
Samuel Pitoiset | edf1af6 | 2019-07-16 16:39:16 +0200 | [diff] [blame] | 823 | if (pdevice->rad_info.chip_class >= GFX10 && info->is_ngg && |
Samuel Pitoiset | 3f50007 | 2019-07-09 08:44:01 +0200 | [diff] [blame] | 824 | (stage == MESA_SHADER_VERTEX || stage == MESA_SHADER_TESS_EVAL || stage == MESA_SHADER_GEOMETRY)) { |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 825 | unsigned gs_vgpr_comp_cnt, es_vgpr_comp_cnt; |
Bas Nieuwenhuizen | 7286865 | 2019-07-11 08:44:15 +0200 | [diff] [blame] | 826 | gl_shader_stage es_stage = stage; |
| 827 | if (stage == MESA_SHADER_GEOMETRY) |
| 828 | es_stage = info->gs.es_type; |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 829 | |
| 830 | /* VGPR5-8: (VertexID, UserVGPR0, UserVGPR1, UserVGPR2 / InstanceID) */ |
Bas Nieuwenhuizen | 7286865 | 2019-07-11 08:44:15 +0200 | [diff] [blame] | 831 | if (es_stage == MESA_SHADER_VERTEX) { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 832 | es_vgpr_comp_cnt = info->vs.needs_instance_id ? 3 : 0; |
Bas Nieuwenhuizen | 7286865 | 2019-07-11 08:44:15 +0200 | [diff] [blame] | 833 | } else if (es_stage == MESA_SHADER_TESS_EVAL) { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 834 | bool enable_prim_id = info->tes.export_prim_id || info->uses_prim_id; |
Samuel Pitoiset | d2a8b63 | 2019-07-09 08:27:30 +0200 | [diff] [blame] | 835 | es_vgpr_comp_cnt = enable_prim_id ? 3 : 2; |
Bas Nieuwenhuizen | 451f030 | 2019-07-19 00:00:03 +0200 | [diff] [blame] | 836 | } else |
| 837 | unreachable("Unexpected ES shader stage"); |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 838 | |
| 839 | bool tes_triangles = stage == MESA_SHADER_TESS_EVAL && |
| 840 | info->tes.primitive_mode >= 4; /* GL_TRIANGLES */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 841 | if (info->uses_invocation_id || stage == MESA_SHADER_VERTEX) { |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 842 | gs_vgpr_comp_cnt = 3; /* VGPR3 contains InvocationID. */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 843 | } else if (info->uses_prim_id) { |
Bas Nieuwenhuizen | 795adbb | 2019-07-08 23:44:32 +0200 | [diff] [blame] | 844 | gs_vgpr_comp_cnt = 2; /* VGPR2 contains PrimitiveID. */ |
| 845 | } else if (info->gs.vertices_in >= 3 || tes_triangles) { |
| 846 | gs_vgpr_comp_cnt = 1; /* VGPR1 contains offsets 2, 3 */ |
| 847 | } else { |
| 848 | gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0, 1 */ |
| 849 | } |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 850 | |
Samuel Pitoiset | e68b55f | 2019-07-12 12:17:16 +0200 | [diff] [blame] | 851 | config_out->rsrc1 |= S_00B228_GS_VGPR_COMP_CNT(gs_vgpr_comp_cnt) | |
| 852 | S_00B228_WGP_MODE(1); |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 853 | config_out->rsrc2 |= S_00B22C_ES_VGPR_COMP_CNT(es_vgpr_comp_cnt) | |
Samuel Pitoiset | ed12be1 | 2019-07-15 18:46:48 +0200 | [diff] [blame] | 854 | S_00B22C_LDS_SIZE(config_in->lds_size) | |
| 855 | S_00B22C_OC_LDS_EN(es_stage == MESA_SHADER_TESS_EVAL); |
Samuel Pitoiset | ee21bd7 | 2019-07-05 08:33:06 +0200 | [diff] [blame] | 856 | } else if (pdevice->rad_info.chip_class >= GFX9 && |
| 857 | stage == MESA_SHADER_GEOMETRY) { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 858 | unsigned es_type = info->gs.es_type; |
Samuel Pitoiset | 4e701cf | 2018-01-09 16:01:10 +0100 | [diff] [blame] | 859 | unsigned gs_vgpr_comp_cnt, es_vgpr_comp_cnt; |
| 860 | |
| 861 | if (es_type == MESA_SHADER_VERTEX) { |
Samuel Pitoiset | d8b079e | 2019-06-26 15:11:03 +0200 | [diff] [blame] | 862 | /* VGPR0-3: (VertexID, InstanceID / StepRate0, ...) */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 863 | if (info->vs.needs_instance_id) { |
Samuel Pitoiset | ea337c8 | 2019-07-23 11:52:36 +0200 | [diff] [blame] | 864 | es_vgpr_comp_cnt = pdevice->rad_info.chip_class >= GFX10 ? 3 : 1; |
| 865 | } else { |
| 866 | es_vgpr_comp_cnt = 0; |
| 867 | } |
Samuel Pitoiset | 4e701cf | 2018-01-09 16:01:10 +0100 | [diff] [blame] | 868 | } else if (es_type == MESA_SHADER_TESS_EVAL) { |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 869 | es_vgpr_comp_cnt = info->uses_prim_id ? 3 : 2; |
Samuel Pitoiset | 4e701cf | 2018-01-09 16:01:10 +0100 | [diff] [blame] | 870 | } else { |
Bas Nieuwenhuizen | 0f89f9b | 2018-01-17 23:23:02 +0100 | [diff] [blame] | 871 | unreachable("invalid shader ES type"); |
Samuel Pitoiset | 4e701cf | 2018-01-09 16:01:10 +0100 | [diff] [blame] | 872 | } |
Samuel Pitoiset | 2670ebb | 2017-12-20 20:56:57 +0100 | [diff] [blame] | 873 | |
| 874 | /* If offsets 4, 5 are used, GS_VGPR_COMP_CNT is ignored and |
| 875 | * VGPR[0:4] are always loaded. |
| 876 | */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 877 | if (info->uses_invocation_id) { |
Samuel Pitoiset | 2670ebb | 2017-12-20 20:56:57 +0100 | [diff] [blame] | 878 | gs_vgpr_comp_cnt = 3; /* VGPR3 contains InvocationID. */ |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 879 | } else if (info->uses_prim_id) { |
Samuel Pitoiset | 2670ebb | 2017-12-20 20:56:57 +0100 | [diff] [blame] | 880 | gs_vgpr_comp_cnt = 2; /* VGPR2 contains PrimitiveID. */ |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 881 | } else if (info->gs.vertices_in >= 3) { |
Samuel Pitoiset | b462ceb | 2018-01-05 17:18:52 +0100 | [diff] [blame] | 882 | gs_vgpr_comp_cnt = 1; /* VGPR1 contains offsets 2, 3 */ |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 883 | } else { |
Samuel Pitoiset | b462ceb | 2018-01-05 17:18:52 +0100 | [diff] [blame] | 884 | gs_vgpr_comp_cnt = 0; /* VGPR0 contains offsets 0, 1 */ |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 885 | } |
Samuel Pitoiset | 2670ebb | 2017-12-20 20:56:57 +0100 | [diff] [blame] | 886 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 887 | config_out->rsrc1 |= S_00B228_GS_VGPR_COMP_CNT(gs_vgpr_comp_cnt); |
| 888 | config_out->rsrc2 |= S_00B22C_ES_VGPR_COMP_CNT(es_vgpr_comp_cnt) | |
Bas Nieuwenhuizen | 7469516 | 2019-06-30 01:47:30 +0200 | [diff] [blame] | 889 | S_00B22C_OC_LDS_EN(es_type == MESA_SHADER_TESS_EVAL); |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 890 | } else if (pdevice->rad_info.chip_class >= GFX9 && |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 891 | stage == MESA_SHADER_TESS_CTRL) { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 892 | config_out->rsrc1 |= S_00B428_LS_VGPR_COMP_CNT(vgpr_comp_cnt); |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 893 | } else { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 894 | config_out->rsrc1 |= S_00B128_VGPR_COMP_CNT(vgpr_comp_cnt); |
Samuel Pitoiset | 3a410f0 | 2018-05-11 09:46:46 +0200 | [diff] [blame] | 895 | } |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 896 | } |
| 897 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 898 | struct radv_shader_variant * |
| 899 | radv_shader_variant_create(struct radv_device *device, |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 900 | const struct radv_shader_binary *binary, |
| 901 | bool keep_shader_info) |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 902 | { |
| 903 | struct ac_shader_config config = {0}; |
| 904 | struct ac_rtld_binary rtld_binary = {0}; |
| 905 | struct radv_shader_variant *variant = calloc(1, sizeof(struct radv_shader_variant)); |
| 906 | if (!variant) |
| 907 | return NULL; |
| 908 | |
| 909 | variant->ref_count = 1; |
| 910 | |
| 911 | if (binary->type == RADV_BINARY_TYPE_RTLD) { |
Samuel Pitoiset | 5387667 | 2019-09-03 13:01:54 +0200 | [diff] [blame] | 912 | struct ac_rtld_symbol lds_symbols[2]; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 913 | unsigned num_lds_symbols = 0; |
| 914 | const char *elf_data = (const char *)((struct radv_shader_binary_rtld *)binary)->data; |
| 915 | size_t elf_size = ((struct radv_shader_binary_rtld *)binary)->elf_size; |
| 916 | |
| 917 | if (device->physical_device->rad_info.chip_class >= GFX9 && |
Samuel Pitoiset | 5387667 | 2019-09-03 13:01:54 +0200 | [diff] [blame] | 918 | (binary->stage == MESA_SHADER_GEOMETRY || binary->info.is_ngg) && |
| 919 | !binary->is_gs_copy_shader) { |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 920 | /* We add this symbol even on LLVM <= 8 to ensure that |
| 921 | * shader->config.lds_size is set correctly below. |
| 922 | */ |
| 923 | struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++]; |
| 924 | sym->name = "esgs_ring"; |
Samuel Pitoiset | a2a68d5 | 2019-09-18 09:58:54 +0200 | [diff] [blame] | 925 | sym->size = binary->info.ngg_info.esgs_ring_size; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 926 | sym->align = 64 * 1024; |
Samuel Pitoiset | 5387667 | 2019-09-03 13:01:54 +0200 | [diff] [blame] | 927 | } |
Samuel Pitoiset | 5bbcb3f | 2019-07-11 08:44:16 +0200 | [diff] [blame] | 928 | |
Samuel Pitoiset | 5387667 | 2019-09-03 13:01:54 +0200 | [diff] [blame] | 929 | if (binary->info.is_ngg && |
| 930 | binary->stage == MESA_SHADER_GEOMETRY) { |
| 931 | struct ac_rtld_symbol *sym = &lds_symbols[num_lds_symbols++]; |
| 932 | sym->name = "ngg_emit"; |
| 933 | sym->size = binary->info.ngg_info.ngg_emit_size * 4; |
| 934 | sym->align = 4; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 935 | } |
Samuel Pitoiset | ea38565 | 2019-07-30 18:32:42 +0200 | [diff] [blame] | 936 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 937 | struct ac_rtld_open_info open_info = { |
| 938 | .info = &device->physical_device->rad_info, |
| 939 | .shader_type = binary->stage, |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 940 | .wave_size = binary->info.wave_size, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 941 | .num_parts = 1, |
| 942 | .elf_ptrs = &elf_data, |
| 943 | .elf_sizes = &elf_size, |
| 944 | .num_shared_lds_symbols = num_lds_symbols, |
| 945 | .shared_lds_symbols = lds_symbols, |
| 946 | }; |
| 947 | |
| 948 | if (!ac_rtld_open(&rtld_binary, open_info)) { |
| 949 | free(variant); |
| 950 | return NULL; |
| 951 | } |
| 952 | |
| 953 | if (!ac_rtld_read_config(&rtld_binary, &config)) { |
| 954 | ac_rtld_close(&rtld_binary); |
| 955 | free(variant); |
| 956 | return NULL; |
| 957 | } |
| 958 | |
| 959 | if (rtld_binary.lds_size > 0) { |
| 960 | unsigned alloc_granularity = device->physical_device->rad_info.chip_class >= GFX7 ? 512 : 256; |
| 961 | config.lds_size = align(rtld_binary.lds_size, alloc_granularity) / alloc_granularity; |
| 962 | } |
| 963 | |
| 964 | variant->code_size = rtld_binary.rx_size; |
Connor Abbott | 5dadbab | 2019-08-29 17:15:46 +0200 | [diff] [blame] | 965 | variant->exec_size = rtld_binary.exec_size; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 966 | } else { |
| 967 | assert(binary->type == RADV_BINARY_TYPE_LEGACY); |
| 968 | config = ((struct radv_shader_binary_legacy *)binary)->config; |
Connor Abbott | 5dadbab | 2019-08-29 17:15:46 +0200 | [diff] [blame] | 969 | variant->code_size = radv_get_shader_binary_size(((struct radv_shader_binary_legacy *)binary)->code_size); |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 970 | variant->exec_size = ((struct radv_shader_binary_legacy *)binary)->exec_size; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 971 | } |
| 972 | |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 973 | variant->info = binary->info; |
| 974 | radv_postprocess_config(device->physical_device, &config, &binary->info, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 975 | binary->stage, &variant->config); |
| 976 | |
| 977 | void *dest_ptr = radv_alloc_shader_memory(device, variant); |
| 978 | |
| 979 | if (binary->type == RADV_BINARY_TYPE_RTLD) { |
| 980 | struct radv_shader_binary_rtld* bin = (struct radv_shader_binary_rtld *)binary; |
| 981 | struct ac_rtld_upload_info info = { |
| 982 | .binary = &rtld_binary, |
| 983 | .rx_va = radv_buffer_get_va(variant->bo) + variant->bo_offset, |
| 984 | .rx_ptr = dest_ptr, |
| 985 | }; |
| 986 | |
| 987 | if (!ac_rtld_upload(&info)) { |
| 988 | radv_shader_variant_destroy(device, variant); |
| 989 | ac_rtld_close(&rtld_binary); |
| 990 | return NULL; |
| 991 | } |
| 992 | |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 993 | if (keep_shader_info || |
Samuel Pitoiset | 9343c93 | 2019-07-23 09:55:24 +0200 | [diff] [blame] | 994 | (device->instance->debug_flags & RADV_DEBUG_DUMP_SHADERS)) { |
Timothy Arceri | a20a9d0 | 2019-07-17 14:20:55 +1000 | [diff] [blame] | 995 | const char *disasm_data; |
| 996 | size_t disasm_size; |
| 997 | if (!ac_rtld_get_section_by_name(&rtld_binary, ".AMDGPU.disasm", &disasm_data, &disasm_size)) { |
| 998 | radv_shader_variant_destroy(device, variant); |
| 999 | ac_rtld_close(&rtld_binary); |
| 1000 | return NULL; |
| 1001 | } |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1002 | |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1003 | variant->ir_string = bin->llvm_ir_size ? strdup((const char*)(bin->data + bin->elf_size)) : NULL; |
Timothy Arceri | a20a9d0 | 2019-07-17 14:20:55 +1000 | [diff] [blame] | 1004 | variant->disasm_string = malloc(disasm_size + 1); |
| 1005 | memcpy(variant->disasm_string, disasm_data, disasm_size); |
| 1006 | variant->disasm_string[disasm_size] = 0; |
| 1007 | } |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1008 | |
| 1009 | ac_rtld_close(&rtld_binary); |
| 1010 | } else { |
| 1011 | struct radv_shader_binary_legacy* bin = (struct radv_shader_binary_legacy *)binary; |
| 1012 | memcpy(dest_ptr, bin->data, bin->code_size); |
| 1013 | |
| 1014 | /* Add end-of-code markers for the UMR disassembler. */ |
| 1015 | uint32_t *ptr32 = (uint32_t *)dest_ptr + bin->code_size / 4; |
| 1016 | for (unsigned i = 0; i < DEBUGGER_NUM_MARKERS; i++) |
| 1017 | ptr32[i] = DEBUGGER_END_OF_CODE_MARKER; |
| 1018 | |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1019 | variant->ir_string = bin->ir_size ? strdup((const char*)(bin->data + bin->code_size)) : NULL; |
| 1020 | variant->disasm_string = bin->disasm_size ? strdup((const char*)(bin->data + bin->code_size + bin->ir_size)) : NULL; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1021 | } |
| 1022 | return variant; |
| 1023 | } |
| 1024 | |
Bas Nieuwenhuizen | 5444d3e | 2019-06-01 20:54:35 +0200 | [diff] [blame] | 1025 | static char * |
| 1026 | radv_dump_nir_shaders(struct nir_shader * const *shaders, |
| 1027 | int shader_count) |
| 1028 | { |
| 1029 | char *data = NULL; |
| 1030 | char *ret = NULL; |
| 1031 | size_t size = 0; |
| 1032 | FILE *f = open_memstream(&data, &size); |
| 1033 | if (f) { |
| 1034 | for (int i = 0; i < shader_count; ++i) |
| 1035 | nir_print_shader(shaders[i], f); |
| 1036 | fclose(f); |
| 1037 | } |
| 1038 | |
| 1039 | ret = malloc(size + 1); |
| 1040 | if (ret) { |
| 1041 | memcpy(ret, data, size); |
| 1042 | ret[size] = 0; |
| 1043 | } |
| 1044 | free(data); |
| 1045 | return ret; |
| 1046 | } |
| 1047 | |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1048 | static struct radv_shader_variant * |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1049 | shader_variant_compile(struct radv_device *device, |
| 1050 | struct radv_shader_module *module, |
| 1051 | struct nir_shader * const *shaders, |
| 1052 | int shader_count, |
| 1053 | gl_shader_stage stage, |
Samuel Pitoiset | a9af11f | 2019-09-03 10:29:19 +0200 | [diff] [blame] | 1054 | struct radv_shader_info *info, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1055 | struct radv_nir_compiler_options *options, |
| 1056 | bool gs_copy_shader, |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 1057 | bool keep_shader_info, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1058 | bool use_aco, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1059 | struct radv_shader_binary **binary_out) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1060 | { |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1061 | enum radeon_family chip_family = device->physical_device->rad_info.family; |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1062 | struct radv_shader_binary *binary = NULL; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1063 | |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1064 | options->family = chip_family; |
| 1065 | options->chip_class = device->physical_device->rad_info.chip_class; |
Samuel Pitoiset | 8ade3e4 | 2018-05-11 16:36:02 +0200 | [diff] [blame] | 1066 | options->dump_shader = radv_can_dump_shader(device, module, gs_copy_shader); |
Samuel Pitoiset | d07edf5 | 2018-03-14 10:28:49 +0100 | [diff] [blame] | 1067 | options->dump_preoptir = options->dump_shader && |
Samuel Pitoiset | 33e6e5e | 2018-01-19 12:12:02 +0100 | [diff] [blame] | 1068 | device->instance->debug_flags & RADV_DEBUG_PREOPTIR; |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1069 | options->record_ir = keep_shader_info; |
Samuel Pitoiset | bfca15e | 2018-06-14 14:28:58 +0200 | [diff] [blame] | 1070 | options->check_ir = device->instance->debug_flags & RADV_DEBUG_CHECKIR; |
Dave Airlie | 010d055 | 2018-02-19 07:14:04 +0000 | [diff] [blame] | 1071 | options->tess_offchip_block_dw_size = device->tess_offchip_block_dw_size; |
Samuel Pitoiset | d8a61d3 | 2018-05-16 16:02:04 +0200 | [diff] [blame] | 1072 | options->address32_hi = device->physical_device->rad_info.address32_hi; |
Samuel Pitoiset | 49f5ddd | 2019-08-23 08:55:53 +0200 | [diff] [blame] | 1073 | options->has_ls_vgpr_init_bug = device->physical_device->rad_info.has_ls_vgpr_init_bug; |
Samuel Pitoiset | a15b3bc | 2019-09-09 10:23:30 +0200 | [diff] [blame] | 1074 | options->use_ngg_streamout = device->physical_device->use_ngg_streamout; |
Bas Nieuwenhuizen | 035406e | 2019-08-04 00:48:05 +0200 | [diff] [blame] | 1075 | |
Bas Nieuwenhuizen | ba8d3c3 | 2019-08-04 01:29:53 +0200 | [diff] [blame] | 1076 | if ((stage == MESA_SHADER_GEOMETRY && !options->key.vs_common_out.as_ngg) || |
| 1077 | gs_copy_shader) |
| 1078 | options->wave_size = 64; |
| 1079 | else if (stage == MESA_SHADER_COMPUTE) |
Bas Nieuwenhuizen | 035406e | 2019-08-04 00:48:05 +0200 | [diff] [blame] | 1080 | options->wave_size = device->physical_device->cs_wave_size; |
| 1081 | else if (stage == MESA_SHADER_FRAGMENT) |
| 1082 | options->wave_size = device->physical_device->ps_wave_size; |
| 1083 | else |
| 1084 | options->wave_size = device->physical_device->ge_wave_size; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1085 | |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1086 | if (!use_aco || options->dump_shader || options->record_ir) |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1087 | ac_init_llvm_once(); |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1088 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1089 | if (use_aco) { |
| 1090 | aco_compile_shader(shader_count, shaders, &binary, info, options); |
| 1091 | binary->info = *info; |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1092 | } else { |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1093 | enum ac_target_machine_options tm_options = 0; |
| 1094 | struct ac_llvm_compiler ac_llvm; |
| 1095 | bool thread_compiler; |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1096 | |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1097 | if (options->supports_spill) |
| 1098 | tm_options |= AC_TM_SUPPORTS_SPILL; |
| 1099 | if (device->instance->perftest_flags & RADV_PERFTEST_SISCHED) |
| 1100 | tm_options |= AC_TM_SISCHED; |
| 1101 | if (options->check_ir) |
| 1102 | tm_options |= AC_TM_CHECK_IR; |
| 1103 | if (device->instance->debug_flags & RADV_DEBUG_NO_LOAD_STORE_OPT) |
| 1104 | tm_options |= AC_TM_NO_LOAD_STORE_OPT; |
| 1105 | |
| 1106 | thread_compiler = !(device->instance->debug_flags & RADV_DEBUG_NOTHREADLLVM); |
| 1107 | radv_init_llvm_compiler(&ac_llvm, |
| 1108 | thread_compiler, |
| 1109 | chip_family, tm_options, |
| 1110 | options->wave_size); |
| 1111 | |
| 1112 | if (gs_copy_shader) { |
| 1113 | assert(shader_count == 1); |
| 1114 | radv_compile_gs_copy_shader(&ac_llvm, *shaders, &binary, |
| 1115 | info, options); |
| 1116 | } else { |
| 1117 | radv_compile_nir_shader(&ac_llvm, &binary, info, |
| 1118 | shaders, shader_count, options); |
| 1119 | } |
| 1120 | |
| 1121 | binary->info = *info; |
| 1122 | radv_destroy_llvm_compiler(&ac_llvm, thread_compiler); |
| 1123 | } |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1124 | |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 1125 | struct radv_shader_variant *variant = radv_shader_variant_create(device, binary, |
| 1126 | keep_shader_info); |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1127 | if (!variant) { |
| 1128 | free(binary); |
| 1129 | return NULL; |
| 1130 | } |
Rhys Perry | ec8ced9 | 2019-09-24 15:25:07 +0100 | [diff] [blame] | 1131 | variant->aco_used = use_aco; |
Samuel Pitoiset | 885d757 | 2017-09-01 13:45:33 +0200 | [diff] [blame] | 1132 | |
Bas Nieuwenhuizen | 5ff651c | 2019-07-01 02:19:13 +0200 | [diff] [blame] | 1133 | if (options->dump_shader) { |
| 1134 | fprintf(stderr, "disasm:\n%s\n", variant->disasm_string); |
| 1135 | } |
| 1136 | |
| 1137 | |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 1138 | if (keep_shader_info) { |
Bas Nieuwenhuizen | 5444d3e | 2019-06-01 20:54:35 +0200 | [diff] [blame] | 1139 | variant->nir_string = radv_dump_nir_shaders(shaders, shader_count); |
Samuel Pitoiset | a2a350a | 2017-09-22 16:44:08 +0200 | [diff] [blame] | 1140 | if (!gs_copy_shader && !module->nir) { |
Samuel Pitoiset | 844ae72 | 2017-09-22 16:56:40 +0200 | [diff] [blame] | 1141 | variant->spirv = (uint32_t *)module->data; |
| 1142 | variant->spirv_size = module->size; |
Samuel Pitoiset | a2a350a | 2017-09-22 16:44:08 +0200 | [diff] [blame] | 1143 | } |
Samuel Pitoiset | 885d757 | 2017-09-01 13:45:33 +0200 | [diff] [blame] | 1144 | } |
| 1145 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1146 | if (binary_out) |
| 1147 | *binary_out = binary; |
| 1148 | else |
| 1149 | free(binary); |
| 1150 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1151 | return variant; |
| 1152 | } |
| 1153 | |
| 1154 | struct radv_shader_variant * |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1155 | radv_shader_variant_compile(struct radv_device *device, |
Samuel Pitoiset | a2a350a | 2017-09-22 16:44:08 +0200 | [diff] [blame] | 1156 | struct radv_shader_module *module, |
Bas Nieuwenhuizen | ce03c11 | 2017-10-16 13:18:02 +0200 | [diff] [blame] | 1157 | struct nir_shader *const *shaders, |
| 1158 | int shader_count, |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1159 | struct radv_pipeline_layout *layout, |
Samuel Pitoiset | fbe6945 | 2018-03-13 14:54:04 +0100 | [diff] [blame] | 1160 | const struct radv_shader_variant_key *key, |
Samuel Pitoiset | a9af11f | 2019-09-03 10:29:19 +0200 | [diff] [blame] | 1161 | struct radv_shader_info *info, |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 1162 | bool keep_shader_info, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1163 | bool use_aco, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1164 | struct radv_shader_binary **binary_out) |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1165 | { |
Samuel Pitoiset | fbe6945 | 2018-03-13 14:54:04 +0100 | [diff] [blame] | 1166 | struct radv_nir_compiler_options options = {0}; |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1167 | |
| 1168 | options.layout = layout; |
| 1169 | if (key) |
| 1170 | options.key = *key; |
| 1171 | |
Timothy Arceri | 7664aaf | 2017-10-11 11:59:20 +1100 | [diff] [blame] | 1172 | options.unsafe_math = !!(device->instance->debug_flags & RADV_DEBUG_UNSAFE_MATH); |
Samuel Pitoiset | 1e86eaf | 2018-05-17 09:56:47 +0200 | [diff] [blame] | 1173 | options.supports_spill = true; |
Bas Nieuwenhuizen | 72e7b7a | 2019-08-02 12:40:17 +0200 | [diff] [blame] | 1174 | options.robust_buffer_access = device->robust_buffer_access; |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1175 | |
Samuel Pitoiset | a9af11f | 2019-09-03 10:29:19 +0200 | [diff] [blame] | 1176 | return shader_variant_compile(device, module, shaders, shader_count, shaders[shader_count - 1]->info.stage, info, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1177 | &options, false, keep_shader_info, use_aco, binary_out); |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1178 | } |
| 1179 | |
| 1180 | struct radv_shader_variant * |
| 1181 | radv_create_gs_copy_shader(struct radv_device *device, |
| 1182 | struct nir_shader *shader, |
Samuel Pitoiset | a9af11f | 2019-09-03 10:29:19 +0200 | [diff] [blame] | 1183 | struct radv_shader_info *info, |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1184 | struct radv_shader_binary **binary_out, |
Bas Nieuwenhuizen | 8874af8 | 2019-05-31 01:06:27 +0200 | [diff] [blame] | 1185 | bool keep_shader_info, |
Samuel Pitoiset | 47efc52 | 2017-09-01 12:09:56 +0200 | [diff] [blame] | 1186 | bool multiview) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1187 | { |
Samuel Pitoiset | fbe6945 | 2018-03-13 14:54:04 +0100 | [diff] [blame] | 1188 | struct radv_nir_compiler_options options = {0}; |
Samuel Pitoiset | 92db23f | 2017-09-01 16:51:12 +0200 | [diff] [blame] | 1189 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1190 | options.key.has_multiview_view_index = multiview; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1191 | |
Bas Nieuwenhuizen | 726a31d | 2019-07-01 01:29:24 +0200 | [diff] [blame] | 1192 | return shader_variant_compile(device, NULL, &shader, 1, MESA_SHADER_VERTEX, |
Daniel Schürmann | a70a998 | 2019-09-17 14:35:22 +0200 | [diff] [blame] | 1193 | info, &options, true, keep_shader_info, false, binary_out); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1194 | } |
| 1195 | |
| 1196 | void |
| 1197 | radv_shader_variant_destroy(struct radv_device *device, |
| 1198 | struct radv_shader_variant *variant) |
| 1199 | { |
| 1200 | if (!p_atomic_dec_zero(&variant->ref_count)) |
| 1201 | return; |
| 1202 | |
| 1203 | mtx_lock(&device->shader_slab_mutex); |
| 1204 | list_del(&variant->slab_list); |
| 1205 | mtx_unlock(&device->shader_slab_mutex); |
| 1206 | |
Bas Nieuwenhuizen | 5444d3e | 2019-06-01 20:54:35 +0200 | [diff] [blame] | 1207 | free(variant->nir_string); |
Samuel Pitoiset | 885d757 | 2017-09-01 13:45:33 +0200 | [diff] [blame] | 1208 | free(variant->disasm_string); |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1209 | free(variant->ir_string); |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1210 | free(variant); |
| 1211 | } |
| 1212 | |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1213 | const char * |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 1214 | radv_get_shader_name(struct radv_shader_info *info, |
Samuel Pitoiset | 2b6a089 | 2019-07-11 18:03:55 +0200 | [diff] [blame] | 1215 | gl_shader_stage stage) |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1216 | { |
| 1217 | switch (stage) { |
Samuel Pitoiset | 2b6a089 | 2019-07-11 18:03:55 +0200 | [diff] [blame] | 1218 | case MESA_SHADER_VERTEX: |
| 1219 | if (info->vs.as_ls) |
| 1220 | return "Vertex Shader as LS"; |
| 1221 | else if (info->vs.as_es) |
| 1222 | return "Vertex Shader as ES"; |
| 1223 | else if (info->is_ngg) |
| 1224 | return "Vertex Shader as ESGS"; |
| 1225 | else |
| 1226 | return "Vertex Shader as VS"; |
| 1227 | case MESA_SHADER_TESS_CTRL: |
| 1228 | return "Tessellation Control Shader"; |
| 1229 | case MESA_SHADER_TESS_EVAL: |
| 1230 | if (info->tes.as_es) |
| 1231 | return "Tessellation Evaluation Shader as ES"; |
| 1232 | else if (info->is_ngg) |
| 1233 | return "Tessellation Evaluation Shader as ESGS"; |
| 1234 | else |
| 1235 | return "Tessellation Evaluation Shader as VS"; |
| 1236 | case MESA_SHADER_GEOMETRY: |
| 1237 | return "Geometry Shader"; |
| 1238 | case MESA_SHADER_FRAGMENT: |
| 1239 | return "Pixel Shader"; |
| 1240 | case MESA_SHADER_COMPUTE: |
| 1241 | return "Compute Shader"; |
Samuel Pitoiset | d4d7773 | 2017-09-01 11:41:18 +0200 | [diff] [blame] | 1242 | default: |
| 1243 | return "Unknown shader"; |
| 1244 | }; |
| 1245 | } |
| 1246 | |
Bas Nieuwenhuizen | 739a288 | 2019-06-01 20:25:47 +0200 | [diff] [blame] | 1247 | unsigned |
| 1248 | radv_get_max_workgroup_size(enum chip_class chip_class, |
| 1249 | gl_shader_stage stage, |
| 1250 | const unsigned *sizes) |
| 1251 | { |
| 1252 | switch (stage) { |
| 1253 | case MESA_SHADER_TESS_CTRL: |
| 1254 | return chip_class >= GFX7 ? 128 : 64; |
| 1255 | case MESA_SHADER_GEOMETRY: |
| 1256 | return chip_class >= GFX9 ? 128 : 64; |
| 1257 | case MESA_SHADER_COMPUTE: |
| 1258 | break; |
| 1259 | default: |
| 1260 | return 0; |
| 1261 | } |
| 1262 | |
| 1263 | unsigned max_workgroup_size = sizes[0] * sizes[1] * sizes[2]; |
| 1264 | return max_workgroup_size; |
| 1265 | } |
Bas Nieuwenhuizen | 290ca0c | 2019-06-01 18:46:21 +0200 | [diff] [blame] | 1266 | |
| 1267 | unsigned |
| 1268 | radv_get_max_waves(struct radv_device *device, |
| 1269 | struct radv_shader_variant *variant, |
| 1270 | gl_shader_stage stage) |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1271 | { |
Timothy Arceri | 9b9ccee | 2019-02-01 22:04:39 +1100 | [diff] [blame] | 1272 | enum chip_class chip_class = device->physical_device->rad_info.chip_class; |
Marek Olšák | ccfcb9d | 2019-05-14 22:16:20 -0400 | [diff] [blame] | 1273 | unsigned lds_increment = chip_class >= GFX7 ? 512 : 256; |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 1274 | uint8_t wave_size = variant->info.wave_size; |
Bas Nieuwenhuizen | 290ca0c | 2019-06-01 18:46:21 +0200 | [diff] [blame] | 1275 | struct ac_shader_config *conf = &variant->config; |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1276 | unsigned max_simd_waves; |
| 1277 | unsigned lds_per_wave = 0; |
| 1278 | |
Marek Olšák | ca43006 | 2019-09-12 19:39:02 -0400 | [diff] [blame] | 1279 | max_simd_waves = device->physical_device->rad_info.max_wave64_per_simd; |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1280 | |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1281 | if (stage == MESA_SHADER_FRAGMENT) { |
| 1282 | lds_per_wave = conf->lds_size * lds_increment + |
Samuel Pitoiset | 83499ac | 2019-09-03 17:39:23 +0200 | [diff] [blame] | 1283 | align(variant->info.ps.num_interp * 48, |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1284 | lds_increment); |
Timothy Arceri | 9b9ccee | 2019-02-01 22:04:39 +1100 | [diff] [blame] | 1285 | } else if (stage == MESA_SHADER_COMPUTE) { |
| 1286 | unsigned max_workgroup_size = |
Bas Nieuwenhuizen | 739a288 | 2019-06-01 20:25:47 +0200 | [diff] [blame] | 1287 | radv_get_max_workgroup_size(chip_class, stage, variant->info.cs.block_size); |
Timothy Arceri | 9b9ccee | 2019-02-01 22:04:39 +1100 | [diff] [blame] | 1288 | lds_per_wave = (conf->lds_size * lds_increment) / |
Samuel Pitoiset | ea38565 | 2019-07-30 18:32:42 +0200 | [diff] [blame] | 1289 | DIV_ROUND_UP(max_workgroup_size, wave_size); |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1290 | } |
| 1291 | |
Rhys Perry | 7453c1a | 2019-10-18 21:13:44 +0100 | [diff] [blame] | 1292 | if (conf->num_sgprs) { |
| 1293 | unsigned sgprs = align(conf->num_sgprs, chip_class >= GFX8 ? 16 : 8); |
Samuel Pitoiset | 2f7bb93 | 2018-04-06 14:06:24 +0200 | [diff] [blame] | 1294 | max_simd_waves = |
| 1295 | MIN2(max_simd_waves, |
Marek Olšák | 0692ae3 | 2019-09-12 19:46:02 -0400 | [diff] [blame] | 1296 | device->physical_device->rad_info.num_physical_sgprs_per_simd / |
Rhys Perry | 7453c1a | 2019-10-18 21:13:44 +0100 | [diff] [blame] | 1297 | sgprs); |
| 1298 | } |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1299 | |
Rhys Perry | 7453c1a | 2019-10-18 21:13:44 +0100 | [diff] [blame] | 1300 | if (conf->num_vgprs) { |
| 1301 | unsigned vgprs = align(conf->num_vgprs, wave_size == 32 ? 8 : 4); |
Samuel Pitoiset | 466aba9 | 2018-04-06 14:10:34 +0200 | [diff] [blame] | 1302 | max_simd_waves = |
| 1303 | MIN2(max_simd_waves, |
Rhys Perry | 7453c1a | 2019-10-18 21:13:44 +0100 | [diff] [blame] | 1304 | RADV_NUM_PHYSICAL_VGPRS / vgprs); |
| 1305 | } |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1306 | |
| 1307 | /* LDS is 64KB per CU (4 SIMDs), divided into 16KB blocks per SIMD |
| 1308 | * that PS can use. |
| 1309 | */ |
| 1310 | if (lds_per_wave) |
| 1311 | max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave); |
| 1312 | |
Bas Nieuwenhuizen | 290ca0c | 2019-06-01 18:46:21 +0200 | [diff] [blame] | 1313 | return max_simd_waves; |
| 1314 | } |
| 1315 | |
| 1316 | static void |
| 1317 | generate_shader_stats(struct radv_device *device, |
| 1318 | struct radv_shader_variant *variant, |
| 1319 | gl_shader_stage stage, |
| 1320 | struct _mesa_string_buffer *buf) |
| 1321 | { |
| 1322 | struct ac_shader_config *conf = &variant->config; |
| 1323 | unsigned max_simd_waves = radv_get_max_waves(device, variant, stage); |
| 1324 | |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1325 | if (stage == MESA_SHADER_FRAGMENT) { |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1326 | _mesa_string_buffer_printf(buf, "*** SHADER CONFIG ***\n" |
| 1327 | "SPI_PS_INPUT_ADDR = 0x%04x\n" |
| 1328 | "SPI_PS_INPUT_ENA = 0x%04x\n", |
| 1329 | conf->spi_ps_input_addr, conf->spi_ps_input_ena); |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1330 | } |
| 1331 | |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1332 | _mesa_string_buffer_printf(buf, "*** SHADER STATS ***\n" |
| 1333 | "SGPRS: %d\n" |
| 1334 | "VGPRS: %d\n" |
| 1335 | "Spilled SGPRs: %d\n" |
| 1336 | "Spilled VGPRs: %d\n" |
Samuel Pitoiset | e96e6f6 | 2018-03-01 22:12:56 +0100 | [diff] [blame] | 1337 | "PrivMem VGPRS: %d\n" |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1338 | "Code Size: %d bytes\n" |
| 1339 | "LDS: %d blocks\n" |
| 1340 | "Scratch: %d bytes per wave\n" |
| 1341 | "Max Waves: %d\n" |
| 1342 | "********************\n\n\n", |
| 1343 | conf->num_sgprs, conf->num_vgprs, |
Samuel Pitoiset | e96e6f6 | 2018-03-01 22:12:56 +0100 | [diff] [blame] | 1344 | conf->spilled_sgprs, conf->spilled_vgprs, |
Connor Abbott | 5dadbab | 2019-08-29 17:15:46 +0200 | [diff] [blame] | 1345 | variant->info.private_mem_vgprs, variant->exec_size, |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1346 | conf->lds_size, conf->scratch_bytes_per_wave, |
| 1347 | max_simd_waves); |
| 1348 | } |
| 1349 | |
| 1350 | void |
| 1351 | radv_shader_dump_stats(struct radv_device *device, |
| 1352 | struct radv_shader_variant *variant, |
| 1353 | gl_shader_stage stage, |
| 1354 | FILE *file) |
| 1355 | { |
| 1356 | struct _mesa_string_buffer *buf = _mesa_string_buffer_create(NULL, 256); |
| 1357 | |
| 1358 | generate_shader_stats(device, variant, stage, buf); |
| 1359 | |
Samuel Pitoiset | 2b6a089 | 2019-07-11 18:03:55 +0200 | [diff] [blame] | 1360 | fprintf(file, "\n%s:\n", radv_get_shader_name(&variant->info, stage)); |
Alex Smith | 134a40d | 2017-10-30 08:38:14 +0000 | [diff] [blame] | 1361 | fprintf(file, "%s", buf->buf); |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1362 | |
| 1363 | _mesa_string_buffer_destroy(buf); |
| 1364 | } |
| 1365 | |
| 1366 | VkResult |
| 1367 | radv_GetShaderInfoAMD(VkDevice _device, |
| 1368 | VkPipeline _pipeline, |
| 1369 | VkShaderStageFlagBits shaderStage, |
| 1370 | VkShaderInfoTypeAMD infoType, |
| 1371 | size_t* pInfoSize, |
| 1372 | void* pInfo) |
| 1373 | { |
| 1374 | RADV_FROM_HANDLE(radv_device, device, _device); |
| 1375 | RADV_FROM_HANDLE(radv_pipeline, pipeline, _pipeline); |
| 1376 | gl_shader_stage stage = vk_to_mesa_shader_stage(shaderStage); |
| 1377 | struct radv_shader_variant *variant = pipeline->shaders[stage]; |
| 1378 | struct _mesa_string_buffer *buf; |
| 1379 | VkResult result = VK_SUCCESS; |
| 1380 | |
| 1381 | /* Spec doesn't indicate what to do if the stage is invalid, so just |
| 1382 | * return no info for this. */ |
| 1383 | if (!variant) |
Bas Nieuwenhuizen | 38933c1 | 2018-05-31 01:06:41 +0200 | [diff] [blame] | 1384 | return vk_error(device->instance, VK_ERROR_FEATURE_NOT_PRESENT); |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1385 | |
| 1386 | switch (infoType) { |
| 1387 | case VK_SHADER_INFO_TYPE_STATISTICS_AMD: |
| 1388 | if (!pInfo) { |
| 1389 | *pInfoSize = sizeof(VkShaderStatisticsInfoAMD); |
| 1390 | } else { |
Marek Olšák | ccfcb9d | 2019-05-14 22:16:20 -0400 | [diff] [blame] | 1391 | unsigned lds_multiplier = device->physical_device->rad_info.chip_class >= GFX7 ? 512 : 256; |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1392 | struct ac_shader_config *conf = &variant->config; |
| 1393 | |
| 1394 | VkShaderStatisticsInfoAMD statistics = {}; |
| 1395 | statistics.shaderStageMask = shaderStage; |
Samuel Pitoiset | 466aba9 | 2018-04-06 14:10:34 +0200 | [diff] [blame] | 1396 | statistics.numPhysicalVgprs = RADV_NUM_PHYSICAL_VGPRS; |
Marek Olšák | 0692ae3 | 2019-09-12 19:46:02 -0400 | [diff] [blame] | 1397 | statistics.numPhysicalSgprs = device->physical_device->rad_info.num_physical_sgprs_per_simd; |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1398 | statistics.numAvailableSgprs = statistics.numPhysicalSgprs; |
| 1399 | |
| 1400 | if (stage == MESA_SHADER_COMPUTE) { |
Bas Nieuwenhuizen | 739a288 | 2019-06-01 20:25:47 +0200 | [diff] [blame] | 1401 | unsigned *local_size = variant->info.cs.block_size; |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1402 | unsigned workgroup_size = local_size[0] * local_size[1] * local_size[2]; |
| 1403 | |
| 1404 | statistics.numAvailableVgprs = statistics.numPhysicalVgprs / |
Eric Engestrom | d85fef1 | 2018-06-15 17:49:08 +0100 | [diff] [blame] | 1405 | ceil((double)workgroup_size / statistics.numPhysicalVgprs); |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1406 | |
| 1407 | statistics.computeWorkGroupSize[0] = local_size[0]; |
| 1408 | statistics.computeWorkGroupSize[1] = local_size[1]; |
| 1409 | statistics.computeWorkGroupSize[2] = local_size[2]; |
| 1410 | } else { |
| 1411 | statistics.numAvailableVgprs = statistics.numPhysicalVgprs; |
| 1412 | } |
| 1413 | |
| 1414 | statistics.resourceUsage.numUsedVgprs = conf->num_vgprs; |
| 1415 | statistics.resourceUsage.numUsedSgprs = conf->num_sgprs; |
| 1416 | statistics.resourceUsage.ldsSizePerLocalWorkGroup = 32768; |
| 1417 | statistics.resourceUsage.ldsUsageSizeInBytes = conf->lds_size * lds_multiplier; |
| 1418 | statistics.resourceUsage.scratchMemUsageInBytes = conf->scratch_bytes_per_wave; |
| 1419 | |
| 1420 | size_t size = *pInfoSize; |
| 1421 | *pInfoSize = sizeof(statistics); |
| 1422 | |
| 1423 | memcpy(pInfo, &statistics, MIN2(size, *pInfoSize)); |
| 1424 | |
| 1425 | if (size < *pInfoSize) |
| 1426 | result = VK_INCOMPLETE; |
| 1427 | } |
| 1428 | |
| 1429 | break; |
| 1430 | case VK_SHADER_INFO_TYPE_DISASSEMBLY_AMD: |
| 1431 | buf = _mesa_string_buffer_create(NULL, 1024); |
| 1432 | |
Samuel Pitoiset | 2b6a089 | 2019-07-11 18:03:55 +0200 | [diff] [blame] | 1433 | _mesa_string_buffer_printf(buf, "%s:\n", radv_get_shader_name(&variant->info, stage)); |
Rhys Perry | 3c966fd | 2019-09-25 11:48:04 +0100 | [diff] [blame] | 1434 | _mesa_string_buffer_printf(buf, "%s\n\n", variant->ir_string); |
Alex Smith | de88979 | 2017-10-27 14:25:05 +0100 | [diff] [blame] | 1435 | _mesa_string_buffer_printf(buf, "%s\n\n", variant->disasm_string); |
| 1436 | generate_shader_stats(device, variant, stage, buf); |
| 1437 | |
| 1438 | /* Need to include the null terminator. */ |
| 1439 | size_t length = buf->length + 1; |
| 1440 | |
| 1441 | if (!pInfo) { |
| 1442 | *pInfoSize = length; |
| 1443 | } else { |
| 1444 | size_t size = *pInfoSize; |
| 1445 | *pInfoSize = length; |
| 1446 | |
| 1447 | memcpy(pInfo, buf->buf, MIN2(size, length)); |
| 1448 | |
| 1449 | if (size < length) |
| 1450 | result = VK_INCOMPLETE; |
| 1451 | } |
| 1452 | |
| 1453 | _mesa_string_buffer_destroy(buf); |
| 1454 | break; |
| 1455 | default: |
| 1456 | /* VK_SHADER_INFO_TYPE_BINARY_AMD unimplemented for now. */ |
| 1457 | result = VK_ERROR_FEATURE_NOT_PRESENT; |
| 1458 | break; |
| 1459 | } |
| 1460 | |
| 1461 | return result; |
Samuel Pitoiset | 80b8d9f | 2017-09-05 15:34:07 +0200 | [diff] [blame] | 1462 | } |