AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

radeonsi: implement optimized unaligned clear/copy_buffer compute shader

Browse Source 
This totally beats CP DMA on Navi31.

Acked-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/30208>
This commit is contained in:
Marek Olšák
2024-07-14 21:05:18 -04:00
committed by Marge Bot
parent fa85b4b49e
commit 2f9201e91b
3 changed files with 315 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/gallium/drivers/radeonsi/si_compute_blit.c
+70 -16
View File
@@ -255,6 +255,18 @@ void si_compute_clear_buffer_rmw(struct si_context *sctx, struct pipe_resource *
1, &sb, 0x1);
}
/**
* This implements a clear/copy_buffer compute shader allowing an arbitrary src_offset, dst_offset,
* and size alignment, so that it can be used as a complete replacement for the typically slower
* CP DMA.
*
* It stores 16B blocks per thread aligned to a 16B offset just like a 16B-aligned clear/copy,
* and it byte-shifts src data by the amount of both src and dst misalignment to get the behavior
* of a totally unaligned clear/copy.
*
* The first and last thread can store less than 16B (up to 1B store granularity) depending on how
* much dst is unaligned.
*/
bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource *dst,
unsigned dst_offset, struct pipe_resource *src,
unsigned src_offset, unsigned size,
@@ -330,27 +342,51 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource
assert(dst->target != PIPE_BUFFER || dst_offset + size <= dst->width0);
assert(!src || src_offset + size <= src->width0);
unsigned num_threads = DIV_ROUND_UP(size, dwords_per_thread * 4);
struct pipe_grid_info info = {};
set_work_size(&info, 64, 1, 1, num_threads, 1, 1);
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;
struct pipe_shader_buffer sb[2] = {};
sb[is_copy].buffer = dst;
sb[is_copy].buffer_offset = dst_offset;
sb[is_copy].buffer_size = size;
/* 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);
if (is_copy) {
sb[0].buffer = src;
sb[0].buffer_offset = src_offset;
sb[0].buffer_size = size;
/* 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)));
for (unsigned i = 0; i < 4; i++)
sctx->cs_user_data[i] = clear_value[i % (clear_value_size / 4)];
/* 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;
@@ -360,6 +396,15 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource
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;
void *shader = _mesa_hash_table_u64_search(sctx->cs_dma_shaders, key.key);
if (!shader) {
@@ -367,6 +412,13 @@ bool si_compute_clear_copy_buffer(struct si_context *sctx, struct pipe_resource
_mesa_hash_table_u64_insert(sctx->cs_dma_shaders, 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;
struct pipe_grid_info info = {};
set_work_size(&info, 64, 1, 1, num_threads, 1, 1);
si_launch_grid_internal_ssbos(sctx, &info, shader, flags, coher, is_copy ? 2 : 1, sb,
is_copy ? 0x2 : 0x1);
return true;
@@ -394,12 +446,14 @@ void si_clear_buffer(struct si_context *sctx, struct pipe_resource *dst,
if (util_lower_clearsize_to_dword(clear_value, (int*)&clear_value_size, &clamped))
clear_value = &clamped;
if (method != SI_CP_DMA_CLEAR_METHOD &&
si_compute_clear_copy_buffer(sctx, dst, offset, NULL, 0, size, clear_value,
clear_value_size, flags, coher, 0,
method == SI_AUTO_SELECT_CLEAR_METHOD))
return;
uint64_t aligned_size = size & ~3ull;
if (aligned_size &&
(method == SI_CP_DMA_CLEAR_METHOD ||
!si_compute_clear_copy_buffer(sctx, dst, offset, NULL, 0, aligned_size, clear_value,
clear_value_size, flags, coher, 0,
method == SI_AUTO_SELECT_CLEAR_METHOD))) {
if (aligned_size) {
assert(clear_value_size == 4);
assert(!(flags & SI_OP_CS_RENDER_COND_ENABLE));
si_cp_dma_clear_buffer(sctx, &sctx->gfx_cs, dst, offset, aligned_size, *clear_value,
src/gallium/drivers/radeonsi/si_pipe.h
+7 -1
View File
@@ -1648,7 +1648,13 @@ union si_cs_clear_copy_buffer_key {
struct {
bool is_clear:1;
unsigned dwords_per_thread:3; /* 1..4 allowed */
unsigned clear_value_size_is_12:1;
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 */
};
uint64_t key;
};
src/gallium/drivers/radeonsi/si_shaderlib_nir.c
+238 -4
View File
@@ -227,6 +227,39 @@ 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)
{
@@ -235,6 +268,11 @@ void *si_create_dma_compute_shader(struct si_context *sctx, union si_cs_clear_co
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");
}
@@ -247,7 +285,11 @@ void *si_create_dma_compute_shader(struct si_context *sctx, union si_cs_clear_co
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 =
key->is_clear ? (key->clear_value_size_is_12 ? 3 : key->dwords_per_thread) : 0;
(key->is_clear ? (key->clear_value_size_is_12 ? 3 : key->dwords_per_thread) : 0);
/* Add the last thread ID value. */
if (key->dst_last_thread_bytes)
b.shader->info.cs.user_data_components_amd = key->is_clear ? 5 : 1;
nir_def *thread_id = ac_get_global_ids(&b, 1, 32);
/* Convert the global thread ID into bytes. */
@@ -277,11 +319,203 @@ void *si_create_dma_compute_shader(struct si_context *sctx, union si_cs_clear_co
value = nir_vec4(&b, vec[0], vec[1], vec[2], vec[0]);
}
} else {
value = nir_load_ssbo(&b, key->dwords_per_thread, 32, nir_imm_int(&b, 0), offset,
.access = ACCESS_RESTRICT);
/* 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_store_ssbo(&b, value, nir_imm_int(&b, !key->is_clear), offset, .access = ACCESS_RESTRICT);
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), key->is_clear ? 4 : 0);
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;
}
}
}
return si_create_shader_state(sctx, b.shader);
}