fix clang warning suppression

CMakeLists.txt

@@ -166,21 +166,11 @@ enable_cppcheck(
         duplicateCondition
         noExplicitConstructor
         passedByValue
-        # preprocessorErrorDirective
+        preprocessorErrorDirective
         shadowVariable
         unusedFunction
         unusedPrivateFunction
         unusedStructMember
-        # Ignore initializer lists in the tests
-       #useInitializationList:*test/*.cpp
-       #*:*src/sqlite/*.cpp
-       #*:*.cl
-       #*:*src/kernels/*.h
-       #knownConditionTrueFalse:*src/kernels/composable_kernel/*/*
-       #redundantAssignment:*src/kernels/composable_kernel/*/*
-       #unreadVariable:*src/kernels/composable_kernel/*/*
-       #unusedScopedObject:*src/kernels/composable_kernel/*/*
-       #wrongPrintfScanfArgNum:*src/kernels/composable_kernel/*/*
         unmatchedSuppression
     FORCE
     SOURCES

composable_kernel/include/utility/amd_address_space.hpp

@@ -2,6 +2,7 @@
 #define CK_AMD_ADDRESS_SPACE_HPP
 
 #include "config.hpp"
+#include "c_style_pointer_cast.hpp"
 
 // Address Space for AMDGCN
 // https://llvm.org/docs/AMDGPUUsage.html#address-space
@@ -21,9 +22,9 @@ template <typename T>
 __device__ T* cast_pointer_to_generic_address_space(T CONSTANT* p)
 {
     // cast a pointer in "Constant" address space (4) to "Generic" address space (0)
-    // only old style cast seems be able to be compiled
-#pragma clang diagnostic ignored "-Wold-style-cast"
+    // only c-style pointer cast seems be able to be compiled
 #pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wold-style-cast"
     return (T*)p; // NOLINT(old-style-cast)
 #pragma clang diagnostic pop
 }
@@ -32,9 +33,9 @@ template <typename T>
 __host__ __device__ T CONSTANT* cast_pointer_to_constant_address_space(T* p)
 {
     // cast a pointer in "Generic" address space (0) to "Constant" address space (4)
-    // only old style cast seems be able to be compiled
-#pragma clang diagnostic ignored "-Wold-style-cast"
+    // only c-style pointer cast seems be able to be compiled
 #pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wold-style-cast"
     return (T CONSTANT*)p; // NOLINT(old-style-cast)
 #pragma clang diagnostic pop
 }

composable_kernel/include/utility/c_style_pointer_cast.hpp

@@ -10,9 +10,10 @@ template <typename PY,
           typename std::enable_if<is_pointer_v<PY> && is_pointer_v<PX>, bool>::type = false>
 __host__ __device__ PY c_style_pointer_cast(PX p_x)
 {
-#pragma clang diagnostic ignored "-Wold-style-cast"
 #pragma clang diagnostic push
-    return (PY)p_x; // NOLINT(old-style-cast)
+#pragma clang diagnostic ignored "-Wold-style-cast"
+#pragma clang diagnostic ignored "-Wcast-align"
+    return (PY)p_x; // NOLINT(old-style-cast, cast-align)
 #pragma clang diagnostic pop
 }
 

host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhwk.hpp

@@ -81,7 +81,7 @@ void device_dynamic_convolution_forward_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhw
     using GemmBBlockTransferDstVectorTensorLengths_K0_N0_N1_K1 = Sequence<1, 1, 1, 1>;
 
     constexpr index_t GemmCThreadTransferDstScalarPerVector_N11 = 4;
-#elif 0
+#elif 1
     // [M, N, K0, K1] = [128, 128, 8, 2] for fp16
     // cdata = 64, BlockSize = 256
     constexpr index_t BlockSize = 256;

host/driver_offline/src/conv_fwd_driver_offline.cpp

@@ -20,12 +20,12 @@
 #include "device_dynamic_convolution_forward_implicit_gemm_v4r4r4_xdlops_nhwc_kyxc_nhwk.hpp"
 
 #define USE_DYNAMIC_MODE 1
-#define USE_CONV_FWD_V4R4_NCHW 1
+#define USE_CONV_FWD_V4R4_NCHW 0
 #define USE_CONV_FWD_V4R4R2_NHWC 1
-#define USE_CONV_FWD_V6R1_NCHW 1
+#define USE_CONV_FWD_V6R1_NCHW 0
 #define USE_CONV_FWD_V5R1_NCHW 0
-#define USE_CONV_FWD_V4R4R2_XDL_NCHW 1
-#define USE_CONV_FWD_V4R4R4_XDL_NHWC 1
+#define USE_CONV_FWD_V4R4R2_XDL_NCHW 0
+#define USE_CONV_FWD_V4R4R4_XDL_NHWC 0
 
 enum ConvForwardAlgo
 {
@@ -126,7 +126,7 @@ int main(int argc, char* argv[])
     const index_t Wo = (Wi + in_left_pad_w + in_right_pad_w - XEff) / conv_stride_w + 1;
 #endif
 
-#if 1
+#if 0
     using in_data_t  = float;
     using acc_data_t = float;
     using out_data_t = float;