.AMD_shader_explicit_vertex_parameter = true,
.AMD_shader_fragment_mask = true,
.AMD_shader_image_load_store_lod = true,
- .AMD_shader_info = true,
.AMD_shader_trinary_minmax = true,
.AMD_texture_gather_bias_lod = true,
#ifdef ANDROID
device->vk.create_sync_for_memory = radv_create_sync_for_memory;
vk_device_set_drm_fd(&device->vk, device->ws->get_fd(device->ws));
- keep_shader_info = device->vk.enabled_extensions.AMD_shader_info;
-
/* With update after bind we can't attach bo's to the command buffer
* from the descriptor set anymore, so we have to use a global BO list.
*/
return spi_ps_input;
}
-VKAPI_ATTR VkResult VKAPI_CALL
-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 *shader = pipeline->shaders[stage];
- VkResult result = VK_SUCCESS;
-
- /* Spec doesn't indicate what to do if the stage is invalid, so just
- * return no info for this. */
- if (!shader)
- return vk_error(device, 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.lds_encode_granularity;
- struct ac_shader_config *conf = &shader->config;
-
- VkShaderStatisticsInfoAMD statistics = {0};
- statistics.shaderStageMask = shaderStage;
- statistics.numPhysicalVgprs =
- device->physical_device->rad_info.num_physical_wave64_vgprs_per_simd;
- statistics.numPhysicalSgprs =
- device->physical_device->rad_info.num_physical_sgprs_per_simd;
- statistics.numAvailableSgprs = statistics.numPhysicalSgprs;
-
- if (stage == MESA_SHADER_COMPUTE) {
- unsigned *local_size = shader->info.cs.block_size;
- unsigned workgroup_size = pipeline->shaders[MESA_SHADER_COMPUTE]->info.workgroup_size;
-
- statistics.numAvailableVgprs =
- statistics.numPhysicalVgprs /
- ceil((double)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: {
- char *out;
- size_t outsize;
- struct u_memstream mem;
- u_memstream_open(&mem, &out, &outsize);
- FILE *const memf = u_memstream_get(&mem);
-
- fprintf(memf, "%s:\n", radv_get_shader_name(&shader->info, stage));
- fprintf(memf, "%s\n\n", shader->ir_string);
- if (shader->disasm_string) {
- fprintf(memf, "%s\n\n", shader->disasm_string);
- }
- radv_dump_shader_stats(device, pipeline, stage, memf);
- u_memstream_close(&mem);
-
- /* Need to include the null terminator. */
- size_t length = outsize + 1;
-
- if (!pInfo) {
- *pInfoSize = length;
- } else {
- size_t size = *pInfoSize;
- *pInfoSize = length;
-
- memcpy(pInfo, out, MIN2(size, length));
-
- if (size < length)
- result = VK_INCOMPLETE;
- }
-
- free(out);
- break;
- }
- default:
- /* VK_SHADER_INFO_TYPE_BINARY_AMD unimplemented for now. */
- result = VK_ERROR_FEATURE_NOT_PRESENT;
- break;
- }
-
- return result;
-}
-
VkResult
radv_dump_shader_stats(struct radv_device *device, struct radv_pipeline *pipeline,
gl_shader_stage stage, FILE *output)