Skip to content

GitLab

Commit 834eb24c authored by Chao Liu's avatar Chao Liu
Browse files

buffer addressing using uint32_t for addr shifting

parent bd4c668e
Hide whitespace changes
Inline Side-by-side
Showing with 51 additions and 53 deletions
composable_kernel/include/utility/amd_buffer_addressing.hpp
View file @ 834eb24c
...@@ -205,10 +205,10 @@ __device__ float amd_buffer_load<float, 1>(const float* p_src_wave, ...@@ -205,10 +205,10 @@ __device__ float amd_buffer_load<float, 1>(const float* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
return __llvm_amdgcn_buffer_load_f32( return __llvm_amdgcn_buffer_load_f32(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
#else #else
return src_thread_data_valid return src_thread_data_valid
...@@ -243,10 +243,10 @@ __device__ float2_t amd_buffer_load<float, 2>(const float* p_src_wave, ...@@ -243,10 +243,10 @@ __device__ float2_t amd_buffer_load<float, 2>(const float* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
return __llvm_amdgcn_buffer_load_f32x2( return __llvm_amdgcn_buffer_load_f32x2(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
} }
...@@ -275,10 +275,10 @@ __device__ float4_t amd_buffer_load<float, 4>(const float* p_src_wave, ...@@ -275,10 +275,10 @@ __device__ float4_t amd_buffer_load<float, 4>(const float* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
return __llvm_amdgcn_buffer_load_f32x4( return __llvm_amdgcn_buffer_load_f32x4(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
} }
...@@ -308,10 +308,10 @@ __device__ half_t amd_buffer_load<half_t, 1>(const half_t* p_src_wave, ...@@ -308,10 +308,10 @@ __device__ half_t amd_buffer_load<half_t, 1>(const half_t* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
return __llvm_amdgcn_buffer_load_f16( return __llvm_amdgcn_buffer_load_f16(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
#else #else
return src_thread_data_valid ? p_src_wave[src_thread_data_offset] : 0; return src_thread_data_valid ? p_src_wave[src_thread_data_offset] : 0;
...@@ -343,10 +343,10 @@ __device__ half2_t amd_buffer_load<half_t, 2>(const half_t* p_src_wave, ...@@ -343,10 +343,10 @@ __device__ half2_t amd_buffer_load<half_t, 2>(const half_t* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float dst_out_tmp = __llvm_amdgcn_buffer_load_f32( float dst_out_tmp = __llvm_amdgcn_buffer_load_f32(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<half2_t*>(&dst_out_tmp); return *reinterpret_cast<half2_t*>(&dst_out_tmp);
...@@ -377,10 +377,10 @@ __device__ half4_t amd_buffer_load<half_t, 4>(const half_t* p_src_wave, ...@@ -377,10 +377,10 @@ __device__ half4_t amd_buffer_load<half_t, 4>(const half_t* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float2_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x2( float2_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x2(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<half4_t*>(&dst_out_tmp); return *reinterpret_cast<half4_t*>(&dst_out_tmp);
...@@ -411,10 +411,10 @@ __device__ half8_t amd_buffer_load<half_t, 8>(const half_t* p_src_wave, ...@@ -411,10 +411,10 @@ __device__ half8_t amd_buffer_load<half_t, 8>(const half_t* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float4_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x4( float4_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x4(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<half8_t*>(&dst_out_tmp); return *reinterpret_cast<half8_t*>(&dst_out_tmp);
...@@ -446,10 +446,10 @@ __device__ ushort amd_buffer_load<ushort, 1>(const ushort* p_src_wave, ...@@ -446,10 +446,10 @@ __device__ ushort amd_buffer_load<ushort, 1>(const ushort* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
return __llvm_amdgcn_buffer_load_bf16( return __llvm_amdgcn_buffer_load_bf16(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
#else #else
...@@ -482,10 +482,10 @@ __device__ ushort2_t amd_buffer_load<ushort, 2>(const ushort* p_src_wave, ...@@ -482,10 +482,10 @@ __device__ ushort2_t amd_buffer_load<ushort, 2>(const ushort* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float dst_out_tmp = __llvm_amdgcn_buffer_load_f32( float dst_out_tmp = __llvm_amdgcn_buffer_load_f32(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<ushort2_t*>(&dst_out_tmp); return *reinterpret_cast<ushort2_t*>(&dst_out_tmp);
...@@ -516,10 +516,10 @@ __device__ ushort4_t amd_buffer_load<ushort, 4>(const ushort* p_src_wave, ...@@ -516,10 +516,10 @@ __device__ ushort4_t amd_buffer_load<ushort, 4>(const ushort* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float2_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x2( float2_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x2(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<ushort4_t*>(&dst_out_tmp); return *reinterpret_cast<ushort4_t*>(&dst_out_tmp);
...@@ -550,10 +550,10 @@ __device__ ushort8_t amd_buffer_load<ushort, 8>(const ushort* p_src_wave, ...@@ -550,10 +550,10 @@ __device__ ushort8_t amd_buffer_load<ushort, 8>(const ushort* p_src_wave,
false, false,
false); false);
#else #else
index_t src_addr_base = src_thread_data_valid ? 0 : 0x7fffffff; uint32_t src_addr_shift = src_thread_data_valid ? 0 : 0x7fffffff;
float4_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x4( float4_t dst_out_tmp = __llvm_amdgcn_buffer_load_f32x4(
src_wave_buffer_resource.data, 0, src_addr_base + src_thread_addr_offset, false, false); src_wave_buffer_resource.data, 0, src_addr_shift + src_thread_addr_offset, false, false);
#endif #endif
return *reinterpret_cast<ushort8_t*>(&dst_out_tmp); return *reinterpret_cast<ushort8_t*>(&dst_out_tmp);
...@@ -586,12 +586,12 @@ __device__ void amd_buffer_store<float, 1>(const float* p_src_thread, ...@@ -586,12 +586,12 @@ __device__ void amd_buffer_store<float, 1>(const float* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32(*p_src_thread, __llvm_amdgcn_buffer_store_f32(*p_src_thread,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -630,12 +630,12 @@ __device__ void amd_buffer_store<float, 2>(const float* p_src_thread, ...@@ -630,12 +630,12 @@ __device__ void amd_buffer_store<float, 2>(const float* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x2(*reinterpret_cast<const float2_t*>(p_src_thread), __llvm_amdgcn_buffer_store_f32x2(*reinterpret_cast<const float2_t*>(p_src_thread),
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -667,12 +667,12 @@ __device__ void amd_buffer_store<float, 4>(const float* p_src_thread, ...@@ -667,12 +667,12 @@ __device__ void amd_buffer_store<float, 4>(const float* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x4(*reinterpret_cast<const float4_t*>(p_src_thread), __llvm_amdgcn_buffer_store_f32x4(*reinterpret_cast<const float4_t*>(p_src_thread),
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -705,12 +705,12 @@ __device__ void amd_buffer_store<half_t, 1>(const half_t* p_src_thread, ...@@ -705,12 +705,12 @@ __device__ void amd_buffer_store<half_t, 1>(const half_t* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f16(*p_src_thread, __llvm_amdgcn_buffer_store_f16(*p_src_thread,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -751,12 +751,12 @@ __device__ void amd_buffer_store<half_t, 2>(const half_t* p_src_thread, ...@@ -751,12 +751,12 @@ __device__ void amd_buffer_store<half_t, 2>(const half_t* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32(*p_src_tmp, __llvm_amdgcn_buffer_store_f32(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -786,18 +786,16 @@ __device__ void amd_buffer_store<half_t, 4>(const half_t* p_src_thread, ...@@ -786,18 +786,16 @@ __device__ void amd_buffer_store<half_t, 4>(const half_t* p_src_thread,
__llvm_amdgcn_buffer_store_f32x2(*p_src_tmp, __llvm_amdgcn_buffer_store_f32x2(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_thread_data_valid dst_thread_data_valid ? dst_thread_addr_offset : 0xffffffff,
? dst_thread_addr_offset,
: 0xffffffff,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x2(*p_src_tmp, __llvm_amdgcn_buffer_store_f32x2(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -831,12 +829,12 @@ __device__ void amd_buffer_store<half_t, 8>(const half_t* p_src_thread, ...@@ -831,12 +829,12 @@ __device__ void amd_buffer_store<half_t, 8>(const half_t* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x4(*p_src_tmp, __llvm_amdgcn_buffer_store_f32x4(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -869,12 +867,12 @@ __device__ void amd_buffer_store<ushort, 1>(const ushort* p_src_thread, ...@@ -869,12 +867,12 @@ __device__ void amd_buffer_store<ushort, 1>(const ushort* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_bf16(*p_src_thread, __llvm_amdgcn_buffer_store_bf16(*p_src_thread,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -915,12 +913,12 @@ __device__ void amd_buffer_store<ushort, 2>(const ushort* p_src_thread, ...@@ -915,12 +913,12 @@ __device__ void amd_buffer_store<ushort, 2>(const ushort* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32(*p_src_tmp, __llvm_amdgcn_buffer_store_f32(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -954,12 +952,12 @@ __device__ void amd_buffer_store<ushort, 4>(const ushort* p_src_thread, ...@@ -954,12 +952,12 @@ __device__ void amd_buffer_store<ushort, 4>(const ushort* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x2(*p_src_tmp, __llvm_amdgcn_buffer_store_f32x2(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -993,12 +991,12 @@ __device__ void amd_buffer_store<ushort, 8>(const ushort* p_src_thread, ...@@ -993,12 +991,12 @@ __device__ void amd_buffer_store<ushort, 8>(const ushort* p_src_thread,
false, false,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_store_f32x4(*p_src_tmp, __llvm_amdgcn_buffer_store_f32x4(*p_src_tmp,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false, false,
false); false);
#endif #endif
...@@ -1029,12 +1027,12 @@ __device__ void amd_buffer_atomic_add<float, 1>(const float* p_src_thread, ...@@ -1029,12 +1027,12 @@ __device__ void amd_buffer_atomic_add<float, 1>(const float* p_src_thread,
dst_thread_data_valid ? dst_thread_addr_offset : 0xffffffff, dst_thread_data_valid ? dst_thread_addr_offset : 0xffffffff,
false); false);
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
__llvm_amdgcn_buffer_atomic_add_f32(*p_src_thread, __llvm_amdgcn_buffer_atomic_add_f32(*p_src_thread,
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset, dst_addr_shift + dst_thread_addr_offset,
false); false);
#endif #endif
} }
...@@ -1068,14 +1066,14 @@ __device__ void amd_buffer_atomic_add<float, 2>(const float* p_src_thread, ...@@ -1068,14 +1066,14 @@ __device__ void amd_buffer_atomic_add<float, 2>(const float* p_src_thread,
false); false);
} }
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
for(index_t i = 0; i < 2; ++i) for(index_t i = 0; i < 2; ++i)
{ {
__llvm_amdgcn_buffer_atomic_add_f32(p_src_thread[i], __llvm_amdgcn_buffer_atomic_add_f32(p_src_thread[i],
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset + dst_addr_shift + dst_thread_addr_offset +
i * sizeof(float), i * sizeof(float),
false); false);
} }
...@@ -1111,14 +1109,14 @@ __device__ void amd_buffer_atomic_add<float, 4>(const float* p_src_thread, ...@@ -1111,14 +1109,14 @@ __device__ void amd_buffer_atomic_add<float, 4>(const float* p_src_thread,
false); false);
} }
#else #else
index_t dst_addr_base = dst_thread_data_valid ? 0 : 0x7fffffff; uint32_t dst_addr_shift = dst_thread_data_valid ? 0 : 0x7fffffff;
for(index_t i = 0; i < 4; ++i) for(index_t i = 0; i < 4; ++i)
{ {
__llvm_amdgcn_buffer_atomic_add_f32(p_src_thread[i], __llvm_amdgcn_buffer_atomic_add_f32(p_src_thread[i],
dst_wave_buffer_resource.data, dst_wave_buffer_resource.data,
0, 0,
dst_addr_base + dst_thread_addr_offset + dst_addr_shift + dst_thread_addr_offset +
i * sizeof(float), i * sizeof(float),
false); false);
} }
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment