From 361266fec72675eab7f6066dd755f9e16532bb49 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Tue, 16 Jul 2024 05:44:40 -0400 Subject: [PATCH] ac/nir: import the clear/copy_buffer compute shader from radeonsi Reviewed-by: Pierre-Eric Pelloux-Prayer Part-of: --- src/amd/common/ac_nir_meta.h | 60 +++ .../common/ac_nir_meta_cs_clear_copy_buffer.c | 490 ++++++++++++++++++ src/amd/common/meson.build | 1 + .../drivers/radeonsi/si_compute_blit.c | 178 ++----- src/gallium/drivers/radeonsi/si_pipe.h | 17 - .../drivers/radeonsi/si_shaderlib_nir.c | 317 ----------- 6 files changed, 588 insertions(+), 475 deletions(-) create mode 100644 src/amd/common/ac_nir_meta_cs_clear_copy_buffer.c diff --git a/src/amd/common/ac_nir_meta.h b/src/amd/common/ac_nir_meta.h index 2e6b562f500..c14932bbdfd 100644 --- a/src/amd/common/ac_nir_meta.h +++ b/src/amd/common/ac_nir_meta.h @@ -135,4 +135,64 @@ ac_prepare_compute_blit(const struct ac_cs_blit_options *options, const struct ac_cs_blit_description *blit, struct ac_cs_blit_dispatches *dispatches); +/* clear_buffer/copy_buffer compute shader. */ +union ac_cs_clear_copy_buffer_key { + struct { + bool is_clear:1; + unsigned dwords_per_thread:3; /* 1..4 allowed */ + bool clear_value_size_is_12:1; + bool src_is_sparse:1; + /* Unaligned clears and copies. */ + unsigned src_align_offset:2; /* how much is the source address unaligned */ + unsigned dst_align_offset:4; /* the first thread shouldn't write this many bytes */ + unsigned dst_last_thread_bytes:4; /* if non-zero, the last thread should write this many bytes */ + bool dst_single_thread_unaligned:1; /* only 1 thread executes, both previous fields apply */ + bool has_start_thread:1; /* whether the first few threads should be skipped, making later + waves start on a 256B boundary */ + }; + uint64_t key; +}; + +struct ac_cs_clear_copy_buffer_options { + const nir_shader_compiler_options *nir_options; + const struct radeon_info *info; + bool print_key; /* print the shader key into stderr */ + bool fail_if_slow; /* fail if a gfx blit is faster, set to false on compute queues */ +}; + +struct ac_cs_clear_copy_buffer_info { + unsigned dst_offset; + unsigned src_offset; + unsigned size; + unsigned clear_value_size; + uint32_t clear_value[4]; + unsigned dwords_per_thread; + bool render_condition_enabled; + bool dst_is_vram; + bool src_is_vram; + bool src_is_sparse; +}; + +struct ac_cs_clear_copy_buffer_dispatch { + union ac_cs_clear_copy_buffer_key shader_key; + uint32_t user_data[6]; /* for nir_intrinsic_load_user_data_amd */ + unsigned num_ssbos; + unsigned workgroup_size; + unsigned num_threads; + + struct { + unsigned offset; + unsigned size; + } ssbo[2]; +}; + +nir_shader * +ac_create_clear_copy_buffer_cs(struct ac_cs_clear_copy_buffer_options *options, + union ac_cs_clear_copy_buffer_key *key); + +bool +ac_prepare_cs_clear_copy_buffer(const struct ac_cs_clear_copy_buffer_options *options, + const struct ac_cs_clear_copy_buffer_info *info, + struct ac_cs_clear_copy_buffer_dispatch *out); + #endif diff --git a/src/amd/common/ac_nir_meta_cs_clear_copy_buffer.c b/src/amd/common/ac_nir_meta_cs_clear_copy_buffer.c new file mode 100644 index 00000000000..ca6e3786734 --- /dev/null +++ b/src/amd/common/ac_nir_meta_cs_clear_copy_buffer.c @@ -0,0 +1,490 @@ +/* + * Copyright 2024 Advanced Micro Devices, Inc. + * + * SPDX-License-Identifier: MIT + */ + +#include "ac_nir_meta.h" +#include "ac_nir_helpers.h" +#include "nir_builder.h" +#include "util/helpers.h" + +/* This is regular load_ssbo with special handling for sparse buffers. Normally, sparse buffer + * loads return 0 for all components if a sparse load starts on a non-resident page, crosses + * the page boundary, and ends on a resident page. For copy_buffer, we want it to return 0 only + * for the portion of the load that's non-resident, and load values for the portion that's + * resident. The workaround is to scalarize such loads and disallow vectorization. + */ +static nir_def * +load_ssbo_sparse(nir_builder *b, unsigned num_components, unsigned bit_size, nir_def *buf, + nir_def *offset, struct _nir_load_ssbo_indices params, bool sparse) +{ + if (sparse && num_components > 1) { + nir_def *vec[NIR_MAX_VEC_COMPONENTS]; + + /* Split the vector load into scalar loads. */ + for (unsigned i = 0; i < num_components; i++) { + unsigned elem_offset = i * bit_size / 8; + unsigned align_offset = (params.align_offset + elem_offset) % params.align_mul; + + vec[i] = nir_load_ssbo(b, 1, bit_size, buf, + nir_iadd_imm(b, offset, elem_offset), + .access = params.access | ACCESS_KEEP_SCALAR, + .align_mul = params.align_mul, + .align_offset = align_offset); + } + return nir_vec(b, vec, num_components); + } else { + return nir_load_ssbo(b, num_components, bit_size, buf, offset, + .access = params.access, + .align_mul = params.align_mul, + .align_offset = params.align_offset); + } +} + +/* Create a compute shader implementing clear_buffer or copy_buffer. */ +nir_shader * +ac_create_clear_copy_buffer_cs(struct ac_cs_clear_copy_buffer_options *options, + union ac_cs_clear_copy_buffer_key *key) +{ + if (options->print_key) { + fprintf(stderr, "Internal shader: dma\n"); + fprintf(stderr, " key.is_clear = %u\n", key->is_clear); + fprintf(stderr, " key.dwords_per_thread = %u\n", key->dwords_per_thread); + fprintf(stderr, " key.clear_value_size_is_12 = %u\n", key->clear_value_size_is_12); + fprintf(stderr, " key.src_is_sparse = %u\n", key->src_is_sparse); + fprintf(stderr, " key.src_align_offset = %u\n", key->src_align_offset); + fprintf(stderr, " key.dst_align_offset = %u\n", key->dst_align_offset); + fprintf(stderr, " key.dst_last_thread_bytes = %u\n", key->dst_last_thread_bytes); + fprintf(stderr, " key.dst_single_thread_unaligned = %u\n", key->dst_single_thread_unaligned); + fprintf(stderr, "\n"); + } + + assert(key->dwords_per_thread && key->dwords_per_thread <= 4); + + nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, options->nir_options, + "clear_copy_buffer_cs"); + b.shader->info.workgroup_size[0] = 64; + b.shader->info.workgroup_size[1] = 1; + b.shader->info.workgroup_size[2] = 1; + b.shader->info.num_ssbos = key->is_clear ? 1 : 2; + b.shader->info.cs.user_data_components_amd = 0; + + if (key->is_clear) { + b.shader->info.cs.user_data_components_amd += + key->clear_value_size_is_12 ? 3 : key->dwords_per_thread; + } + + /* Add the last thread ID value. */ + unsigned last_thread_user_data_index = b.shader->info.cs.user_data_components_amd; + if (key->dst_last_thread_bytes) + b.shader->info.cs.user_data_components_amd++; + + unsigned start_thread_user_data_index = b.shader->info.cs.user_data_components_amd; + if (key->has_start_thread) + b.shader->info.cs.user_data_components_amd++; + + nir_def *thread_id = ac_get_global_ids(&b, 1, 32); + + /* If the clear/copy area is unaligned, we launched extra threads at the beginning to make it + * aligned. Skip those threads here. + */ + nir_if *if_positive = NULL; + if (key->has_start_thread) { + nir_def *start_thread = + nir_channel(&b, nir_load_user_data_amd(&b), start_thread_user_data_index); + thread_id = nir_isub(&b, thread_id, start_thread); + if_positive = nir_push_if(&b, nir_ige_imm(&b, thread_id, 0)); + } + + /* Convert the global thread ID into bytes. */ + nir_def *offset = nir_imul_imm(&b, thread_id, 4 * key->dwords_per_thread); + nir_def *value; + + if (key->is_clear) { + value = nir_trim_vector(&b, nir_load_user_data_amd(&b), key->dwords_per_thread); + + /* We store 4 dwords per thread, but the clear value has 3 dwords. Swizzle it to 4 dwords. + * Storing 4 dwords per thread is faster even when the ALU cost is worse. + */ + if (key->clear_value_size_is_12 && key->dwords_per_thread == 4) { + nir_def *dw_offset = nir_imul_imm(&b, thread_id, key->dwords_per_thread); + nir_def *vec[3]; + + /* Swizzle a 3-component clear value to get a 4-component clear value. Example: + * 0 1 2 3 | 4 5 6 7 | 8 9 10 11 // dw_offset + * | + * V + * 0 1 2 0 | 1 2 0 1 | 2 0 1 2 // clear value component indices + */ + for (unsigned i = 0; i < 3; i++) { + vec[i] = nir_vector_extract(&b, value, + nir_umod_imm(&b, nir_iadd_imm(&b, dw_offset, i), 3)); + } + value = nir_vec4(&b, vec[0], vec[1], vec[2], vec[0]); + } + } else { + /* The hw doesn't support unaligned 32-bit loads, and only supports single-component + * unaligned 1-byte and 2-byte loads. Luckily, we don't have to use single-component loads + * because ac_nir_lower_subdword_load converts 1-byte and 2-byte vector loads with unaligned + * offsets into aligned 32-bit loads by loading an extra dword and then bit-shifting all bits + * to get the expected result. We only have to set bit_size to 8 or 16 and align_offset to + * 1..3 to indicate that this is an unaligned load. align_offset is the amount of + * unalignment. + * + * Since the buffer binding offsets are rounded down to the clear/copy size of the thread + * (i.e. dst_align_offset is subtracted from dst_offset, and src_align_offset is subtracted + * from src_offset), the stores expect the loaded value to be byte-shifted accordingly. + * realign_offset is the amount of byte-shifting we have to do. + */ + assert(util_is_power_of_two_nonzero(key->dwords_per_thread)); + int realign_offset = key->src_align_offset - key->dst_align_offset; + unsigned alignment = (unsigned)realign_offset % 4 == 0 ? 4 : + (unsigned)realign_offset % 2 == 0 ? 2 : 1; + unsigned bit_size = alignment * 8; + unsigned num_comps = key->dwords_per_thread * 4 / alignment; + nir_if *if_first_thread = NULL; + nir_def *value0 = NULL; + + if (realign_offset < 0) { + /* if src_align_offset is less than dst_align_offset, realign_offset is + * negative, which causes the first thread to use a negative buffer offset, which goes + * entirely out of bounds because the offset is treated as unsigned. Instead of that, + * the first thread should load from offset 0 by not loading the bytes before + * the beginning of the buffer. + */ + if_first_thread = nir_push_if(&b, nir_ieq_imm(&b, thread_id, 0)); + { + unsigned num_removed_comps = -realign_offset / alignment; + unsigned num_inbounds_comps = num_comps - num_removed_comps; + + /* Only 8 and 16 component vectors are valid after 5 in NIR. */ + while (!nir_num_components_valid(num_inbounds_comps)) + num_inbounds_comps = util_next_power_of_two(num_inbounds_comps); + + value0 = load_ssbo_sparse(&b, num_inbounds_comps, bit_size, nir_imm_int(&b, 0), offset, + (struct _nir_load_ssbo_indices){ + .access = ACCESS_RESTRICT, + .align_mul = 4, + .align_offset = 0 + }, key->src_is_sparse); + + /* Add the components that we didn't load as undef. */ + nir_def *comps[16]; + assert(num_comps <= ARRAY_SIZE(comps)); + for (unsigned i = 0; i < num_comps; i++) { + if (i < num_removed_comps) + comps[i] = nir_undef(&b, 1, bit_size); + else + comps[i] = nir_channel(&b, value0, i - num_removed_comps); + } + value0 = nir_vec(&b, comps, num_comps); + } + nir_push_else(&b, if_first_thread); + } + + value = load_ssbo_sparse(&b, num_comps, bit_size, nir_imm_int(&b, 0), + nir_iadd_imm(&b, offset, realign_offset), + (struct _nir_load_ssbo_indices){ + .access = ACCESS_RESTRICT, + .align_mul = 4, + .align_offset = (unsigned)realign_offset % 4 + }, key->src_is_sparse); + + + if (if_first_thread) { + nir_pop_if(&b, if_first_thread); + value = nir_if_phi(&b, value0, value); + } + + /* Bitcast the vector to 32 bits. */ + if (value->bit_size != 32) + value = nir_extract_bits(&b, &value, 1, 0, key->dwords_per_thread, 32); + } + + nir_def *dst_buf = nir_imm_int(&b, !key->is_clear); + nir_if *if_first_thread = NULL, *if_last_thread = NULL; + + if (!key->dst_single_thread_unaligned) { + /* dst_align_offset means how many bytes the first thread should skip because the offset of + * the buffer binding is rounded down to the clear/copy size of thread, causing the bytes + * before dst_align_offset to be writable. Above we used realign_offset to byte-shift + * the value to compensate for the rounded-down offset, so that all stores are dword stores + * regardless of the offset/size alignment except that the first thread shouldn't store + * the first dst_align_offset bytes, and the last thread should only store the first + * dst_last_thread_bytes. In both cases, there is a dword that must be only partially + * written by splitting it into 8-bit and 16-bit stores. + */ + if (key->dst_align_offset) { + if_first_thread = nir_push_if(&b, nir_ieq_imm(&b, thread_id, 0)); + { + unsigned local_offset = key->dst_align_offset; + nir_def *first_dword = nir_channel(&b, value, local_offset / 4); + + if (local_offset % 2 == 1) { + nir_store_ssbo(&b, nir_channel(&b, nir_unpack_32_4x8(&b, first_dword), local_offset % 4), + dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), + .access = ACCESS_RESTRICT); + local_offset++; + } + + if (local_offset % 4 == 2) { + nir_store_ssbo(&b, nir_unpack_32_2x16_split_y(&b, first_dword), dst_buf, + nir_iadd_imm_nuw(&b, offset, local_offset), + .access = ACCESS_RESTRICT); + local_offset += 2; + } + + assert(local_offset % 4 == 0); + unsigned num_dw_remaining = key->dwords_per_thread - local_offset / 4; + + if (num_dw_remaining) { + nir_def *dwords = + nir_channels(&b, value, BITFIELD_RANGE(local_offset / 4, num_dw_remaining)); + + nir_store_ssbo(&b, dwords, dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), + .access = ACCESS_RESTRICT); + } + } + nir_push_else(&b, if_first_thread); + } + + if (key->dst_last_thread_bytes) { + nir_def *last_thread_id = + nir_channel(&b, nir_load_user_data_amd(&b), last_thread_user_data_index); + + if_last_thread = nir_push_if(&b, nir_ieq(&b, thread_id, last_thread_id)); + { + unsigned num_dwords = key->dst_last_thread_bytes / 4; + bool write_short = (key->dst_last_thread_bytes - num_dwords * 4) / 2; + bool write_byte = key->dst_last_thread_bytes % 2; + nir_def *last_dword = nir_channel(&b, value, num_dwords); + + if (num_dwords) { + nir_def *dwords = nir_channels(&b, value, BITFIELD_MASK(num_dwords)); + nir_store_ssbo(&b, dwords, dst_buf, offset, .access = ACCESS_RESTRICT); + } + + if (write_short) { + nir_store_ssbo(&b, nir_u2u16(&b, last_dword), dst_buf, + nir_iadd_imm_nuw(&b, offset, num_dwords * 4), + .access = ACCESS_RESTRICT); + } + + if (write_byte) { + nir_store_ssbo(&b, nir_channel(&b, nir_unpack_32_4x8(&b, last_dword), write_short * 2), + dst_buf, nir_iadd_imm_nuw(&b, offset, num_dwords * 4 + write_short * 2), + .access = ACCESS_RESTRICT); + } + } + nir_push_else(&b, if_last_thread); + } + + nir_store_ssbo(&b, value, dst_buf, offset, .access = ACCESS_RESTRICT); + + if (if_last_thread) + nir_pop_if(&b, if_last_thread); + if (if_first_thread) + nir_pop_if(&b, if_first_thread); + } else { + /* This shader only executes a single thread (tiny copy or clear) and it's unaligned at both + * the beginning and the end. Walk the individual dwords/words/bytes that should be written + * to split the store accordingly. + */ + for (unsigned local_offset = key->dst_align_offset; + local_offset < key->dst_last_thread_bytes;) { + unsigned remaining = key->dst_last_thread_bytes - local_offset; + nir_def *src_dword = nir_channel(&b, value, local_offset / 4); + + if (local_offset % 2 == 1 || remaining == 1) { + /* 1-byte store. */ + nir_def *src_dword4x8 = nir_unpack_32_4x8(&b, src_dword); + nir_store_ssbo(&b, nir_channel(&b, src_dword4x8, local_offset % 4), dst_buf, + nir_iadd_imm_nuw(&b, offset, local_offset), .access = ACCESS_RESTRICT); + local_offset++; + } else if (local_offset % 4 == 2 || remaining == 2 || remaining == 3) { + /* 2-byte store. */ + nir_def *src_dword2x16 = nir_unpack_32_2x16(&b, src_dword); + nir_store_ssbo(&b, nir_channel(&b, src_dword2x16, (local_offset / 2) % 2), dst_buf, + nir_iadd_imm_nuw(&b, offset, local_offset), .access = ACCESS_RESTRICT); + local_offset += 2; + } else { + /* 1-N dwords. */ + unsigned dw_size = remaining / 4; + assert(dw_size); + assert(local_offset % 4 == 0); + + nir_store_ssbo(&b, nir_channels(&b, value, BITFIELD_RANGE(local_offset / 4, dw_size)), + dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), + .access = ACCESS_RESTRICT); + local_offset += dw_size * 4; + } + } + } + + if (key->has_start_thread) + nir_pop_if(&b, if_positive); + + return b.shader; +} + +bool +ac_prepare_cs_clear_copy_buffer(const struct ac_cs_clear_copy_buffer_options *options, + const struct ac_cs_clear_copy_buffer_info *info, + struct ac_cs_clear_copy_buffer_dispatch *out) +{ + bool is_copy = info->clear_value_size == 0; + + memset(out, 0, sizeof(*out)); + + /* Expand 1-byte and 2-byte clear values to a dword. */ + int clear_value_size = info->clear_value_size; + const uint32_t *clear_value = info->clear_value; + uint32_t tmp_clear_value; + + if (!is_copy) { + if (util_lower_clearsize_to_dword(clear_value, &clear_value_size, &tmp_clear_value)) + clear_value = &tmp_clear_value; + + assert(clear_value_size % 4 == 0); + } + + unsigned dwords_per_thread = info->dwords_per_thread; + if (!dwords_per_thread) { + /* Set default optimal settings. */ + dwords_per_thread = info->size <= 64 * 1024 ? 2 : 4; + + if (!is_copy) { + if (clear_value_size == 12) { + /* Clearing 4 dwords per thread with a 3-dword clear value is faster with big sizes. */ + dwords_per_thread = info->size <= 4096 ? 3 : 4; + } else { + /* dwords_per_thread must be at least the size of the clear value. */ + dwords_per_thread = MAX2(dwords_per_thread, clear_value_size / 4); + } + } + } + + /* Validate dwords_per_thread. */ + if (dwords_per_thread > 4) { + assert(!"dwords_per_thread must be <= 4"); + return false; /* invalid value */ + } + + if (clear_value_size > dwords_per_thread * 4) { + assert(!"clear_value_size must be <= dwords_per_thread"); + return false; /* invalid value */ + } + + if (clear_value_size == 12 && info->dst_offset % 4) { + assert(!"if clear_value_size == 12, dst_offset must be aligned to 4"); + return false; /* invalid value */ + } + + /* This doesn't fail very often because the only possible fallback is CP DMA, which doesn't + * support the render condition. + */ + if (options->fail_if_slow && !info->render_condition_enabled && options->info->has_cp_dma && + !options->info->cp_sdma_ge_use_system_memory_scope) { + if (is_copy) { + /* Only use compute for large VRAM copies on dGPUs. */ + if (info->size <= 8192 || !options->info->has_dedicated_vram || !info->dst_is_vram || + !info->src_is_vram) + return false; + } else { + /* Buffer clear. + * + * CP DMA clears are terribly slow with GTT on GFX6-8, which can be encountered with any + * buffer due to BO evictions, so never use CP DMA clears on GFX6-8. On GFX9+, use CP DMA + * clears if the size is small. + */ + if (options->info->gfx_level >= GFX9 && clear_value_size <= 4 && info->size <= 4096) + return false; + } + } + + unsigned dst_align_offset = info->dst_offset % (dwords_per_thread * 4); + unsigned dst_offset_bound = info->dst_offset - dst_align_offset; + unsigned src_align_offset = is_copy ? info->src_offset % 4 : 0; + unsigned num_user_data_terms = 0; + + /* Set the clear value in user data SGPRs. */ + if (!is_copy) { + assert(clear_value_size >= 4 && clear_value_size <= 16 && + (clear_value_size == 12 || util_is_power_of_two_or_zero(clear_value_size))); + + /* Since the clear value may start on an unaligned offset and we just pass user SGPRs + * to dword stores as-is, we need to byte-shift the clear value to that offset and + * replicate it because 1 invocation stores up to 4 dwords from user SGPRs regardless of + * the clear value size. + */ + num_user_data_terms = clear_value_size == 12 ? 3 : dwords_per_thread; + unsigned user_data_size = num_user_data_terms * 4; + + memcpy(out->user_data, + (uint8_t*)clear_value + clear_value_size - dst_align_offset % clear_value_size, + dst_align_offset % clear_value_size); + unsigned offset = dst_align_offset % clear_value_size; + + while (offset + clear_value_size <= user_data_size) { + memcpy((uint8_t*)out->user_data + offset, clear_value, clear_value_size); + offset += clear_value_size; + } + + if (offset < user_data_size) + memcpy((uint8_t*)out->user_data + offset, clear_value, user_data_size - offset); + } + + out->shader_key.key = 0; + + out->shader_key.is_clear = !is_copy; + assert(dwords_per_thread && dwords_per_thread <= 4); + out->shader_key.dwords_per_thread = dwords_per_thread; + out->shader_key.clear_value_size_is_12 = !is_copy && clear_value_size == 12; + out->shader_key.src_is_sparse = info->src_is_sparse; + out->shader_key.src_align_offset = src_align_offset; + out->shader_key.dst_align_offset = dst_align_offset; + + if ((dst_align_offset + info->size) % 4) + out->shader_key.dst_last_thread_bytes = (dst_align_offset + info->size) % (dwords_per_thread * 4); + + unsigned num_threads = DIV_ROUND_UP(dst_align_offset + info->size, dwords_per_thread * 4); + out->shader_key.dst_single_thread_unaligned = num_threads == 1 && dst_align_offset && + out->shader_key.dst_last_thread_bytes; + + /* start_thread offsets threads to make sure all non-zero waves start clearing/copying from + * the beginning a 256B block and clear/copy whole 256B blocks. Clearing/copying a 256B block + * partially for each wave is inefficient, which happens when dst_offset isn't aligned to 256. + * Clearing/copying whole 256B blocks per wave isn't possible if dwords_per_thread isn't 2^n. + */ + unsigned start_thread = + dst_offset_bound % 256 && util_is_power_of_two_nonzero(dwords_per_thread) ? + DIV_ROUND_UP(256 - dst_offset_bound % 256, dwords_per_thread * 4) : 0; + out->shader_key.has_start_thread = start_thread != 0; + + /* Set the value of the last thread ID, so that the shader knows which thread is the last one. */ + if (out->shader_key.dst_last_thread_bytes) + out->user_data[num_user_data_terms++] = num_threads - 1; + if (out->shader_key.has_start_thread) + out->user_data[num_user_data_terms++] = start_thread; + + /* We need to bind whole dwords because of how we compute voffset. The bytes that shouldn't + * be written are not written by the shader. + */ + out->ssbo[is_copy].offset = dst_offset_bound; + out->ssbo[is_copy].size = align(dst_align_offset + info->size, 4); + + if (is_copy) { + /* Since unaligned copies use 32-bit loads, any dword that's partially covered by the copy + * range must be fully covered, so that the 32-bit loads succeed. + */ + out->ssbo[0].offset = info->src_offset - src_align_offset; + out->ssbo[0].size = align(src_align_offset + info->size, 4); + assert(out->ssbo[0].offset % 4 == 0 && out->ssbo[0].size % 4 == 0); + } + + out->num_ssbos = is_copy ? 2 : 1; + out->workgroup_size = 64; + out->num_threads = start_thread + num_threads; + return true; +} diff --git a/src/amd/common/meson.build b/src/amd/common/meson.build index bb9934e0a77..74329ffa5dd 100644 --- a/src/amd/common/meson.build +++ b/src/amd/common/meson.build @@ -98,6 +98,7 @@ amd_common_files = files( 'ac_nir_lower_ps.c', 'ac_nir_meta.h', 'ac_nir_meta_cs_blit.c', + 'ac_nir_meta_cs_clear_copy_buffer.c', 'ac_nir_meta_ps_resolve.c', 'amd_family.c', 'ac_parse_ib.c', diff --git a/src/gallium/drivers/radeonsi/si_compute_blit.c b/src/gallium/drivers/radeonsi/si_compute_blit.c index e1dd7ea427e..90339c56671 100644 --- a/src/gallium/drivers/radeonsi/si_compute_blit.c +++ b/src/gallium/drivers/radeonsi/si_compute_blit.c @@ -274,164 +274,60 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource unsigned flags, enum si_coherency coher, unsigned dwords_per_thread, bool fail_if_slow) { + assert(dst->target != PIPE_BUFFER || dst_offset + size <= dst->width0); + assert(!src || src_offset + size <= src->width0); bool is_copy = src != NULL; - uint32_t tmp_clear_value; si_improve_sync_flags(sctx, dst, src, &flags); - if (!is_copy) { - if (util_lower_clearsize_to_dword(clear_value, (int*)&clear_value_size, &tmp_clear_value)) - clear_value = &tmp_clear_value; + struct ac_cs_clear_copy_buffer_options options = { + .nir_options = sctx->screen->nir_options, + .info = &sctx->screen->info, + .print_key = si_can_dump_shader(sctx->screen, MESA_SHADER_COMPUTE, SI_DUMP_SHADER_KEY), + .fail_if_slow = fail_if_slow, + }; - assert(clear_value_size % 4 == 0); - } + struct ac_cs_clear_copy_buffer_info info = { + .dst_offset = dst_offset, + .src_offset = src_offset, + .size = size, + .clear_value_size = is_copy ? 0 : clear_value_size, + .dwords_per_thread = dwords_per_thread, + .render_condition_enabled = flags & SI_OP_CS_RENDER_COND_ENABLE, + .dst_is_vram = si_resource(dst)->domains & RADEON_DOMAIN_VRAM, + .src_is_vram = src && si_resource(src)->domains & RADEON_DOMAIN_VRAM, + .src_is_sparse = src && src->flags & PIPE_RESOURCE_FLAG_SPARSE, + }; + memcpy(info.clear_value, clear_value, clear_value_size); - if (!dwords_per_thread) { - /* Set default optimal settings. */ - dwords_per_thread = size <= 64 * 1024 ? 2 : 4; + struct ac_cs_clear_copy_buffer_dispatch dispatch; - if (!is_copy) { - if (clear_value_size == 12) { - /* Clearing 4 dwords per thread with a 3-dword clear value is faster with big sizes. */ - dwords_per_thread = size <= 4096 ? 3 : 4; - } else { - /* dwords_per_thread must be at least the size of the clear value. */ - dwords_per_thread = MAX2(dwords_per_thread, clear_value_size / 4); - } - } - } - - /* Validate dwords_per_thread. */ - if (dwords_per_thread > 4) { - assert(!"dwords_per_thread must be <= 4"); - return false; /* invalid value */ - } - - if (clear_value_size > dwords_per_thread * 4) { - assert(!"clear_value_size must be <= dwords_per_thread"); - return false; /* invalid value */ - } - - if (clear_value_size == 12 && dst_offset % 4) { - assert(!"if clear_value_size == 12, dst_offset must be aligned to 4"); - return false; /* invalid value */ - } - - /* This doesn't fail very often because the only possible fallback is CP DMA, which doesn't - * support the render condition. - */ - if (fail_if_slow && !(flags & SI_OP_CS_RENDER_COND_ENABLE) && sctx->screen->info.has_cp_dma && - !sctx->screen->info.cp_sdma_ge_use_system_memory_scope) { - if (is_copy) { - /* Only use compute for large VRAM copies on dGPUs. */ - if (size <= 8192 || !sctx->screen->info.has_dedicated_vram || - !(si_resource(dst)->domains & RADEON_DOMAIN_VRAM) || - !(si_resource(src)->domains & RADEON_DOMAIN_VRAM)) - return false; - } else { - /* Buffer clear. - * - * CP DMA clears are terribly slow with GTT on GFX6-8, which can be encountered with any - * buffer due to BO evictions, so never use CP DMA clears on GFX6-8. On GFX9+, use CP DMA - * clears if the size is small. - */ - if (sctx->gfx_level >= GFX9 && clear_value_size <= 4 && size <= 4096) - return false; - } - } - - assert(dst->target != PIPE_BUFFER || dst_offset + size <= dst->width0); - assert(!src || src_offset + size <= src->width0); - unsigned dst_align_offset = dst_offset % (dwords_per_thread * 4); - unsigned dst_offset_bound = dst_offset - dst_align_offset; - unsigned src_align_offset = is_copy ? src_offset % 4 : 0; - unsigned num_user_data_terms = 0; + if (!ac_prepare_cs_clear_copy_buffer(&options, &info, &dispatch)) + return false; struct pipe_shader_buffer sb[2] = {}; - sb[is_copy].buffer = dst; - /* We need to bind whole dwords because of how we compute voffset. The bytes that shouldn't - * be written are not written by the shader. - */ - sb[is_copy].buffer_offset = dst_offset_bound; - sb[is_copy].buffer_size = align(dst_align_offset + size, 4); + for (unsigned i = 0; i < 2; i++) { + sb[i].buffer_offset = dispatch.ssbo[i].offset; + sb[i].buffer_size = dispatch.ssbo[i].size; + } - if (is_copy) { + if (is_copy) sb[0].buffer = src; - /* Since unaligned copies use 32-bit loads, any dword that's partially covered by the copy - * range must be fully covered, so that the 32-bit loads succeed. - */ - sb[0].buffer_offset = src_offset - src_align_offset; - sb[0].buffer_size = align(src_align_offset + size, 4); - assert(sb[0].buffer_offset % 4 == 0 && sb[0].buffer_size % 4 == 0); - } else { - assert(clear_value_size >= 4 && clear_value_size <= 16 && - (clear_value_size == 12 || util_is_power_of_two_or_zero(clear_value_size))); + sb[is_copy].buffer = dst; - /* Since the clear value may start on an unaligned offset and we just pass user SGPRs - * to dword stores as-is, we need to byte-shift the clear value to that offset and - * replicate it because 1 invocation stores up to 4 dwords from user SGPRs regardless of - * the clear value size. - */ - num_user_data_terms = clear_value_size == 12 ? 3 : dwords_per_thread; - unsigned user_data_size = num_user_data_terms * 4; - - memcpy(sctx->cs_user_data, - (uint8_t*)clear_value + clear_value_size - dst_align_offset % clear_value_size, - dst_align_offset % clear_value_size); - unsigned offset = dst_align_offset % clear_value_size; - - while (offset + clear_value_size <= user_data_size) { - memcpy((uint8_t*)sctx->cs_user_data + offset, clear_value, clear_value_size); - offset += clear_value_size; - } - - if (offset < user_data_size) - memcpy((uint8_t*)sctx->cs_user_data + offset, clear_value, user_data_size - offset); - } - - union si_cs_clear_copy_buffer_key key; - key.key = 0; - - key.is_clear = !is_copy; - assert(dwords_per_thread && dwords_per_thread <= 4); - key.dwords_per_thread = dwords_per_thread; - key.clear_value_size_is_12 = !is_copy && clear_value_size == 12; - key.src_is_sparse = src ? src->flags & PIPE_RESOURCE_FLAG_SPARSE : false; - key.src_align_offset = src_align_offset; - key.dst_align_offset = dst_align_offset; - - if ((dst_align_offset + size) % 4) - key.dst_last_thread_bytes = (dst_align_offset + size) % (dwords_per_thread * 4); - - unsigned num_threads = DIV_ROUND_UP(dst_align_offset + size, dwords_per_thread * 4); - key.dst_single_thread_unaligned = num_threads == 1 && dst_align_offset && key.dst_last_thread_bytes; - - /* start_thread offsets threads to make sure all non-zero waves start clearing/copying from - * the beginning a 256B block and clear/copy whole 256B blocks. Clearing/copying a 256B block - * partially for each wave is inefficient, which happens when dst_offset isn't aligned to 256. - * Clearing/copying whole 256B blocks per wave isn't possible if dwords_per_thread isn't 2^n. - */ - unsigned start_thread = - dst_offset_bound % 256 && util_is_power_of_two_nonzero(dwords_per_thread) ? - DIV_ROUND_UP(256 - dst_offset_bound % 256, dwords_per_thread * 4) : 0; - key.has_start_thread = start_thread != 0; - - void *shader = _mesa_hash_table_u64_search(sctx->cs_dma_shaders, key.key); + void *shader = _mesa_hash_table_u64_search(sctx->cs_dma_shaders, dispatch.shader_key.key); if (!shader) { - shader = si_create_dma_compute_shader(sctx, &key); - _mesa_hash_table_u64_insert(sctx->cs_dma_shaders, key.key, shader); + shader = si_create_shader_state(sctx, ac_create_clear_copy_buffer_cs(&options, + &dispatch.shader_key)); + _mesa_hash_table_u64_insert(sctx->cs_dma_shaders, dispatch.shader_key.key, shader); } - /* Set the value of the last thread ID, so that the shader knows which thread is the last one. */ - if (key.dst_last_thread_bytes) - sctx->cs_user_data[num_user_data_terms++] = num_threads - 1; - if (key.has_start_thread) - sctx->cs_user_data[num_user_data_terms++] = start_thread; + memcpy(sctx->cs_user_data, dispatch.user_data, sizeof(dispatch.user_data)); - struct pipe_grid_info info = {}; - set_work_size(&info, 64, 1, 1, start_thread + num_threads, 1, 1); + struct pipe_grid_info grid = {}; + set_work_size(&grid, dispatch.workgroup_size, 1, 1, dispatch.num_threads, 1, 1); - si_launch_grid_internal_ssbos(sctx, &info, shader, flags, coher, is_copy ? 2 : 1, sb, + si_launch_grid_internal_ssbos(sctx, &grid, shader, flags, coher, dispatch.num_ssbos, sb, is_copy ? 0x2 : 0x1); return true; } diff --git a/src/gallium/drivers/radeonsi/si_pipe.h b/src/gallium/drivers/radeonsi/si_pipe.h index 1898fb3c6db..a8d5fe80f82 100644 --- a/src/gallium/drivers/radeonsi/si_pipe.h +++ b/src/gallium/drivers/radeonsi/si_pipe.h @@ -1644,22 +1644,6 @@ void si_suspend_queries(struct si_context *sctx); void si_resume_queries(struct si_context *sctx); /* si_shaderlib_nir.c */ -union si_cs_clear_copy_buffer_key { - struct { - bool is_clear:1; - unsigned dwords_per_thread:3; /* 1..4 allowed */ - bool clear_value_size_is_12:1; - bool src_is_sparse:1; - /* Unaligned clears and copies. */ - unsigned src_align_offset:2; /* how much is the source address unaligned */ - unsigned dst_align_offset:4; /* the first thread shouldn't write this many bytes */ - unsigned dst_last_thread_bytes:4; /* if non-zero, the last thread should write this many bytes */ - bool dst_single_thread_unaligned:1; /* only 1 thread executes, both previous fields apply */ - bool has_start_thread:1; /* whether the first few threads should be skipped, making later - waves start on a 256B boundary */ - }; - uint64_t key; -}; void *si_create_shader_state(struct si_context *sctx, struct nir_shader *nir); void *si_create_dcc_retile_cs(struct si_context *sctx, struct radeon_surf *surf); @@ -1668,7 +1652,6 @@ void *si_create_passthrough_tcs(struct si_context *sctx); void *si_clear_image_dcc_single_shader(struct si_context *sctx, bool is_msaa, unsigned wg_dim); void *si_get_blitter_vs(struct si_context *sctx, enum blitter_attrib_type type, unsigned num_layers); -void *si_create_dma_compute_shader(struct si_context *sctx, union si_cs_clear_copy_buffer_key *key); void *si_create_ubyte_to_ushort_compute_shader(struct si_context *sctx); void *si_create_clear_buffer_rmw_cs(struct si_context *sctx); void *si_create_fmask_expand_cs(struct si_context *sctx, unsigned num_samples, bool is_array); diff --git a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c index 3296927039a..50803254ec1 100644 --- a/src/gallium/drivers/radeonsi/si_shaderlib_nir.c +++ b/src/gallium/drivers/radeonsi/si_shaderlib_nir.c @@ -227,323 +227,6 @@ void *si_create_ubyte_to_ushort_compute_shader(struct si_context *sctx) return si_create_shader_state(sctx, b.shader); } -/* This is regular load_ssbo with special handling for sparse buffers. Normally, sparse buffer - * loads return 0 for all components if a sparse load starts on a non-resident page, crosses - * the page boundary, and ends on a resident page. For copy_buffer, we want it to return 0 only - * for the portion of the load that's non-resident, and load values for the portion that's - * resident. The workaround is to scalarize such loads and disallow vectorization. - */ -static nir_def * -load_ssbo_sparse(nir_builder *b, unsigned num_components, unsigned bit_size, nir_def *buf, - nir_def *offset, struct _nir_load_ssbo_indices params, bool sparse) -{ - if (sparse && num_components > 1) { - nir_def *vec[NIR_MAX_VEC_COMPONENTS]; - - /* Split the vector load into scalar loads. */ - for (unsigned i = 0; i < num_components; i++) { - unsigned elem_offset = i * bit_size / 8; - unsigned align_offset = (params.align_offset + elem_offset) % params.align_mul; - - vec[i] = nir_load_ssbo(b, 1, bit_size, buf, - nir_iadd_imm(b, offset, elem_offset), - .access = params.access | ACCESS_KEEP_SCALAR, - .align_mul = params.align_mul, - .align_offset = align_offset); - } - return nir_vec(b, vec, num_components); - } else { - return nir_load_ssbo(b, num_components, bit_size, buf, offset, - .access = params.access, - .align_mul = params.align_mul, - .align_offset = params.align_offset); - } -} - -/* Create a compute shader implementing clear_buffer or copy_buffer. */ -void *si_create_dma_compute_shader(struct si_context *sctx, union si_cs_clear_copy_buffer_key *key) -{ - if (si_can_dump_shader(sctx->screen, MESA_SHADER_COMPUTE, SI_DUMP_SHADER_KEY)) { - fprintf(stderr, "Internal shader: dma\n"); - fprintf(stderr, " key.is_clear = %u\n", key->is_clear); - fprintf(stderr, " key.dwords_per_thread = %u\n", key->dwords_per_thread); - fprintf(stderr, " key.clear_value_size_is_12 = %u\n", key->clear_value_size_is_12); - fprintf(stderr, " key.src_is_sparse = %u\n", key->src_is_sparse); - fprintf(stderr, " key.src_align_offset = %u\n", key->src_align_offset); - fprintf(stderr, " key.dst_align_offset = %u\n", key->dst_align_offset); - fprintf(stderr, " key.dst_last_thread_bytes = %u\n", key->dst_last_thread_bytes); - fprintf(stderr, " key.dst_single_thread_unaligned = %u\n", key->dst_single_thread_unaligned); - fprintf(stderr, "\n"); - } - - assert(key->dwords_per_thread && key->dwords_per_thread <= 4); - - nir_builder b = nir_builder_init_simple_shader(MESA_SHADER_COMPUTE, sctx->screen->nir_options, - "create_dma_compute"); - b.shader->info.workgroup_size[0] = 64; - b.shader->info.workgroup_size[1] = 1; - b.shader->info.workgroup_size[2] = 1; - b.shader->info.num_ssbos = key->is_clear ? 1 : 2; - b.shader->info.cs.user_data_components_amd = 0; - - if (key->is_clear) { - b.shader->info.cs.user_data_components_amd += - key->clear_value_size_is_12 ? 3 : key->dwords_per_thread; - } - - /* Add the last thread ID value. */ - unsigned last_thread_user_data_index = b.shader->info.cs.user_data_components_amd; - if (key->dst_last_thread_bytes) - b.shader->info.cs.user_data_components_amd++; - - unsigned start_thread_user_data_index = b.shader->info.cs.user_data_components_amd; - if (key->has_start_thread) - b.shader->info.cs.user_data_components_amd++; - - nir_def *thread_id = ac_get_global_ids(&b, 1, 32); - - /* If the clear/copy area is unaligned, we launched extra threads at the beginning to make it - * aligned. Skip those threads here. - */ - nir_if *if_positive = NULL; - if (key->has_start_thread) { - nir_def *start_thread = - nir_channel(&b, nir_load_user_data_amd(&b), start_thread_user_data_index); - thread_id = nir_isub(&b, thread_id, start_thread); - if_positive = nir_push_if(&b, nir_ige_imm(&b, thread_id, 0)); - } - - /* Convert the global thread ID into bytes. */ - nir_def *offset = nir_imul_imm(&b, thread_id, 4 * key->dwords_per_thread); - nir_def *value; - - if (key->is_clear) { - value = nir_trim_vector(&b, nir_load_user_data_amd(&b), key->dwords_per_thread); - - /* We store 4 dwords per thread, but the clear value has 3 dwords. Swizzle it to 4 dwords. - * Storing 4 dwords per thread is faster even when the ALU cost is worse. - */ - if (key->clear_value_size_is_12 && key->dwords_per_thread == 4) { - nir_def *dw_offset = nir_imul_imm(&b, thread_id, key->dwords_per_thread); - nir_def *vec[3]; - - /* Swizzle a 3-component clear value to get a 4-component clear value. Example: - * 0 1 2 3 | 4 5 6 7 | 8 9 10 11 // dw_offset - * | - * V - * 0 1 2 0 | 1 2 0 1 | 2 0 1 2 // clear value component indices - */ - for (unsigned i = 0; i < 3; i++) { - vec[i] = nir_vector_extract(&b, value, - nir_umod_imm(&b, nir_iadd_imm(&b, dw_offset, i), 3)); - } - value = nir_vec4(&b, vec[0], vec[1], vec[2], vec[0]); - } - } else { - /* The hw doesn't support unaligned 32-bit loads, and only supports single-component - * unaligned 1-byte and 2-byte loads. Luckily, we don't have to use single-component loads - * because ac_nir_lower_subdword_load converts 1-byte and 2-byte vector loads with unaligned - * offsets into aligned 32-bit loads by loading an extra dword and then bit-shifting all bits - * to get the expected result. We only have to set bit_size to 8 or 16 and align_offset to - * 1..3 to indicate that this is an unaligned load. align_offset is the amount of - * unalignment. - * - * Since the buffer binding offsets are rounded down to the clear/copy size of the thread - * (i.e. dst_align_offset is subtracted from dst_offset, and src_align_offset is subtracted - * from src_offset), the stores expect the loaded value to be byte-shifted accordingly. - * realign_offset is the amount of byte-shifting we have to do. - */ - assert(util_is_power_of_two_nonzero(key->dwords_per_thread)); - int realign_offset = key->src_align_offset - key->dst_align_offset; - unsigned alignment = (unsigned)realign_offset % 4 == 0 ? 4 : - (unsigned)realign_offset % 2 == 0 ? 2 : 1; - unsigned bit_size = alignment * 8; - unsigned num_comps = key->dwords_per_thread * 4 / alignment; - nir_if *if_first_thread = NULL; - nir_def *value0 = NULL; - - if (realign_offset < 0) { - /* if src_align_offset is less than dst_align_offset, realign_offset is - * negative, which causes the first thread to use a negative buffer offset, which goes - * entirely out of bounds because the offset is treated as unsigned. Instead of that, - * the first thread should load from offset 0 by not loading the bytes before - * the beginning of the buffer. - */ - if_first_thread = nir_push_if(&b, nir_ieq_imm(&b, thread_id, 0)); - { - unsigned num_removed_comps = -realign_offset / alignment; - unsigned num_inbounds_comps = num_comps - num_removed_comps; - - /* Only 8 and 16 component vectors are valid after 5 in NIR. */ - while (!nir_num_components_valid(num_inbounds_comps)) - num_inbounds_comps = util_next_power_of_two(num_inbounds_comps); - - value0 = load_ssbo_sparse(&b, num_inbounds_comps, bit_size, nir_imm_int(&b, 0), offset, - (struct _nir_load_ssbo_indices){ - .access = ACCESS_RESTRICT, - .align_mul = 4, - .align_offset = 0 - }, key->src_is_sparse); - - /* Add the components that we didn't load as undef. */ - nir_def *comps[16]; - assert(num_comps <= ARRAY_SIZE(comps)); - for (unsigned i = 0; i < num_comps; i++) { - if (i < num_removed_comps) - comps[i] = nir_undef(&b, 1, bit_size); - else - comps[i] = nir_channel(&b, value0, i - num_removed_comps); - } - value0 = nir_vec(&b, comps, num_comps); - } - nir_push_else(&b, if_first_thread); - } - - value = load_ssbo_sparse(&b, num_comps, bit_size, nir_imm_int(&b, 0), - nir_iadd_imm(&b, offset, realign_offset), - (struct _nir_load_ssbo_indices){ - .access = ACCESS_RESTRICT, - .align_mul = 4, - .align_offset = (unsigned)realign_offset % 4 - }, key->src_is_sparse); - - - if (if_first_thread) { - nir_pop_if(&b, if_first_thread); - value = nir_if_phi(&b, value0, value); - } - - /* Bitcast the vector to 32 bits. */ - if (value->bit_size != 32) - value = nir_extract_bits(&b, &value, 1, 0, key->dwords_per_thread, 32); - } - - nir_def *dst_buf = nir_imm_int(&b, !key->is_clear); - nir_if *if_first_thread = NULL, *if_last_thread = NULL; - - if (!key->dst_single_thread_unaligned) { - /* dst_align_offset means how many bytes the first thread should skip because the offset of - * the buffer binding is rounded down to the clear/copy size of thread, causing the bytes - * before dst_align_offset to be writable. Above we used realign_offset to byte-shift - * the value to compensate for the rounded-down offset, so that all stores are dword stores - * regardless of the offset/size alignment except that the first thread shouldn't store - * the first dst_align_offset bytes, and the last thread should only store the first - * dst_last_thread_bytes. In both cases, there is a dword that must be only partially - * written by splitting it into 8-bit and 16-bit stores. - */ - if (key->dst_align_offset) { - if_first_thread = nir_push_if(&b, nir_ieq_imm(&b, thread_id, 0)); - { - unsigned local_offset = key->dst_align_offset; - nir_def *first_dword = nir_channel(&b, value, local_offset / 4); - - if (local_offset % 2 == 1) { - nir_store_ssbo(&b, nir_channel(&b, nir_unpack_32_4x8(&b, first_dword), local_offset % 4), - dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), - .access = ACCESS_RESTRICT); - local_offset++; - } - - if (local_offset % 4 == 2) { - nir_store_ssbo(&b, nir_unpack_32_2x16_split_y(&b, first_dword), dst_buf, - nir_iadd_imm_nuw(&b, offset, local_offset), - .access = ACCESS_RESTRICT); - local_offset += 2; - } - - assert(local_offset % 4 == 0); - unsigned num_dw_remaining = key->dwords_per_thread - local_offset / 4; - - if (num_dw_remaining) { - nir_def *dwords = - nir_channels(&b, value, BITFIELD_RANGE(local_offset / 4, num_dw_remaining)); - - nir_store_ssbo(&b, dwords, dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), - .access = ACCESS_RESTRICT); - } - } - nir_push_else(&b, if_first_thread); - } - - if (key->dst_last_thread_bytes) { - nir_def *last_thread_id = - nir_channel(&b, nir_load_user_data_amd(&b), last_thread_user_data_index); - - if_last_thread = nir_push_if(&b, nir_ieq(&b, thread_id, last_thread_id)); - { - unsigned num_dwords = key->dst_last_thread_bytes / 4; - bool write_short = (key->dst_last_thread_bytes - num_dwords * 4) / 2; - bool write_byte = key->dst_last_thread_bytes % 2; - nir_def *last_dword = nir_channel(&b, value, num_dwords); - - if (num_dwords) { - nir_def *dwords = nir_channels(&b, value, BITFIELD_MASK(num_dwords)); - nir_store_ssbo(&b, dwords, dst_buf, offset, .access = ACCESS_RESTRICT); - } - - if (write_short) { - nir_store_ssbo(&b, nir_u2u16(&b, last_dword), dst_buf, - nir_iadd_imm_nuw(&b, offset, num_dwords * 4), - .access = ACCESS_RESTRICT); - } - - if (write_byte) { - nir_store_ssbo(&b, nir_channel(&b, nir_unpack_32_4x8(&b, last_dword), write_short * 2), - dst_buf, nir_iadd_imm_nuw(&b, offset, num_dwords * 4 + write_short * 2), - .access = ACCESS_RESTRICT); - } - } - nir_push_else(&b, if_last_thread); - } - - nir_store_ssbo(&b, value, dst_buf, offset, .access = ACCESS_RESTRICT); - - if (if_last_thread) - nir_pop_if(&b, if_last_thread); - if (if_first_thread) - nir_pop_if(&b, if_first_thread); - } else { - /* This shader only executes a single thread (tiny copy or clear) and it's unaligned at both - * the beginning and the end. Walk the individual dwords/words/bytes that should be written - * to split the store accordingly. - */ - for (unsigned local_offset = key->dst_align_offset; - local_offset < key->dst_last_thread_bytes;) { - unsigned remaining = key->dst_last_thread_bytes - local_offset; - nir_def *src_dword = nir_channel(&b, value, local_offset / 4); - - if (local_offset % 2 == 1 || remaining == 1) { - /* 1-byte store. */ - nir_def *src_dword4x8 = nir_unpack_32_4x8(&b, src_dword); - nir_store_ssbo(&b, nir_channel(&b, src_dword4x8, local_offset % 4), dst_buf, - nir_iadd_imm_nuw(&b, offset, local_offset), .access = ACCESS_RESTRICT); - local_offset++; - } else if (local_offset % 4 == 2 || remaining == 2 || remaining == 3) { - /* 2-byte store. */ - nir_def *src_dword2x16 = nir_unpack_32_2x16(&b, src_dword); - nir_store_ssbo(&b, nir_channel(&b, src_dword2x16, (local_offset / 2) % 2), dst_buf, - nir_iadd_imm_nuw(&b, offset, local_offset), .access = ACCESS_RESTRICT); - local_offset += 2; - } else { - /* 1-N dwords. */ - unsigned dw_size = remaining / 4; - assert(dw_size); - assert(local_offset % 4 == 0); - - nir_store_ssbo(&b, nir_channels(&b, value, BITFIELD_RANGE(local_offset / 4, dw_size)), - dst_buf, nir_iadd_imm_nuw(&b, offset, local_offset), - .access = ACCESS_RESTRICT); - local_offset += dw_size * 4; - } - } - } - - if (key->has_start_thread) - nir_pop_if(&b, if_positive); - - return si_create_shader_state(sctx, b.shader); -} - /* Load samples from the image, and copy them to the same image. This looks like * a no-op, but it's not. Loads use FMASK, while stores don't, so samples are * reordered to match expanded FMASK.