diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp index 3112109167..1ab41f87bb 100644 --- a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp +++ b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp @@ -1247,6 +1247,13 @@ llvm_amdgcn_raw_buffer_store_i16x4(int16x4_t vdata, index_t soffset, index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v4i16"); +CK_TILE_DEVICE_EXTERN void +llvm_amdgcn_raw_buffer_store_i16x8(int16x8_t vdata, + int32x4_t rsrc, + index_t voffset, + index_t soffset, + index_t glc_slc) __asm("llvm.amdgcn.raw.buffer.store.v8i16"); + // buffer store i32 CK_TILE_DEVICE_EXTERN void llvm_amdgcn_raw_buffer_store_i32(int32_t vdata, @@ -2246,19 +2253,12 @@ CK_TILE_DEVICE void amd_buffer_store_impl(const thread_buffer src_thread_d } else if constexpr(N == 8) { - llvm_amdgcn_raw_buffer_store_i16x4( - src_thread_data.template get_as()[number<0>{}], + llvm_amdgcn_raw_buffer_store_i16x8( + bit_cast(src_thread_data), dst_wave_buffer_resource, dst_thread_addr_offset, dst_wave_addr_offset, static_cast(coherence)); - - llvm_amdgcn_raw_buffer_store_i16x4( - src_thread_data.template get_as()[number<1>{}], - dst_wave_buffer_resource, - dst_thread_addr_offset, - dst_wave_addr_offset + 4 * sizeof(bf16_t), - static_cast(coherence)); } } else if constexpr(std::is_same::value) diff --git a/include/ck_tile/core/arch/amd_buffer_coherence.hpp b/include/ck_tile/core/arch/amd_buffer_coherence.hpp index ee10a3d190..00fc71f077 100644 --- a/include/ck_tile/core/arch/amd_buffer_coherence.hpp +++ b/include/ck_tile/core/arch/amd_buffer_coherence.hpp @@ -139,6 +139,7 @@ enum struct amd_buffer_coherence_enum // Other archs compatiblity DEVICE_NT0 = 0, SYSTEM_NT0 = 0, + GROUP_NT0 = 0, DEVICE_NT1 = glc, SYSTEM_NT1 = slc, DEVICE = 0, // Required for template parsing (see GFX11 comment re: Two-Phase Lookup) diff --git a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp index 9b4dc30eb2..4bf4f252c3 100644 --- a/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp +++ b/include/ck_tile/ops/epilogue/cshuffle_epilogue.hpp @@ -341,17 +341,17 @@ struct CShuffleEpilogue "LDS row stride must be 4B-aligned for bank-word padding logic"); // calculate how many elements to pad to avoid bank conflict #if defined(__gfx950__) || defined(__gfx125__) -#if defined(__gfx950__) - constexpr index_t ElemsPer4B = BytesPerBank / ck_tile::gcd(BytesPerBank, DataTypeSize); - constexpr auto ToWords = [](index_t elems) constexpr { - return (elems * DataTypeSize) / BytesPerBank; - }; - constexpr index_t BaseWords = ToWords(BaseStrideElems); - constexpr index_t PadWords = ((BaseWords % 2) == 0) ? 1 : 0; - constexpr auto PaddingAmount = PadWords * ElemsPer4B; -#else +// #if defined(__gfx950__) +// constexpr index_t ElemsPer4B = BytesPerBank / ck_tile::gcd(BytesPerBank, DataTypeSize); +// constexpr auto ToWords = [](index_t elems) constexpr { +// return (elems * DataTypeSize) / BytesPerBank; +// }; +// constexpr index_t BaseWords = ToWords(BaseStrideElems); +// constexpr index_t PadWords = ((BaseWords % 2) == 0) ? 1 : 0; +// constexpr auto PaddingAmount = PadWords * ElemsPer4B; +// #else constexpr auto PaddingAmount = VectorLen; -#endif +// #endif constexpr auto lds_block_desc_0 = make_naive_tensor_descriptor( make_tuple(number{}, diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp index 823c4eef32..e73600655f 100644 --- a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp +++ b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_eight_waves_base.hpp @@ -493,11 +493,11 @@ struct GemmPipelineAgBgCrEightWavesImplBase : public GemmPipelineAgBgCrImplBase< if constexpr(HasHotLoop && TailNum == TailNumber::Even) { asm volatile(";; Even Tail Start ;;"); - __builtin_amdgcn_s_barrier(); + // __builtin_amdgcn_s_barrier(); main_body(I0, I1); - __builtin_amdgcn_s_barrier(); + // __builtin_amdgcn_s_barrier(); asm volatile(";; Even Tail End ;;"); - __builtin_amdgcn_s_barrier(); + // __builtin_amdgcn_s_barrier(); } constexpr int tic = HasHotLoop ? (TailNum == TailNumber::Odd ? 0 : 1) : 1 - N_LOOP % 2; diff --git a/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp b/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp index 32b83951ed..f2e483c2ee 100644 --- a/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp +++ b/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp @@ -520,7 +520,9 @@ struct QuantGemmMultiDKernel } else { - return make_naive_tensor_view( + return make_naive_tensor_view( a_ptr, make_tuple(k_size, kargs.M), make_tuple(kargs.stride_A, 1), @@ -679,7 +681,9 @@ struct QuantGemmMultiDKernel } else if constexpr(kQuantType == QuantType::RowColQuant) { - return make_naive_tensor_view( + return make_naive_tensor_view( aq_ptr, make_tuple(kargs.M, kargs.N), make_tuple(1, 0), // broadcasting over n @@ -829,7 +833,9 @@ struct QuantGemmMultiDKernel } else { - return make_naive_tensor_view( + return make_naive_tensor_view( b_ptr, make_tuple(kargs.N, k_size), make_tuple(kargs.stride_B, 1), @@ -906,7 +912,9 @@ struct QuantGemmMultiDKernel const auto& bq_tensor_view = [&]() { if constexpr(kQuantType == QuantType::RowColQuant) { - return make_naive_tensor_view( + return make_naive_tensor_view( bq_ptr, make_tuple(kargs.M, kargs.N), make_tuple(0, 1), // broadcasting over m @@ -1111,7 +1119,9 @@ struct QuantGemmMultiDKernel const auto& ds_tensor_view = generate_tuple( [&](auto i) { using DDataType_ = remove_cvref_t>; - return make_tensor_view( + return make_tensor_view( static_cast(ds_ptr[i]), ds_desc[i]); }, number{}); @@ -1187,7 +1197,9 @@ struct QuantGemmMultiDKernel const auto& c_tensor_view = [&]() { if constexpr(std::is_same_v) { - return make_naive_tensor_view( + return make_naive_tensor_view( c_ptr, make_tuple(kargs.M, kargs.N), make_tuple(kargs.stride_C, 1), @@ -1196,7 +1208,9 @@ struct QuantGemmMultiDKernel } else { - return make_naive_tensor_view( + return make_naive_tensor_view( c_ptr, make_tuple(kargs.M, kargs.N), make_tuple(1, kargs.stride_C),