/* * Copyright 2013 Advanced Micro Devices, Inc. * * SPDX-License-Identifier: MIT */ #include "ac_rtld.h" #include "nir/tgsi_to_nir.h" #include "si_build_pm4.h" #include "si_shader_internal.h" #include "util/u_async_debug.h" #include "util/u_memory.h" #include "util/u_upload_mgr.h" #include "si_tracepoints.h" #include "nir.h" #define COMPUTE_DBG(sscreen, fmt, args...) \ do { \ if ((sscreen->shader_debug_flags & DBG(COMPUTE))) \ fprintf(stderr, fmt, ##args); \ } while (0); /* Asynchronous compute shader compilation. */ static void si_create_compute_state_async(void *job, void *gdata, int thread_index) { struct si_compute *program = (struct si_compute *)job; struct si_shader_selector *sel = &program->sel; struct si_shader *shader = &program->shader; struct ac_llvm_compiler **compiler; struct util_debug_callback *debug = &sel->compiler_ctx_state.debug; struct si_screen *sscreen = sel->screen; assert(!debug->debug_message || debug->async); assert(thread_index >= 0); assert(thread_index < ARRAY_SIZE(sscreen->compiler)); compiler = &sscreen->compiler[thread_index]; si_nir_scan_shader(sscreen, sel->nir, &sel->info, false); if (!sel->nir->info.use_aco_amd && !*compiler) *compiler = si_create_llvm_compiler(sscreen); si_get_active_slot_masks(sscreen, &sel->info, &sel->active_const_and_shader_buffers, &sel->active_samplers_and_images); program->shader.is_monolithic = true; program->shader.wave_size = si_determine_wave_size(sscreen, &program->shader); /* Variable block sizes need 10 bits (1 + log2(SI_MAX_VARIABLE_THREADS_PER_BLOCK)) per dim. * We pack them into a single user SGPR. */ unsigned user_sgprs = SI_NUM_RESOURCE_SGPRS + (sel->info.uses_grid_size ? 3 : 0) + (sel->info.uses_variable_block_size ? 1 : 0) + sel->nir->info.cs.user_data_components_amd; /* Fast path for compute shaders - some descriptors passed via user SGPRs. */ /* Shader buffers in user SGPRs. */ for (unsigned i = 0; i < MIN2(3, sel->nir->info.num_ssbos) && user_sgprs <= 12; i++) { user_sgprs = align(user_sgprs, 4); if (i == 0) sel->cs_shaderbufs_sgpr_index = user_sgprs; user_sgprs += 4; sel->cs_num_shaderbufs_in_user_sgprs++; } /* Images in user SGPRs. */ unsigned non_fmask_images = BITFIELD_MASK(sel->nir->info.num_images); /* Remove images with FMASK from the bitmask. We only care about the first * 3 anyway, so we can take msaa_images[0] and ignore the rest. */ if (sscreen->info.gfx_level < GFX11) non_fmask_images &= ~sel->nir->info.msaa_images[0]; for (unsigned i = 0; i < 3 && non_fmask_images & (1 << i); i++) { unsigned num_sgprs = BITSET_TEST(sel->nir->info.image_buffers, i) ? 4 : 8; if (align(user_sgprs, num_sgprs) + num_sgprs > 16) break; user_sgprs = align(user_sgprs, num_sgprs); if (i == 0) sel->cs_images_sgpr_index = user_sgprs; user_sgprs += num_sgprs; sel->cs_num_images_in_user_sgprs++; } sel->cs_images_num_sgprs = user_sgprs - sel->cs_images_sgpr_index; assert(user_sgprs <= 16); unsigned char ir_sha1_cache_key[20]; si_get_ir_cache_key(sel, false, false, shader->wave_size, ir_sha1_cache_key); /* Try to load the shader from the shader cache. */ simple_mtx_lock(&sscreen->shader_cache_mutex); if (si_shader_cache_load_shader(sscreen, ir_sha1_cache_key, shader)) { simple_mtx_unlock(&sscreen->shader_cache_mutex); shader->complete_shader_binary_size = si_get_shader_binary_size(sscreen, shader); if (!si_shader_binary_upload(sscreen, shader, 0)) program->shader.compilation_failed = true; si_shader_dump_stats_for_shader_db(sscreen, shader, debug); si_shader_dump(sscreen, shader, debug, stderr, true); } else { simple_mtx_unlock(&sscreen->shader_cache_mutex); if (!si_create_shader_variant(sscreen, *compiler, &program->shader, debug)) { program->shader.compilation_failed = true; return; } shader->config.rsrc1 = S_00B848_VGPRS((shader->config.num_vgprs - 1) / ((shader->wave_size == 32 || sscreen->info.wave64_vgpr_alloc_granularity == 8) ? 8 : 4)) | S_00B848_DX10_CLAMP(sscreen->info.gfx_level < GFX12) | S_00B848_MEM_ORDERED(si_shader_mem_ordered(shader)) | S_00B848_FLOAT_MODE(shader->config.float_mode) | /* This is needed for CWSR, but it causes halts to work differently. */ S_00B848_PRIV(sscreen->info.gfx_level == GFX11); if (sscreen->info.gfx_level < GFX10) { shader->config.rsrc1 |= S_00B848_SGPRS((shader->config.num_sgprs - 1) / 8); } shader->config.rsrc2 = S_00B84C_USER_SGPR(user_sgprs) | S_00B84C_SCRATCH_EN(shader->config.scratch_bytes_per_wave > 0) | S_00B84C_TGID_X_EN(sel->info.uses_block_id[0]) | S_00B84C_TGID_Y_EN(sel->info.uses_block_id[1]) | S_00B84C_TGID_Z_EN(sel->info.uses_block_id[2]) | S_00B84C_TG_SIZE_EN(sel->info.uses_tg_size) | S_00B84C_TIDIG_COMP_CNT(sel->info.uses_thread_id[2] ? 2 : sel->info.uses_thread_id[1] ? 1 : 0) | S_00B84C_LDS_SIZE(shader->config.lds_size); /* COMPUTE_PGM_RSRC3 is only present on GFX10+ and GFX940+. */ shader->config.rsrc3 = S_00B8A0_SHARED_VGPR_CNT(shader->config.num_shared_vgprs / 8); if (sscreen->info.gfx_level >= GFX12) shader->config.rsrc3 |= S_00B8A0_INST_PREF_SIZE_GFX12(si_get_shader_prefetch_size(shader)); else if (sscreen->info.gfx_level >= GFX11) shader->config.rsrc3 |= S_00B8A0_INST_PREF_SIZE_GFX11(si_get_shader_prefetch_size(shader)); simple_mtx_lock(&sscreen->shader_cache_mutex); si_shader_cache_insert_shader(sscreen, ir_sha1_cache_key, shader, true); simple_mtx_unlock(&sscreen->shader_cache_mutex); } ralloc_free(sel->nir); sel->nir = NULL; } static void *si_create_compute_state(struct pipe_context *ctx, const struct pipe_compute_state *cso) { struct si_context *sctx = (struct si_context *)ctx; struct si_screen *sscreen = (struct si_screen *)ctx->screen; struct si_compute *program = CALLOC_STRUCT(si_compute); struct si_shader_selector *sel = &program->sel; pipe_reference_init(&sel->base.reference, 1); sel->stage = MESA_SHADER_COMPUTE; sel->screen = sscreen; simple_mtx_init(&sel->mutex, mtx_plain); sel->const_and_shader_buf_descriptors_index = si_const_and_shader_buffer_descriptors_idx(PIPE_SHADER_COMPUTE); sel->sampler_and_images_descriptors_index = si_sampler_and_image_descriptors_idx(PIPE_SHADER_COMPUTE); sel->info.base.shared_size = cso->static_shared_mem; program->shader.selector = &program->sel; if (cso->ir_type == PIPE_SHADER_IR_TGSI) { sel->nir = tgsi_to_nir(cso->prog, ctx->screen, true); } else { assert(cso->ir_type == PIPE_SHADER_IR_NIR); sel->nir = (struct nir_shader *)cso->prog; } sel->nir->info.shared_size = cso->static_shared_mem; if (si_can_dump_shader(sscreen, sel->stage, SI_DUMP_INIT_NIR)) nir_print_shader(sel->nir, stderr); sel->compiler_ctx_state.debug = sctx->debug; sel->compiler_ctx_state.is_debug_context = sctx->is_debug; p_atomic_inc(&sscreen->num_shaders_created); si_schedule_initial_compile(sctx, MESA_SHADER_COMPUTE, &sel->ready, &sel->compiler_ctx_state, program, si_create_compute_state_async); return program; } static void si_get_compute_state_info(struct pipe_context *ctx, void *state, struct pipe_compute_state_object_info *info) { struct si_compute *program = (struct si_compute *)state; struct si_shader_selector *sel = &program->sel; /* Wait because we need the compilation to finish first */ util_queue_fence_wait(&sel->ready); uint8_t wave_size = program->shader.wave_size; info->private_memory = DIV_ROUND_UP(program->shader.config.scratch_bytes_per_wave, wave_size); info->preferred_simd_size = wave_size; info->simd_sizes = wave_size; info->max_threads = si_get_max_workgroup_size(&program->shader); } static void si_bind_compute_state(struct pipe_context *ctx, void *state) { struct si_context *sctx = (struct si_context *)ctx; struct si_compute *program = (struct si_compute *)state; struct si_shader_selector *sel = &program->sel; sctx->cs_shader_state.program = program; if (!program) return; /* Wait because we need active slot usage masks. */ util_queue_fence_wait(&sel->ready); si_update_common_shader_state(sctx, sel, PIPE_SHADER_COMPUTE); sctx->compute_shaderbuf_sgprs_dirty = true; sctx->compute_image_sgprs_dirty = true; if (unlikely((sctx->screen->debug_flags & DBG(SQTT)) && sctx->sqtt)) { uint32_t pipeline_code_hash = _mesa_hash_data_with_seed( program->shader.binary.code_buffer, program->shader.binary.code_size, 0); if (!si_sqtt_pipeline_is_registered(sctx->sqtt, pipeline_code_hash)) { /* Short lived fake pipeline: we don't need to reupload the compute shaders, * as we do for the gfx ones so just create a temp pipeline to be able to * call si_sqtt_register_pipeline, and then drop it. */ struct si_sqtt_fake_pipeline pipeline = { 0 }; pipeline.code_hash = pipeline_code_hash; pipeline.bo = program->shader.bo; si_sqtt_register_pipeline(sctx, &pipeline, NULL); } si_sqtt_describe_pipeline_bind(sctx, pipeline_code_hash, 1); } } static void si_set_global_binding(struct pipe_context *ctx, unsigned first, unsigned n, struct pipe_resource **resources, uint32_t **handles) { unsigned i; struct si_context *sctx = (struct si_context *)ctx; if (first + n > sctx->max_global_buffers) { unsigned old_max = sctx->max_global_buffers; sctx->max_global_buffers = first + n; sctx->global_buffers = realloc( sctx->global_buffers, sctx->max_global_buffers * sizeof(sctx->global_buffers[0])); if (!sctx->global_buffers) { mesa_loge("failed to allocate compute global_buffers"); return; } memset(&sctx->global_buffers[old_max], 0, (sctx->max_global_buffers - old_max) * sizeof(sctx->global_buffers[0])); } if (!resources) { for (i = 0; i < n; i++) { pipe_resource_reference(&sctx->global_buffers[first + i], NULL); } return; } for (i = 0; i < n; i++) { uint64_t va; uint32_t offset; pipe_resource_reference(&sctx->global_buffers[first + i], resources[i]); va = si_resource(resources[i])->gpu_address; offset = util_le32_to_cpu(*handles[i]); va += offset; va = util_cpu_to_le64(va); memcpy(handles[i], &va, sizeof(va)); } } static bool si_setup_compute_scratch_buffer(struct si_context *sctx, struct si_shader *shader) { uint64_t scratch_bo_size = sctx->compute_scratch_buffer ? sctx->compute_scratch_buffer->b.b.width0 : 0; uint64_t scratch_needed = (uint64_t)sctx->max_seen_compute_scratch_bytes_per_wave * sctx->screen->info.max_scratch_waves; assert(scratch_needed); if (scratch_bo_size < scratch_needed) { si_resource_reference(&sctx->compute_scratch_buffer, NULL); sctx->compute_scratch_buffer = si_aligned_buffer_create(&sctx->screen->b, PIPE_RESOURCE_FLAG_UNMAPPABLE | SI_RESOURCE_FLAG_DRIVER_INTERNAL | SI_RESOURCE_FLAG_DISCARDABLE, PIPE_USAGE_DEFAULT, scratch_needed, sctx->screen->info.pte_fragment_size); if (!sctx->compute_scratch_buffer) return false; } /* Set the scratch address in the shader binary. */ if (!sctx->screen->info.has_scratch_base_registers) { uint64_t scratch_va = sctx->compute_scratch_buffer->gpu_address; if (shader->scratch_va != scratch_va) { if (!si_shader_binary_upload(sctx->screen, shader, scratch_va)) return false; shader->scratch_va = scratch_va; } } return true; } static bool si_switch_compute_shader(struct si_context *sctx, struct si_compute *program, struct si_shader *shader, bool *prefetch, unsigned variable_shared_size) { struct radeon_cmdbuf *cs = &sctx->gfx_cs; const struct ac_shader_config *config = &shader->config; unsigned rsrc2; *prefetch = false; if (sctx->cs_shader_state.emitted_program == program && sctx->cs_shader_state.variable_shared_size == variable_shared_size) return true; /* copy rsrc2 so we don't have to change it inside the si_shader object */ rsrc2 = config->rsrc2; /* only do this for OpenCL */ if (variable_shared_size) { unsigned shared_size = program->sel.info.base.shared_size + variable_shared_size; unsigned lds_blocks = 0; if (sctx->gfx_level <= GFX6) { lds_blocks += align(shared_size, 256) >> 8; } else { lds_blocks += align(shared_size, 512) >> 9; } /* TODO: use si_multiwave_lds_size_workaround */ assert(lds_blocks <= 0xFF); rsrc2 &= C_00B84C_LDS_SIZE; rsrc2 |= S_00B84C_LDS_SIZE(lds_blocks); } if (config->scratch_bytes_per_wave) { /* Prevent race conditions for accesses to shader->scratch_va and shader->bo, which * can change when scratch_va is updated. Any accesses to shader->bo must also be inside * the lock. * * TODO: This lock could be removed if the scratch address was passed via user SGPRs instead * of the shader binary. */ if (!sctx->screen->info.has_scratch_base_registers) simple_mtx_lock(&shader->selector->mutex); /* Update max_seen_compute_scratch_bytes_per_wave and compute_tmpring_size. */ si_get_scratch_tmpring_size(sctx, config->scratch_bytes_per_wave, true, &sctx->compute_tmpring_size); if (!si_setup_compute_scratch_buffer(sctx, shader)) return false; radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, sctx->compute_scratch_buffer, RADEON_USAGE_READWRITE | RADEON_PRIO_SCRATCH_BUFFER); } uint64_t shader_va = shader->bo->gpu_address; radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, shader->bo, RADEON_USAGE_READ | RADEON_PRIO_SHADER_BINARY); /* shader->bo can't be used after this if the scratch address is inserted into the shader * binary. */ if (config->scratch_bytes_per_wave && !sctx->screen->info.has_scratch_base_registers) simple_mtx_unlock(&shader->selector->mutex); if (sctx->gfx_level >= GFX12) { gfx12_push_compute_sh_reg(R_00B830_COMPUTE_PGM_LO, shader_va >> 8); gfx12_opt_push_compute_sh_reg(R_00B848_COMPUTE_PGM_RSRC1, SI_TRACKED_COMPUTE_PGM_RSRC1, config->rsrc1); gfx12_opt_push_compute_sh_reg(R_00B84C_COMPUTE_PGM_RSRC2, SI_TRACKED_COMPUTE_PGM_RSRC2, rsrc2); gfx12_opt_push_compute_sh_reg(R_00B8A0_COMPUTE_PGM_RSRC3, SI_TRACKED_COMPUTE_PGM_RSRC3, config->rsrc3); gfx12_opt_push_compute_sh_reg(R_00B860_COMPUTE_TMPRING_SIZE, SI_TRACKED_COMPUTE_TMPRING_SIZE, sctx->compute_tmpring_size); if (config->scratch_bytes_per_wave) { gfx12_opt_push_compute_sh_reg(R_00B840_COMPUTE_DISPATCH_SCRATCH_BASE_LO, SI_TRACKED_COMPUTE_DISPATCH_SCRATCH_BASE_LO, sctx->compute_scratch_buffer->gpu_address >> 8); gfx12_opt_push_compute_sh_reg(R_00B844_COMPUTE_DISPATCH_SCRATCH_BASE_HI, SI_TRACKED_COMPUTE_DISPATCH_SCRATCH_BASE_HI, sctx->compute_scratch_buffer->gpu_address >> 40); } } else if (sctx->screen->info.has_set_sh_pairs_packed) { gfx11_push_compute_sh_reg(R_00B830_COMPUTE_PGM_LO, shader_va >> 8); gfx11_opt_push_compute_sh_reg(R_00B848_COMPUTE_PGM_RSRC1, SI_TRACKED_COMPUTE_PGM_RSRC1, config->rsrc1); gfx11_opt_push_compute_sh_reg(R_00B84C_COMPUTE_PGM_RSRC2, SI_TRACKED_COMPUTE_PGM_RSRC2, rsrc2); gfx11_opt_push_compute_sh_reg(R_00B8A0_COMPUTE_PGM_RSRC3, SI_TRACKED_COMPUTE_PGM_RSRC3, config->rsrc3); gfx11_opt_push_compute_sh_reg(R_00B860_COMPUTE_TMPRING_SIZE, SI_TRACKED_COMPUTE_TMPRING_SIZE, sctx->compute_tmpring_size); if (config->scratch_bytes_per_wave) { gfx11_opt_push_compute_sh_reg(R_00B840_COMPUTE_DISPATCH_SCRATCH_BASE_LO, SI_TRACKED_COMPUTE_DISPATCH_SCRATCH_BASE_LO, sctx->compute_scratch_buffer->gpu_address >> 8); gfx11_opt_push_compute_sh_reg(R_00B844_COMPUTE_DISPATCH_SCRATCH_BASE_HI, SI_TRACKED_COMPUTE_DISPATCH_SCRATCH_BASE_HI, sctx->compute_scratch_buffer->gpu_address >> 40); } } else { radeon_begin(cs); radeon_set_sh_reg(R_00B830_COMPUTE_PGM_LO, shader_va >> 8); radeon_opt_set_sh_reg2(R_00B848_COMPUTE_PGM_RSRC1, SI_TRACKED_COMPUTE_PGM_RSRC1, config->rsrc1, rsrc2); radeon_opt_set_sh_reg(R_00B860_COMPUTE_TMPRING_SIZE, SI_TRACKED_COMPUTE_TMPRING_SIZE, sctx->compute_tmpring_size); if (config->scratch_bytes_per_wave && sctx->screen->info.has_scratch_base_registers) { radeon_opt_set_sh_reg2(R_00B840_COMPUTE_DISPATCH_SCRATCH_BASE_LO, SI_TRACKED_COMPUTE_DISPATCH_SCRATCH_BASE_LO, sctx->compute_scratch_buffer->gpu_address >> 8, sctx->compute_scratch_buffer->gpu_address >> 40); } if (sctx->gfx_level >= GFX10) { radeon_opt_set_sh_reg(R_00B8A0_COMPUTE_PGM_RSRC3, SI_TRACKED_COMPUTE_PGM_RSRC3, config->rsrc3); } radeon_end(); } COMPUTE_DBG(sctx->screen, "COMPUTE_PGM_RSRC1: 0x%08x " "COMPUTE_PGM_RSRC2: 0x%08x\n", config->rsrc1, config->rsrc2); sctx->cs_shader_state.emitted_program = program; sctx->cs_shader_state.variable_shared_size = variable_shared_size; *prefetch = true; return true; } static void si_setup_nir_user_data(struct si_context *sctx, const struct pipe_grid_info *info) { struct si_compute *program = sctx->cs_shader_state.program; struct si_shader_selector *sel = &program->sel; struct radeon_cmdbuf *cs = &sctx->gfx_cs; unsigned grid_size_reg = R_00B900_COMPUTE_USER_DATA_0 + 4 * SI_NUM_RESOURCE_SGPRS; unsigned block_size_reg = grid_size_reg + /* 12 bytes = 3 dwords. */ 12 * sel->info.uses_grid_size; unsigned cs_user_data_reg = block_size_reg + 4 * program->sel.info.uses_variable_block_size; if (sel->info.uses_grid_size && info->indirect) { for (unsigned i = 0; i < 3; ++i) { si_cp_copy_data(sctx, &sctx->gfx_cs, COPY_DATA_REG, NULL, (grid_size_reg >> 2) + i, COPY_DATA_SRC_MEM, si_resource(info->indirect), info->indirect_offset + 4 * i); } } if (sctx->gfx_level >= GFX12) { if (sel->info.uses_grid_size && !info->indirect) { gfx12_push_compute_sh_reg(grid_size_reg, info->grid[0]); gfx12_push_compute_sh_reg(grid_size_reg + 4, info->grid[1]); gfx12_push_compute_sh_reg(grid_size_reg + 8, info->grid[2]); } if (sel->info.uses_variable_block_size) { uint32_t value = info->block[0] | (info->block[1] << 10) | (info->block[2] << 20); gfx12_push_compute_sh_reg(block_size_reg, value); } if (sel->info.base.cs.user_data_components_amd) { unsigned num = sel->info.base.cs.user_data_components_amd; for (unsigned i = 0; i < num; i++) gfx12_push_compute_sh_reg(cs_user_data_reg + i * 4, sctx->cs_user_data[i]); } } else if (sctx->screen->info.has_set_sh_pairs_packed) { if (sel->info.uses_grid_size && !info->indirect) { gfx11_push_compute_sh_reg(grid_size_reg, info->grid[0]); gfx11_push_compute_sh_reg(grid_size_reg + 4, info->grid[1]); gfx11_push_compute_sh_reg(grid_size_reg + 8, info->grid[2]); } if (sel->info.uses_variable_block_size) { uint32_t value = info->block[0] | (info->block[1] << 10) | (info->block[2] << 20); gfx11_push_compute_sh_reg(block_size_reg, value); } if (sel->info.base.cs.user_data_components_amd) { unsigned num = sel->info.base.cs.user_data_components_amd; for (unsigned i = 0; i < num; i++) gfx11_push_compute_sh_reg(cs_user_data_reg + i * 4, sctx->cs_user_data[i]); } } else { radeon_begin(cs); if (sel->info.uses_grid_size && !info->indirect) { radeon_set_sh_reg_seq(grid_size_reg, 3); radeon_emit(info->grid[0]); radeon_emit(info->grid[1]); radeon_emit(info->grid[2]); } if (sel->info.uses_variable_block_size) { uint32_t value = info->block[0] | (info->block[1] << 10) | (info->block[2] << 20); radeon_set_sh_reg(block_size_reg, value); } if (sel->info.base.cs.user_data_components_amd) { unsigned num = sel->info.base.cs.user_data_components_amd; radeon_set_sh_reg_seq(cs_user_data_reg, num); radeon_emit_array(sctx->cs_user_data, num); } radeon_end(); } } static bool si_get_2d_interleave_size(const struct pipe_grid_info *info, unsigned *log_x, unsigned *log_y) { /* The following code produces this behavior: * * WG size | WG block/SE | Thread block/SE * ( 1, 32) = 32 | (16, 1) = 16 | ( 16, 32) = 512 * ( 2, 16) = 32 | ( 8, 2) = 16 | ( 16, 32) = 512 * ( 2, 32) = 64 | (16, 1) = 16 | ( 32, 32) = 1024 * ( 4, 8) = 32 | ( 4, 4) = 16 | ( 16, 32) = 512 * ( 4, 16) = 64 | ( 8, 2) = 16 | ( 32, 32) = 1024 * ( 4, 32) = 128 | ( 8, 1) = 8 | ( 32, 32) = 1024 * ( 8, 4) = 32 | ( 2, 8) = 16 | ( 16, 32) = 512 * ( 8, 8) = 64 | ( 4, 4) = 16 | ( 32, 32) = 1024 * ( 8, 16) = 128 | ( 4, 2) = 8 | ( 32, 32) = 1024 * ( 8, 32) = 256 | ( 4, 1) = 4 | ( 32, 32) = 1024 * (16, 2) = 32 | ( 1, 16) = 16 | ( 16, 32) = 512 * (16, 4) = 64 | ( 2, 8) = 16 | ( 32, 32) = 1024 * (16, 8) = 128 | ( 2, 4) = 8 | ( 32, 32) = 1024 * (16, 16) = 256 | ( 2, 2) = 4 | ( 32, 32) = 1024 * (16, 32) = 512 | ( 2, 1) = 2 | ( 32, 32) = 1024 * (32, 1) = 32 | ( 1, 16) = 16 | ( 32, 16) = 512 * (32, 2) = 64 | ( 1, 16) = 16 | ( 32, 32) = 1024 * (32, 4) = 128 | ( 1, 8) = 8 | ( 32, 32) = 1024 * (32, 8) = 256 | ( 1, 4) = 4 | ( 32, 32) = 1024 * (32, 16) = 512 | ( 1, 2) = 2 | ( 32, 32) = 1024 * * For 3D workgroups, the total 2D thread count is divided by Z. * Example with Z=8, showing only a 2D slice of the grid: * * WG size | WG block/SE | Thread block/SE * ( 1, 32) = 32 | ( 4, 1) = 4 | ( 4, 32) = 128 * ( 2, 16) = 32 | ( 4, 1) = 4 | ( 8, 16) = 128 * ( 2, 32) = 64 | ( 2, 1) = 2 | ( 4, 32) = 128 * ( 4, 8) = 32 | ( 2, 2) = 4 | ( 8, 16) = 128 * ( 4, 16) = 64 | ( 2, 1) = 2 | ( 8, 16) = 128 * ( 8, 4) = 32 | ( 1, 4) = 4 | ( 8, 16) = 128 * ( 8, 8) = 64 | ( 1, 2) = 2 | ( 8, 16) = 128 * (16, 2) = 32 | ( 1, 4) = 4 | ( 16, 8) = 128 * (16, 4) = 64 | ( 1, 2) = 2 | ( 16, 8) = 128 * (32, 1) = 32 | ( 1, 4) = 4 | ( 32, 4) = 128 * (32, 2) = 64 | ( 1, 2) = 2 | ( 32, 4) = 128 * * It tries to find a WG block size that corresponds to (N, N) or (N, 2*N) threads, * but it's limited by the maximum WGs/SE, which is 16, and the number of threads/SE, * which we set to 1024. */ unsigned max_threads_per_se = 1024; unsigned threads_per_threadgroup = info->block[0] * info->block[1] * info->block[2]; unsigned workgroups_per_se = MIN2(max_threads_per_se / threads_per_threadgroup, 16); unsigned log_workgroups_per_se = util_logbase2(workgroups_per_se); if (!log_workgroups_per_se) return false; assert(log_workgroups_per_se <= 4); *log_x = MIN2(log_workgroups_per_se, 4); *log_y = log_workgroups_per_se - *log_x; while (*log_x > 0 && *log_y < 4 && info->block[0] * (1 << *log_x) > info->block[1] * (1 << *log_y)) { (*log_x)--; (*log_y)++; } assert(*log_x + *log_y <= 4); return true; } static void si_emit_dispatch_packets(struct si_context *sctx, const struct pipe_grid_info *info) { struct si_screen *sscreen = sctx->screen; struct radeon_cmdbuf *cs = &sctx->gfx_cs; bool render_cond_bit = sctx->render_cond_enabled; unsigned threads_per_threadgroup = info->block[0] * info->block[1] * info->block[2]; unsigned waves_per_threadgroup = DIV_ROUND_UP(threads_per_threadgroup, sctx->cs_shader_state.program->shader.wave_size); unsigned threadgroups_per_cu = 1; if (sctx->gfx_level >= GFX10 && waves_per_threadgroup == 1) threadgroups_per_cu = 2; if (unlikely(sctx->sqtt_enabled)) { if (info->indirect) { si_sqtt_write_event_marker(sctx, &sctx->gfx_cs, EventCmdDispatchIndirect, UINT_MAX, UINT_MAX, UINT_MAX); } else { si_write_event_with_dims_marker(sctx, &sctx->gfx_cs, EventCmdDispatch, info->grid[0], info->grid[1], info->grid[2]); } } radeon_begin(cs); unsigned compute_resource_limits = ac_get_compute_resource_limits(&sscreen->info, waves_per_threadgroup, sctx->cs_max_waves_per_sh, threadgroups_per_cu); if (sctx->gfx_level >= GFX12) { gfx12_opt_push_compute_sh_reg(R_00B854_COMPUTE_RESOURCE_LIMITS, SI_TRACKED_COMPUTE_RESOURCE_LIMITS, compute_resource_limits); } else if (sctx->screen->info.has_set_sh_pairs_packed) { gfx11_opt_push_compute_sh_reg(R_00B854_COMPUTE_RESOURCE_LIMITS, SI_TRACKED_COMPUTE_RESOURCE_LIMITS, compute_resource_limits); } else { radeon_opt_set_sh_reg(R_00B854_COMPUTE_RESOURCE_LIMITS, SI_TRACKED_COMPUTE_RESOURCE_LIMITS, compute_resource_limits); } unsigned dispatch_initiator = S_00B800_COMPUTE_SHADER_EN(1) | S_00B800_FORCE_START_AT_000(1) | /* If the KMD allows it (there is a KMD hw register for it), * allow launching waves out-of-order. (same as Vulkan) * Not available in gfx940. */ S_00B800_ORDER_MODE(!sctx->cs_shader_state.program->sel.info.uses_atomic_ordered_add && sctx->gfx_level >= GFX7 && (sctx->family < CHIP_GFX940 || sctx->screen->info.has_graphics)) | S_00B800_CS_W32_EN(sctx->cs_shader_state.program->shader.wave_size == 32); const uint *last_block = info->last_block; bool partial_block_en = last_block[0] || last_block[1] || last_block[2]; uint32_t num_threads[3]; if (sctx->gfx_level >= GFX12) { num_threads[0] = S_00B81C_NUM_THREAD_FULL_GFX12(info->block[0]); num_threads[1] = S_00B820_NUM_THREAD_FULL_GFX12(info->block[1]); } else { num_threads[0] = S_00B81C_NUM_THREAD_FULL_GFX6(info->block[0]); num_threads[1] = S_00B820_NUM_THREAD_FULL_GFX6(info->block[1]); } num_threads[2] = S_00B824_NUM_THREAD_FULL(info->block[2]); if (partial_block_en) { unsigned partial[3]; /* If no partial_block, these should be an entire block size, not 0. */ partial[0] = last_block[0] ? last_block[0] : info->block[0]; partial[1] = last_block[1] ? last_block[1] : info->block[1]; partial[2] = last_block[2] ? last_block[2] : info->block[2]; num_threads[0] |= S_00B81C_NUM_THREAD_PARTIAL(partial[0]); num_threads[1] |= S_00B820_NUM_THREAD_PARTIAL(partial[1]); num_threads[2] |= S_00B824_NUM_THREAD_PARTIAL(partial[2]); dispatch_initiator |= S_00B800_PARTIAL_TG_EN(1); } if (sctx->gfx_level >= GFX12) { /* Set PING_PONG_EN for every other dispatch. * Only allowed on a gfx queue, and PARTIAL_TG_EN and USE_THREAD_DIMENSIONS must be 0. */ if (sctx->is_gfx_queue && !partial_block_en && !sctx->cs_shader_state.program->sel.info.uses_atomic_ordered_add) { dispatch_initiator |= S_00B800_PING_PONG_EN(sctx->compute_ping_pong_launch); sctx->compute_ping_pong_launch ^= 1; } /* Thread tiling within a workgroup. */ switch (sctx->cs_shader_state.program->shader.selector->info.base.derivative_group) { case DERIVATIVE_GROUP_LINEAR: break; case DERIVATIVE_GROUP_QUADS: num_threads[0] |= S_00B81C_INTERLEAVE_BITS_X(1); /* 2x2 */ num_threads[1] |= S_00B820_INTERLEAVE_BITS_Y(1); break; case DERIVATIVE_GROUP_NONE: /* These are the only legal combinations. */ if (info->block[0] % 8 == 0 && info->block[1] % 8 == 0) { num_threads[0] |= S_00B81C_INTERLEAVE_BITS_X(3); /* 8x8 */ num_threads[1] |= S_00B820_INTERLEAVE_BITS_Y(3); } else if (info->block[0] % 4 == 0 && info->block[1] % 8 == 0) { num_threads[0] |= S_00B81C_INTERLEAVE_BITS_X(2); /* 4x8 */ num_threads[1] |= S_00B820_INTERLEAVE_BITS_Y(3); } else if (info->block[0] % 4 == 0 && info->block[1] % 4 == 0) { num_threads[0] |= S_00B81C_INTERLEAVE_BITS_X(2); /* 4x4 */ num_threads[1] |= S_00B820_INTERLEAVE_BITS_Y(2); } else if (info->block[0] % 2 == 0 && info->block[1] % 2 == 0) { num_threads[0] |= S_00B81C_INTERLEAVE_BITS_X(1); /* 2x2 */ num_threads[1] |= S_00B820_INTERLEAVE_BITS_Y(1); } break; } /* How many threads should go to 1 SE before moving onto the next if INTERLEAVE_2D_EN == 0. * Only these values are valid: 0 (disabled), 64, 128, 256, 512 * 64 = RT, 256 = non-RT (run benchmarks to be sure) */ unsigned dispatch_interleave = S_00B8BC_INTERLEAVE_1D(256); unsigned log_x, log_y; /* Launch a 2D subgrid on each SE instead of a 1D subgrid. If enabled, INTERLEAVE_1D is * ignored and each SE gets 1 subgrid up to a certain number of threads. * * Constraints: * - Only supported by the gfx queue. * - Max 16 workgroups per SE can be launched, max 4 in each dimension. * - PARTIAL_TG_EN, USE_THREAD_DIMENSIONS, and ORDERED_APPEND_ENBL must be 0. * - COMPUTE_START_X/Y are in units of 2D subgrids, not workgroups * (program COMPUTE_START_X to start_x >> log_x, COMPUTE_START_Y to start_y >> log_y). */ if (sctx->is_gfx_queue && !partial_block_en && (info->indirect || info->grid[1] >= 4) && MIN2(info->block[0], info->block[1]) >= 4 && si_get_2d_interleave_size(info, &log_x, &log_y)) { dispatch_interleave = S_00B8BC_INTERLEAVE_1D(1) || /* 1D is disabled */ S_00B8BC_INTERLEAVE_2D_X_SIZE(log_x) | S_00B8BC_INTERLEAVE_2D_Y_SIZE(log_y); dispatch_initiator |= S_00B800_INTERLEAVE_2D_EN(1); } if (sctx->is_gfx_queue) { radeon_opt_set_sh_reg_idx(R_00B8BC_COMPUTE_DISPATCH_INTERLEAVE, SI_TRACKED_COMPUTE_DISPATCH_INTERLEAVE, 2, dispatch_interleave); } else { gfx12_opt_push_compute_sh_reg(R_00B8BC_COMPUTE_DISPATCH_INTERLEAVE, SI_TRACKED_COMPUTE_DISPATCH_INTERLEAVE, dispatch_interleave); } } if (sctx->gfx_level >= GFX12) { gfx12_opt_push_compute_sh_reg(R_00B81C_COMPUTE_NUM_THREAD_X, SI_TRACKED_COMPUTE_NUM_THREAD_X, num_threads[0]); gfx12_opt_push_compute_sh_reg(R_00B820_COMPUTE_NUM_THREAD_Y, SI_TRACKED_COMPUTE_NUM_THREAD_Y, num_threads[1]); gfx12_opt_push_compute_sh_reg(R_00B824_COMPUTE_NUM_THREAD_Z, SI_TRACKED_COMPUTE_NUM_THREAD_Z, num_threads[2]); } else if (sctx->screen->info.has_set_sh_pairs_packed) { gfx11_opt_push_compute_sh_reg(R_00B81C_COMPUTE_NUM_THREAD_X, SI_TRACKED_COMPUTE_NUM_THREAD_X, num_threads[0]); gfx11_opt_push_compute_sh_reg(R_00B820_COMPUTE_NUM_THREAD_Y, SI_TRACKED_COMPUTE_NUM_THREAD_Y, num_threads[1]); gfx11_opt_push_compute_sh_reg(R_00B824_COMPUTE_NUM_THREAD_Z, SI_TRACKED_COMPUTE_NUM_THREAD_Z, num_threads[2]); } else { radeon_opt_set_sh_reg3(R_00B81C_COMPUTE_NUM_THREAD_X, SI_TRACKED_COMPUTE_NUM_THREAD_X, num_threads[0], num_threads[1], num_threads[2]); } if (sctx->gfx_level >= GFX12 || sctx->screen->info.has_set_sh_pairs_packed) { radeon_end(); si_emit_buffered_compute_sh_regs(sctx); radeon_begin_again(cs); } if (info->indirect) { uint64_t base_va = si_resource(info->indirect)->gpu_address; radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, si_resource(info->indirect), RADEON_USAGE_READ | RADEON_PRIO_DRAW_INDIRECT); radeon_emit(PKT3(PKT3_SET_BASE, 2, 0) | PKT3_SHADER_TYPE_S(1)); radeon_emit(1); radeon_emit(base_va); radeon_emit(base_va >> 32); unsigned pkt = PKT3_DISPATCH_INDIRECT; if (sctx->gfx_level >= GFX12 && G_00B800_INTERLEAVE_2D_EN(dispatch_initiator)) pkt = PKT3_DISPATCH_INDIRECT_INTERLEAVED; radeon_emit(PKT3(pkt, 1, render_cond_bit) | PKT3_SHADER_TYPE_S(1)); radeon_emit(info->indirect_offset); radeon_emit(dispatch_initiator); } else { unsigned pkt = PKT3_DISPATCH_DIRECT; if (sctx->gfx_level >= GFX12 && G_00B800_INTERLEAVE_2D_EN(dispatch_initiator)) pkt = PKT3_DISPATCH_DIRECT_INTERLEAVED; radeon_emit(PKT3(pkt, 3, render_cond_bit) | PKT3_SHADER_TYPE_S(1)); radeon_emit(info->grid[0]); radeon_emit(info->grid[1]); radeon_emit(info->grid[2]); radeon_emit(dispatch_initiator); } if (unlikely(sctx->sqtt_enabled && sctx->gfx_level >= GFX9)) radeon_event_write(V_028A90_THREAD_TRACE_MARKER); radeon_end(); } static bool si_check_needs_implicit_sync(struct si_context *sctx, uint32_t usage) { /* If the compute shader is going to read from a texture/image written by a * previous draw, we must wait for its completion before continuing. * Buffers and image stores (from the draw) are not taken into consideration * because that's the app responsibility. * * The OpenGL 4.6 spec says: * * buffer object and texture stores performed by shaders are not * automatically synchronized * * TODO: Bindless textures are not handled, and thus are not synchronized. */ struct si_shader_info *info = &sctx->cs_shader_state.program->sel.info; struct si_samplers *samplers = &sctx->samplers[PIPE_SHADER_COMPUTE]; unsigned mask = samplers->enabled_mask & info->base.textures_used; while (mask) { int i = u_bit_scan(&mask); struct si_sampler_view *sview = (struct si_sampler_view *)samplers->views[i]; struct si_resource *res = si_resource(sview->base.texture); if (sctx->ws->cs_is_buffer_referenced(&sctx->gfx_cs, res->buf, usage)) return true; } struct si_images *images = &sctx->images[PIPE_SHADER_COMPUTE]; mask = BITFIELD_MASK(info->base.num_images) & images->enabled_mask; while (mask) { int i = u_bit_scan(&mask); struct pipe_image_view *sview = &images->views[i]; struct si_resource *res = si_resource(sview->resource); if (sctx->ws->cs_is_buffer_referenced(&sctx->gfx_cs, res->buf, usage)) return true; } return false; } static void si_launch_grid(struct pipe_context *ctx, const struct pipe_grid_info *info) { struct si_context *sctx = (struct si_context *)ctx; struct si_screen *sscreen = sctx->screen; struct si_compute *program = sctx->cs_shader_state.program; if (program->shader.compilation_failed) return; bool cs_regalloc_hang = sscreen->info.has_cs_regalloc_hang_bug && info->block[0] * info->block[1] * info->block[2] > 256; if (cs_regalloc_hang) { sctx->barrier_flags |= SI_BARRIER_SYNC_PS | SI_BARRIER_SYNC_CS; si_mark_atom_dirty(sctx, &sctx->atoms.s.barrier); } si_check_dirty_buffers_textures(sctx); if (sctx->is_gfx_queue) { if (sctx->num_draw_calls_sh_coherent.with_cb != sctx->num_draw_calls || sctx->num_draw_calls_sh_coherent.with_db != sctx->num_draw_calls) { bool sync_cb = sctx->force_shader_coherency.with_cb || si_check_needs_implicit_sync(sctx, RADEON_USAGE_CB_NEEDS_IMPLICIT_SYNC); bool sync_db = sctx->gfx_level >= GFX12 && (sctx->force_shader_coherency.with_db || si_check_needs_implicit_sync(sctx, RADEON_USAGE_DB_NEEDS_IMPLICIT_SYNC)); si_fb_barrier_after_rendering(sctx, (sync_cb ? SI_FB_BARRIER_SYNC_CB : 0) | (sync_db ? SI_FB_BARRIER_SYNC_DB : 0)); if (sync_cb) sctx->num_draw_calls_sh_coherent.with_cb = sctx->num_draw_calls; if (sync_db) sctx->num_draw_calls_sh_coherent.with_db = sctx->num_draw_calls; } if (sctx->gfx_level < GFX11) gfx6_decompress_textures(sctx, 1 << PIPE_SHADER_COMPUTE); else if (sctx->gfx_level < GFX12) gfx11_decompress_textures(sctx, 1 << PIPE_SHADER_COMPUTE); } if (info->indirect) { /* Indirect buffers are read through L2 on GFX9-GFX11, but not other hw. */ if ((sctx->gfx_level <= GFX8 || sscreen->info.cp_sdma_ge_use_system_memory_scope) && si_resource(info->indirect)->L2_cache_dirty) { sctx->barrier_flags |= SI_BARRIER_WB_L2 | SI_BARRIER_PFP_SYNC_ME; si_mark_atom_dirty(sctx, &sctx->atoms.s.barrier); si_resource(info->indirect)->L2_cache_dirty = false; } } si_need_gfx_cs_space(sctx, 0, 0); /* If we're using a secure context, determine if cs must be secure or not */ if (unlikely(radeon_uses_secure_bos(sctx->ws))) { bool secure = si_compute_resources_check_encrypted(sctx); if (secure != sctx->ws->cs_is_secure(&sctx->gfx_cs)) { si_flush_gfx_cs(sctx, RADEON_FLUSH_ASYNC_START_NEXT_GFX_IB_NOW | RADEON_FLUSH_TOGGLE_SECURE_SUBMISSION, NULL); } } if (unlikely(sctx->perfetto_enabled)) trace_si_begin_compute(&sctx->trace); if (sctx->bo_list_add_all_compute_resources) si_compute_resources_add_all_to_bo_list(sctx); /* Skipping setting redundant registers on compute queues breaks compute. */ if (!sctx->is_gfx_queue) { BITSET_CLEAR_RANGE(sctx->tracked_regs.reg_saved_mask, SI_FIRST_TRACKED_OTHER_REG, SI_NUM_ALL_TRACKED_REGS - 1); } /* First emit registers. */ bool prefetch; if (!si_switch_compute_shader(sctx, program, &program->shader, &prefetch, info->variable_shared_mem)) return; si_emit_compute_shader_pointers(sctx); /* Global buffers */ for (unsigned i = 0; i < sctx->max_global_buffers; i++) { struct si_resource *buffer = si_resource(sctx->global_buffers[i]); if (!buffer) { continue; } radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, buffer, RADEON_USAGE_READWRITE | RADEON_PRIO_SHADER_RW_BUFFER); } for (unsigned i = 0; i < info->num_globals; i++) { struct si_resource *buffer = si_resource(info->globals[i]); if (!buffer) { continue; } radeon_add_to_buffer_list(sctx, &sctx->gfx_cs, buffer, RADEON_USAGE_READWRITE | RADEON_PRIO_SHADER_RW_BUFFER); } /* Registers that are not read from memory should be set before this: */ si_emit_barrier_direct(sctx); if (sctx->is_gfx_queue && si_is_atom_dirty(sctx, &sctx->atoms.s.render_cond)) { sctx->atoms.s.render_cond.emit(sctx, -1); si_set_atom_dirty(sctx, &sctx->atoms.s.render_cond, false); } /* Prefetch the compute shader to L2. */ if (sctx->gfx_level >= GFX7 && sctx->screen->info.has_cp_dma && prefetch) si_cp_dma_prefetch(sctx, &program->shader.bo->b.b, 0, program->shader.bo->b.b.width0); si_setup_nir_user_data(sctx, info); si_emit_dispatch_packets(sctx, info); if (unlikely(sctx->current_saved_cs)) { si_trace_emit(sctx); si_log_compute_state(sctx, sctx->log); } if (sctx->gfx_level < GFX12) { /* Mark displayable DCC as dirty for bound images. */ unsigned display_dcc_store_mask = sctx->images[PIPE_SHADER_COMPUTE].display_dcc_store_mask & BITFIELD_MASK(program->sel.info.base.num_images); while (display_dcc_store_mask) { struct si_texture *tex = (struct si_texture *) sctx->images[PIPE_SHADER_COMPUTE].views[u_bit_scan(&display_dcc_store_mask)].resource; si_mark_display_dcc_dirty(sctx, tex); } /* TODO: Bindless images don't set displayable_dcc_dirty after image stores. */ } sctx->compute_is_busy = true; sctx->num_compute_calls++; if (unlikely(sctx->perfetto_enabled)) trace_si_end_compute(&sctx->trace, info->grid[0], info->grid[1], info->grid[2]); if (cs_regalloc_hang) { sctx->barrier_flags |= SI_BARRIER_SYNC_CS; si_mark_atom_dirty(sctx, &sctx->atoms.s.barrier); } } void si_destroy_compute(struct si_compute *program) { struct si_shader_selector *sel = &program->sel; util_queue_drop_job(&sel->screen->shader_compiler_queue, &sel->ready); util_queue_fence_destroy(&sel->ready); si_shader_destroy(&program->shader); ralloc_free(program->sel.nir); simple_mtx_destroy(&sel->mutex); FREE(program); } static void si_delete_compute_state(struct pipe_context *ctx, void *state) { struct si_compute *program = (struct si_compute *)state; struct si_context *sctx = (struct si_context *)ctx; if (!state) return; if (program == sctx->cs_shader_state.program) sctx->cs_shader_state.program = NULL; if (program == sctx->cs_shader_state.emitted_program) sctx->cs_shader_state.emitted_program = NULL; si_compute_reference(&program, NULL); } void si_init_compute_functions(struct si_context *sctx) { sctx->b.create_compute_state = si_create_compute_state; sctx->b.delete_compute_state = si_delete_compute_state; sctx->b.bind_compute_state = si_bind_compute_state; sctx->b.get_compute_state_info = si_get_compute_state_info; sctx->b.set_global_binding = si_set_global_binding; sctx->b.launch_grid = si_launch_grid; #if 0 /* test for si_get_2d_interleave_size */ static bool visited = false; if (visited) return; visited = true; struct pipe_grid_info info = {}; info.grid[0] = info.grid[1] = info.grid[2] = 1024; info.block[2] = 1; for (unsigned block_3d = 0; block_3d < 2; block_3d++) { printf(" WG size | WG block/SE | Thread block/SE\n"); for (unsigned x = 1; x <= 32; x *= 2) { for (unsigned y = 1; y <= 32; y *= 2) { info.block[0] = x; info.block[1] = y; info.block[2] = block_3d ? 8 : 1; if ((x * y) % 32) continue; unsigned log_x, log_y; if (!si_get_2d_interleave_size(&info, &log_x, &log_y)) continue; printf(" (%2u, %2u) = %3u | (%2u, %2u) = %2u | (%3u,%3u) = %u\n", info.block[0], info.block[1], info.block[0] * info.block[1], 1 << log_x, 1 << log_y, (1 << log_x) * (1 << log_y), info.block[0] * (1 << log_x), info.block[1] * (1 << log_y), info.block[0] * (1 << log_x) * info.block[1] * (1 << log_y)); } } } #endif }