+static uint32_t
+get_total_sgprs(struct radv_device *device)
+{
+ if (device->physical_device->rad_info.chip_class >= VI)
+ return 800;
+ else
+ return 512;
+}
+
+static void
+generate_shader_stats(struct radv_device *device,
+ struct radv_shader_variant *variant,
+ gl_shader_stage stage,
+ struct _mesa_string_buffer *buf)
+{
+ unsigned lds_increment = device->physical_device->rad_info.chip_class >= CIK ? 512 : 256;
+ struct ac_shader_config *conf;
+ unsigned max_simd_waves;
+ unsigned lds_per_wave = 0;
+
+ switch (device->physical_device->rad_info.family) {
+ /* These always have 8 waves: */
+ case CHIP_POLARIS10:
+ case CHIP_POLARIS11:
+ case CHIP_POLARIS12:
+ max_simd_waves = 8;
+ break;
+ default:
+ max_simd_waves = 10;
+ }
+
+ conf = &variant->config;
+
+ if (stage == MESA_SHADER_FRAGMENT) {
+ lds_per_wave = conf->lds_size * lds_increment +
+ align(variant->info.fs.num_interp * 48,
+ lds_increment);
+ }
+
+ if (conf->num_sgprs)
+ max_simd_waves = MIN2(max_simd_waves, get_total_sgprs(device) / conf->num_sgprs);
+
+ if (conf->num_vgprs)
+ max_simd_waves = MIN2(max_simd_waves, 256 / conf->num_vgprs);
+
+ /* LDS is 64KB per CU (4 SIMDs), divided into 16KB blocks per SIMD
+ * that PS can use.
+ */
+ if (lds_per_wave)
+ max_simd_waves = MIN2(max_simd_waves, 16384 / lds_per_wave);
+
+ if (stage == MESA_SHADER_FRAGMENT) {
+ _mesa_string_buffer_printf(buf, "*** SHADER CONFIG ***\n"
+ "SPI_PS_INPUT_ADDR = 0x%04x\n"
+ "SPI_PS_INPUT_ENA = 0x%04x\n",
+ conf->spi_ps_input_addr, conf->spi_ps_input_ena);
+ }
+
+ _mesa_string_buffer_printf(buf, "*** SHADER STATS ***\n"
+ "SGPRS: %d\n"
+ "VGPRS: %d\n"
+ "Spilled SGPRs: %d\n"
+ "Spilled VGPRs: %d\n"
+ "Code Size: %d bytes\n"
+ "LDS: %d blocks\n"
+ "Scratch: %d bytes per wave\n"
+ "Max Waves: %d\n"
+ "********************\n\n\n",
+ conf->num_sgprs, conf->num_vgprs,
+ conf->spilled_sgprs, conf->spilled_vgprs, variant->code_size,
+ conf->lds_size, conf->scratch_bytes_per_wave,
+ max_simd_waves);
+}
+
+void
+radv_shader_dump_stats(struct radv_device *device,
+ struct radv_shader_variant *variant,
+ gl_shader_stage stage,
+ FILE *file)
+{
+ struct _mesa_string_buffer *buf = _mesa_string_buffer_create(NULL, 256);
+
+ generate_shader_stats(device, variant, stage, buf);
+
+ fprintf(file, "\n%s:\n", radv_get_shader_name(variant, stage));
+ fprintf(file, "%s", buf->buf);
+
+ _mesa_string_buffer_destroy(buf);
+}
+
+VkResult
+radv_GetShaderInfoAMD(VkDevice _device,
+ VkPipeline _pipeline,
+ VkShaderStageFlagBits shaderStage,
+ VkShaderInfoTypeAMD infoType,
+ size_t* pInfoSize,
+ void* pInfo)
+{
+ RADV_FROM_HANDLE(radv_device, device, _device);
+ RADV_FROM_HANDLE(radv_pipeline, pipeline, _pipeline);
+ gl_shader_stage stage = vk_to_mesa_shader_stage(shaderStage);
+ struct radv_shader_variant *variant = pipeline->shaders[stage];
+ struct _mesa_string_buffer *buf;
+ VkResult result = VK_SUCCESS;
+
+ /* Spec doesn't indicate what to do if the stage is invalid, so just
+ * return no info for this. */
+ if (!variant)
+ return vk_error(VK_ERROR_FEATURE_NOT_PRESENT);
+
+ switch (infoType) {
+ case VK_SHADER_INFO_TYPE_STATISTICS_AMD:
+ if (!pInfo) {
+ *pInfoSize = sizeof(VkShaderStatisticsInfoAMD);
+ } else {
+ unsigned lds_multiplier = device->physical_device->rad_info.chip_class >= CIK ? 512 : 256;
+ struct ac_shader_config *conf = &variant->config;
+
+ VkShaderStatisticsInfoAMD statistics = {};
+ statistics.shaderStageMask = shaderStage;
+ statistics.numPhysicalVgprs = 256;
+ statistics.numPhysicalSgprs = get_total_sgprs(device);
+ statistics.numAvailableSgprs = statistics.numPhysicalSgprs;
+
+ if (stage == MESA_SHADER_COMPUTE) {
+ unsigned *local_size = variant->nir->info.cs.local_size;
+ unsigned workgroup_size = local_size[0] * local_size[1] * local_size[2];
+
+ statistics.numAvailableVgprs = statistics.numPhysicalVgprs /
+ ceil(workgroup_size / statistics.numPhysicalVgprs);
+
+ statistics.computeWorkGroupSize[0] = local_size[0];
+ statistics.computeWorkGroupSize[1] = local_size[1];
+ statistics.computeWorkGroupSize[2] = local_size[2];
+ } else {
+ statistics.numAvailableVgprs = statistics.numPhysicalVgprs;
+ }
+
+ statistics.resourceUsage.numUsedVgprs = conf->num_vgprs;
+ statistics.resourceUsage.numUsedSgprs = conf->num_sgprs;
+ statistics.resourceUsage.ldsSizePerLocalWorkGroup = 32768;
+ statistics.resourceUsage.ldsUsageSizeInBytes = conf->lds_size * lds_multiplier;
+ statistics.resourceUsage.scratchMemUsageInBytes = conf->scratch_bytes_per_wave;
+
+ size_t size = *pInfoSize;
+ *pInfoSize = sizeof(statistics);
+
+ memcpy(pInfo, &statistics, MIN2(size, *pInfoSize));
+
+ if (size < *pInfoSize)
+ result = VK_INCOMPLETE;
+ }
+
+ break;
+ case VK_SHADER_INFO_TYPE_DISASSEMBLY_AMD:
+ buf = _mesa_string_buffer_create(NULL, 1024);
+
+ _mesa_string_buffer_printf(buf, "%s:\n", radv_get_shader_name(variant, stage));
+ _mesa_string_buffer_printf(buf, "%s\n\n", variant->disasm_string);
+ generate_shader_stats(device, variant, stage, buf);
+
+ /* Need to include the null terminator. */
+ size_t length = buf->length + 1;
+
+ if (!pInfo) {
+ *pInfoSize = length;
+ } else {
+ size_t size = *pInfoSize;
+ *pInfoSize = length;
+
+ memcpy(pInfo, buf->buf, MIN2(size, length));
+
+ if (size < length)
+ result = VK_INCOMPLETE;
+ }
+
+ _mesa_string_buffer_destroy(buf);
+ break;
+ default:
+ /* VK_SHADER_INFO_TYPE_BINARY_AMD unimplemented for now. */
+ result = VK_ERROR_FEATURE_NOT_PRESENT;
+ break;
+ }
+
+ return result;
+}