diff --git a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp index 6f0d1ba7fd..e8a3064de6 100644 --- a/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp +++ b/example/01_gemm/gemm_xdl_fp16_pk_i4_v3.cpp @@ -45,13 +45,6 @@ using DeviceGemmV2Instance = // clang-format on -using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceGemm; using ReferenceGemmInstance = ck::tensor_operation::host::ReferenceGemm #include -#include "ck/utility/common_header.hpp" +#include "ck/host_utility/device_prop.hpp" +#include "ck/host_utility/flush_cache.hpp" +#include "ck/host_utility/kernel_launch.hpp" #include "ck/tensor_description/tensor_descriptor.hpp" #include "ck/tensor_description/tensor_descriptor_helper.hpp" -#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/device/device_gemm_v2.hpp" #include "ck/tensor_operation/gpu/device/gemm_specialization.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp" -#include "ck/host_utility/device_prop.hpp" -#include "ck/host_utility/kernel_launch.hpp" -#include "ck/host_utility/flush_cache.hpp" +#include "ck/utility/common_header.hpp" namespace ck { namespace tensor_operation { diff --git a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp index a39d39a18f..826cad5996 100644 --- a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp +++ b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp @@ -41,6 +41,32 @@ __host__ __device__ inline half4_t pki4_to_half4(int q) res.template AsType()(Number<1>{}) = amd_assembly_pk_fma_f16( bit_cast(hi), bit_cast(MUL), bit_cast(ADD)); + return res.template AsType()[Number<0>{}]; +} + +__host__ __device__ inline half4_t pki4_to_half4_scale(int q, const ck::half2_t& scale) +{ + const int LO = 0x000f000f; + const int HI = 0x00f000f0; + const int EX = 0x64006400; + + // Extract the two int4 at low bit and create two fp16 number. + int lo = amd_assembly_and_or_b32(q, LO, EX); + // Extract the two int4 at hight bit and create two fp16 number. + int hi = amd_assembly_and_or_b32(q, HI, EX); + + const int SUB = 0xE408E408; // half2 {-1032, -1032} + const int MUL = 0x2c002c00; // half2 {1 / 16, 1 / 16} + const int ADD = 0xd480d480; // half2 {-72, -72} + + vector_type res; + + res.template AsType()(Number<0>{}) = + amd_assembly_pk_add_f16(bit_cast(lo), bit_cast(SUB)); + + res.template AsType()(Number<1>{}) = amd_assembly_pk_fma_f16( + bit_cast(hi), bit_cast(MUL), bit_cast(ADD)); + asm volatile("v_pk_mul_f16 %0, %1, %2" : "=v"(res.template AsType()(Number<0>{})) : "v"(res.template AsType()(Number<0>{})), "v"(scale)); @@ -52,64 +78,6 @@ __host__ __device__ inline half4_t pki4_to_half4(int q) return res.template AsType()[Number<0>{}]; } -// Further fuse the scale into inline assembly, sanity failed -#if 0 -__host__ __device__ inline half4_t pki4_to_half4_scale(int q, const ck::half_t& scale) -{ - constexpr int LO = 0x000f000f; - constexpr int HI = 0x00f000f0; - constexpr int EX = 0x64006400; - // Guarantee that the `(a & b) | c` operations are LOP3s. - // int lo = lop3<(0xf0 & 0xcc) | 0xaa>(q, LO, EX); - // int hi = lop3<(0xf0 & 0xcc) | 0xaa>(q, HI, EX); - int lo = amd_assembly_and_or_b32(q, LO, EX); - int hi = amd_assembly_and_or_b32(q, HI, EX); - // We want signed int4 outputs, hence we fuse the `-8` symmetric zero point - // directly into `SUB` and `ADD`. - // constexpr int SUB = 0xE408E408; //-8 - // constexpr int MUL = 0x2c002c00; // 1/16 - // constexpr int ADD = 0xd480d480; //-79 - constexpr half_t SUB = bit_cast(static_cast(0xE408)); - constexpr half_t MUL = bit_cast(static_cast(0x2c00)); - constexpr half_t ADD = bit_cast(static_cast(0xd480)); - - vector_type scale_2; - scale_2.template AsType()(Number<0>{}) = scale; - scale_2.template AsType()(Number<1>{}) = scale; - vector_type sub_2; - sub_2.template AsType()(Number<0>{}) = SUB * scale; - sub_2.template AsType()(Number<1>{}) = SUB * scale; - vector_type mul_2; - mul_2.template AsType()(Number<0>{}) = MUL * scale; - mul_2.template AsType()(Number<1>{}) = MUL * scale; - vector_type add_2; - add_2.template AsType()(Number<0>{}) = ADD * scale; - add_2.template AsType()(Number<1>{}) = ADD * scale; - - vector_type res; - - res.template AsType()(Number<0>{}) = - amd_assembly_pk_fma_f16(bit_cast(lo), - scale_2.template AsType()(Number<0>{}), - sub_2.template AsType()(Number<0>{})); - - res.template AsType()(Number<1>{}) = - amd_assembly_pk_fma_f16(bit_cast(hi), - mul_2.template AsType()(Number<0>{}), - add_2.template AsType()(Number<0>{})); - - // asm volatile("v_pk_mul_f16 %0, %1, %2" - // : "=v"(res.template AsType()(Number<0>{})) - // : "v"(res.template AsType()(Number<0>{})), "v"(scale)); - - // asm volatile("v_pk_mul_f16 %0, %1, %2" - // : "=v"(res.template AsType()(Number<1>{})) - // : "v"(res.template AsType()(Number<1>{})), "v"(scale)); - - return res.template AsType()[Number<0>{}]; -} -#endif - __host__ __device__ inline half2_t pki4_to_half2(pk_i4_t q) { #if 1 diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp index d928760b17..9ba468529b 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp @@ -3,16 +3,16 @@ #pragma once -#include "ck/utility/common_header.hpp" #include "ck/tensor_description/multi_index_transform_helper.hpp" #include "ck/tensor_description/tensor_descriptor.hpp" #include "ck/tensor_description/tensor_descriptor_helper.hpp" -#include "ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp" #include "ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_b_scale_selector.hpp" #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp" #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v6r1.hpp" -#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" +#include "ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp" +#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp" +#include "ck/utility/common_header.hpp" namespace ck { @@ -29,7 +29,7 @@ template __global__ void #if CK_USE_LAUNCH_BOUNDS - __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy) +__launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy) #endif // __attribute__((amdgpu_waves_per_eu(1, 1))) kernel_gemm_xdl_cshuffle_v3(typename GridwiseGemm::Argument karg) @@ -59,7 +59,7 @@ template __global__ void #if CK_USE_LAUNCH_BOUNDS - __launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy) +__launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy) #endif // __attribute__((amdgpu_waves_per_eu(1, 1))) kernel_gemm_xdl_cshuffle_v3_2lds(typename GridwiseGemm::Argument karg) @@ -1926,7 +1926,7 @@ struct GridwiseGemm_xdl_cshuffle_v3 1, false>( b_scale_grid_desc_bn_ak, - make_multi_index(block_n_id * NPerBlock / ScaleBlockN + b_thread_offset_n, + make_multi_index(block_n_id * NPerBlock / ScaleBlockN + b_thread_offset_n, b_thread_offset_k / ScaleBlockK)); constexpr auto b_scale_thread_slice_copy_step = diff --git a/include/ck/utility/amd_inline_asm.hpp b/include/ck/utility/amd_inline_asm.hpp index 8e281a929f..113f3af4ae 100644 --- a/include/ck/utility/amd_inline_asm.hpp +++ b/include/ck/utility/amd_inline_asm.hpp @@ -4,8 +4,8 @@ #ifndef CK_AMD_INLINE_ASM_HPP #define CK_AMD_INLINE_ASM_HPP -#include "data_type.hpp" #include "c_style_pointer_cast.hpp" +#include "data_type.hpp" // TODO: deprecate all amd_assembly_outer_product_xxx diff --git a/library/include/ck/library/tensor_operation_instance/gpu/gemm_b_scale.hpp b/library/include/ck/library/tensor_operation_instance/gpu/gemm_b_scale.hpp index e235b047a2..93eed31bc5 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/gemm_b_scale.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/gemm_b_scale.hpp @@ -3,12 +3,12 @@ #pragma once -#include -#include #include "ck/ck.hpp" -#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_scale.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" +#include +#include #include "ck/library/tensor_operation_instance/device_operation_instance_factory.hpp" diff --git a/library/src/tensor_operation_instance/gpu/gemm_b_scale/device_gemm_b_scale_xdl_f16_i4_f16/device_gemm_b_scale_xdl_f16_i4_f16_mk_nk_mn.hpp b/library/src/tensor_operation_instance/gpu/gemm_b_scale/device_gemm_b_scale_xdl_f16_i4_f16/device_gemm_b_scale_xdl_f16_i4_f16_mk_nk_mn.hpp index 6641961b6f..1822296f53 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_b_scale/device_gemm_b_scale_xdl_f16_i4_f16/device_gemm_b_scale_xdl_f16_i4_f16_mk_nk_mn.hpp +++ b/library/src/tensor_operation_instance/gpu/gemm_b_scale/device_gemm_b_scale_xdl_f16_i4_f16/device_gemm_b_scale_xdl_f16_i4_f16_mk_nk_mn.hpp @@ -2,9 +2,9 @@ // Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved. #include "ck/ck.hpp" -#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/device/gemm_specialization.hpp" #include "ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_scale.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp" diff --git a/profiler/include/profiler/profile_gemm_b_scale_impl.hpp b/profiler/include/profiler/profile_gemm_b_scale_impl.hpp index 6e8237a00e..c52e65b647 100644 --- a/profiler/include/profiler/profile_gemm_b_scale_impl.hpp +++ b/profiler/include/profiler/profile_gemm_b_scale_impl.hpp @@ -8,18 +8,18 @@ #include #include "ck/ck.hpp" -#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/device/impl/device_gemm_xdl_cshuffle_v3_b_scale.hpp" +#include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/library/tensor_operation_instance/gpu/gemm_b_scale.hpp" +#include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp" #include "ck/library/utility/check_err.hpp" #include "ck/library/utility/device_memory.hpp" #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" #include "ck/library/utility/literals.hpp" -#include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp" namespace ck { namespace profiler { diff --git a/profiler/src/profile_gemm_b_scale.cpp b/profiler/src/profile_gemm_b_scale.cpp index f6deaaa89b..623b33c941 100644 --- a/profiler/src/profile_gemm_b_scale.cpp +++ b/profiler/src/profile_gemm_b_scale.cpp @@ -1,10 +1,10 @@ // SPDX-License-Identifier: MIT // Copyright (c) 2023-2024, Advanced Micro Devices, Inc. All rights reserved. +#include +#include #include #include -#include -#include #include "profiler/profile_gemm_b_scale_impl.hpp" #include "profiler_operation_registry.hpp"