From ddd49ec9e7f829c3a4ca92de22206c708d8d94a7 Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Tue, 10 Aug 2021 06:20:24 +0000 Subject: [PATCH] fix clang warning suppression --- CMakeLists.txt | 12 +----------- .../include/utility/amd_address_space.hpp | 9 +++++---- .../include/utility/c_style_pointer_cast.hpp | 5 +++-- ...ard_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhwk.hpp | 2 +- host/driver_offline/src/conv_fwd_driver_offline.cpp | 10 +++++----- 5 files changed, 15 insertions(+), 23 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e5a5258328..6581bae8b9 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/composable_kernel/include/utility/amd_address_space.hpp b/composable_kernel/include/utility/amd_address_space.hpp index a8010f951c..24c95b27af 100644 --- a/composable_kernel/include/utility/amd_address_space.hpp +++ b/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 __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 __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 } diff --git a/composable_kernel/include/utility/c_style_pointer_cast.hpp b/composable_kernel/include/utility/c_style_pointer_cast.hpp index f4b4b09a76..29cd3d07ca 100644 --- a/composable_kernel/include/utility/c_style_pointer_cast.hpp +++ b/composable_kernel/include/utility/c_style_pointer_cast.hpp @@ -10,9 +10,10 @@ template && is_pointer_v, 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 } diff --git a/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhwk.hpp b/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhwk.hpp index 667150317e..64557b642e 100644 --- a/host/driver_offline/include/device_dynamic_convolution_forward_implicit_gemm_v4r4r2_dlops_nhwc_kyxc_nhwk.hpp +++ b/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; diff --git a/host/driver_offline/src/conv_fwd_driver_offline.cpp b/host/driver_offline/src/conv_fwd_driver_offline.cpp index 3358d5b98c..f88d1b831e 100644 --- a/host/driver_offline/src/conv_fwd_driver_offline.cpp +++ b/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;