From b18d739672d7db2044ea31f783cbd93122dd2c65 Mon Sep 17 00:00:00 2001 From: zjing14 Date: Tue, 11 Apr 2023 07:44:43 -0500 Subject: [PATCH] add a marco to turn on/off denorm fix (off by default) (#673) * add a marco to turn off denorm fix by default * expose the marco --------- Co-authored-by: root [ROCm/composable_kernel commit: c54f8bcc25ace7b8d9ee86ddeb72738c87f908bb] --- include/ck/ck.hpp | 5 +++++ .../gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp | 2 +- 2 files changed, 6 insertions(+), 1 deletion(-) diff --git a/include/ck/ck.hpp b/include/ck/ck.hpp index 9853f19a47..036ca24a44 100644 --- a/include/ck/ck.hpp +++ b/include/ck/ck.hpp @@ -168,6 +168,11 @@ // flag to enable (1) or disable (0) the debugging output in some kernels #define DEBUG_LOG 0 +// denorm test fix, required to work around dissue +#ifndef CK_WORKAROUND_DENORM_FIX +#define CK_WORKAROUND_DENORM_FIX 0 +#endif + namespace ck { enum struct InMemoryDataOperationEnum diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp index d49c96f86c..98a71a7c24 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp @@ -96,7 +96,7 @@ struct GridwiseGemmMultipleD_xdl_cshuffle // we convert fp16->fp32->bf16 and execute bf16 mfma instruction // when mfma if fixed, remove this section and update // ABDataTypeAdjusted -> ABDataType throughout this file -#if defined(__gfx90a__) +#if CK_WORKAROUND_DENORM_FIX && defined(__gfx90a__) using ABDataTypeAdjusted = conditional_t, ck::bhalf_t, ABDataType>; #else