ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 11:16:59 +00:00
Code Issues Packages Projects Releases Wiki Activity

[buffer store/load] Replace memcpy with __builtin_memcpy in cast_to_amdgpu_buffer_rsrc_t

Browse Source 
We used memcpy to implement a bitcast of the opaque type
amdgcn_buffer_rsrc. However, hip's implementation of memcpy did not
allow the compiler to infer that the result of the copy of a uniform value
was also uniform.

This resulted in a waterfall loop over every value that the copy could
take (and a loss in performance).

When we use __builtin_memcpy, the optimizer correctly handles the
uniform copy.

Solves SWDEV-537500
This commit is contained in:
Juan Manuel Martinez Caamaño
2025-06-16 17:32:50 +02:00
parent a0f4db8d9c
commit 1810074bd7
1 changed files with 1 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
include/ck_tile/core/arch/amd_buffer_addressing.hpp
View File

@@ -92,7 +92,7 @@ CK_TILE_DEVICE __amdgpu_buffer_rsrc_t cast_to_amdgpu_buffer_rsrc_t(int32x4_t res
{
__amdgpu_buffer_rsrc_t as_rsrc;
static_assert(sizeof(res) == sizeof(as_rsrc) && "Size of buffer resource should match");
memcpy(&as_rsrc, &res, sizeof(res));
__builtin_memcpy(&as_rsrc, &res, sizeof(res));
return as_rsrc;
}
#endif