diff --git a/projects/miopen/cmake/EnableCompilerWarnings.cmake b/projects/miopen/cmake/EnableCompilerWarnings.cmake index 10a610f1c4a7..c08a3611e329 100644 --- a/projects/miopen/cmake/EnableCompilerWarnings.cmake +++ b/projects/miopen/cmake/EnableCompilerWarnings.cmake @@ -24,64 +24,29 @@ # ################################################################################ -set(__default_cxx_compile_options +set(__cxx_compile_options + -Werror -Wall -Wextra - -Wcomment - -Wendif-labels - -Wformat - -Winit-self - -Wreturn-type - -Wsequence-point - -Wswitch - -Wtrigraphs + # Additional warnings not included in -Wall/-Wextra -Wundef - -Wuninitialized -Wunreachable-code - -Wno-ignored-qualifiers - -Wno-sign-compare + -Wmissing-noreturn + -Wshadow + -Wsuggest-override + -Wold-style-cast + # TODO: Working to enable these warnings. Each requires code cleanup first. + # -Wconversion # ~1000+ implicit narrowing/sign conversions to fix + # -Wdouble-promotion # implicit float-to-double promotions + # Suppress specific warnings -- working to remove these by fixing the code + -Wno-c++11-narrowing # ~40 instances: narrowing in brace init (batchnorm, ck_impl, addkernels) + -Wno-sign-compare # ~1000+ instances: signed/unsigned comparisons throughout codebase + -Wno-deprecated-declarations # 2 deprecated MIOpen APIs still have callers ) set(__clang_cxx_compile_options - -Weverything - -Wno-c++98-compat - -Wno-c++98-compat-pedantic - -Wno-conversion - -Wno-double-promotion - -Wno-exit-time-destructors - -Wno-extra-semi - -Wno-extra-semi-stmt - -Wno-missing-prototypes - -Wno-padded -Wno-unused-command-line-argument - -Wno-weak-vtables - -Wno-covered-switch-default - -Wno-unsafe-buffer-usage - -Wno-deprecated-declarations - -Wno-global-constructors - -Wno-reserved-identifier - -Wno-deprecated - -Wno-old-style-cast - -Wno-c++11-narrowing - -Wno-switch-enum - -Wno-suggest-override - -Wno-nonportable-system-include-path - -Wno-documentation - -Wmissing-noreturn) - -if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL "19") - list(APPEND __clang_cxx_compile_options - -Wno-unique-object-duplication - -Wno-switch-default) -endif() - -if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL "23") - list(APPEND __clang_cxx_compile_options - -Wno-lifetime-safety - -Wno-lifetime-safety-suggestions - -Wno-lifetime-safety-intra-tu-suggestions - -Wno-lifetime-safety-cross-tu-suggestions) -endif() +) if(WIN32) list(APPEND __clang_cxx_compile_options @@ -90,8 +55,9 @@ if(WIN32) endif() add_compile_options( - "$<$,$>:${__default_cxx_compile_options};${__clang_cxx_compile_options}>" + "$<$:${__cxx_compile_options}>" + "$<$,$>:${__clang_cxx_compile_options}>" ) -unset(__default_cxx_compile_options) +unset(__cxx_compile_options) unset(__clang_cxx_compile_options) diff --git a/projects/miopen/driver/kthvalue_driver.hpp b/projects/miopen/driver/kthvalue_driver.hpp index 75f7e5b535b2..44010c173107 100644 --- a/projects/miopen/driver/kthvalue_driver.hpp +++ b/projects/miopen/driver/kthvalue_driver.hpp @@ -294,7 +294,7 @@ int KthvalueDriver::RunForwardGPU() outputDesc, output_dev->GetMem(), indicesDesc, - (size_t*)indices_dev->GetMem(), + static_cast(indices_dev->GetMem()), k, dim, keepDim); diff --git a/projects/miopen/driver/rocrand_wrapper.cpp b/projects/miopen/driver/rocrand_wrapper.cpp index 2a7d910c4be7..fe8ba5fe9687 100644 --- a/projects/miopen/driver/rocrand_wrapper.cpp +++ b/projects/miopen/driver/rocrand_wrapper.cpp @@ -26,7 +26,14 @@ #include "rocrand_wrapper.hpp" +// Suppress warnings from rocrand's own headers: in TheRock builds rocrand/stage/include +// is injected globally as -I (not -isystem), so rocrand headers are not treated as system +// includes and their internal issues trigger our warnings. +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wold-style-cast" +#pragma clang diagnostic ignored "-Wundef" #include +#pragma clang diagnostic pop namespace gpumemrand { diff --git a/projects/miopen/src/CMakeLists.txt b/projects/miopen/src/CMakeLists.txt index 9e6f401b7506..cd4ee6ab6dcb 100644 --- a/projects/miopen/src/CMakeLists.txt +++ b/projects/miopen/src/CMakeLists.txt @@ -1010,10 +1010,10 @@ endif() if(hipblaslt_FOUND) target_link_libraries( MIOpen PRIVATE roc::hipblaslt ) - target_include_directories( MIOpen INTERFACE $ ) + target_include_directories( MIOpen SYSTEM INTERFACE $ ) endif() -target_include_directories( MIOpen PRIVATE ${ROCRAND_INCLUDE_DIR} ) +target_include_directories( MIOpen SYSTEM PRIVATE ${ROCRAND_INCLUDE_DIR} ) # For backward compatibility with ROCm 5.3 # Build with library libMLIRMIOpen diff --git a/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp b/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp index bc3a840e35e6..2c2da8e56644 100644 --- a/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp +++ b/projects/miopen/src/ck_impl/implicitgemm_ck_util.hpp @@ -433,7 +433,8 @@ template -bool IsCKArgsSupported(const ProblemDescriptionType& problem, const std::string& kernel_id) +bool IsCKArgsSupported([[maybe_unused]] const ProblemDescriptionType& problem, + [[maybe_unused]] const std::string& kernel_id) { #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL if(!kernel_id.empty()) @@ -687,7 +688,7 @@ ConvSolution InitAnyInvokerFactory(const ProblemDescriptionType& problem, } template -OutElemOp GetOutElementOp(const miopen::fusion::ActivationOpInvokeParam& activationOp) +OutElemOp GetOutElementOp([[maybe_unused]] const miopen::fusion::ActivationOpInvokeParam& activationOp) { #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL auto activationMode = activationOp.activMode; @@ -893,12 +894,12 @@ template -ConvSolution InitInvokerFactoryNCHW(const ExecutionContext& ctx, - const miopen::conv::ProblemDescription& problem, - const std::string& kernel_id, - const Input1TposeOp& input1_op, - const Input2TposeOp& input2_op, - const OutputTposeOp& output_op) +ConvSolution InitInvokerFactoryNCHW([[maybe_unused]] const ExecutionContext& ctx, + [[maybe_unused]] const miopen::conv::ProblemDescription& problem, + [[maybe_unused]] const std::string& kernel_id, + [[maybe_unused]] const Input1TposeOp& input1_op, + [[maybe_unused]] const Input2TposeOp& input2_op, + [[maybe_unused]] const OutputTposeOp& output_op) { assert(problem.IsLayoutDefault()); @@ -1063,7 +1064,7 @@ template ConvSolution InitInvokerFactoryNHWC(const ExecutionContext&, - const ProblemDescriptionType& problem, + [[maybe_unused]] const ProblemDescriptionType& problem, const std::string& kernel_id) { ConvSolution result; @@ -1270,10 +1271,10 @@ ConvSolution InitInvokerFactoryWrwNCHW(const ExecutionContext& ctx, template ConvSolution -MakeSolutionGroupConvImplicitGemmXdlops(const miopen::conv::ProblemDescription& problem, - InvokerFactoryMakerNCHW&& invoker_factory_maker_ncdhw, - InvokerFactoryMakerNHWC&& invoker_factory_maker_ndhwc, - const bool use_tf32 = false) +MakeSolutionGroupConvImplicitGemmXdlops([[maybe_unused]] const miopen::conv::ProblemDescription& problem, + [[maybe_unused]] InvokerFactoryMakerNCHW&& invoker_factory_maker_ncdhw, + [[maybe_unused]] InvokerFactoryMakerNHWC&& invoker_factory_maker_ndhwc, + [[maybe_unused]] const bool use_tf32 = false) { #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL diff --git a/projects/miopen/src/env.cpp b/projects/miopen/src/env.cpp index 24892ce7b3dd..5ad2a88f1eb5 100644 --- a/projects/miopen/src/env.cpp +++ b/projects/miopen/src/env.cpp @@ -35,7 +35,9 @@ #include #ifdef _WIN32 +#ifndef WIN32_LEAN_AND_MEAN #define WIN32_LEAN_AND_MEAN +#endif #include #endif diff --git a/projects/miopen/src/include/miopen/batchnorm/common_spatial.hpp b/projects/miopen/src/include/miopen/batchnorm/common_spatial.hpp index af563db96a41..e7e59436a279 100644 --- a/projects/miopen/src/include/miopen/batchnorm/common_spatial.hpp +++ b/projects/miopen/src/include/miopen/batchnorm/common_spatial.hpp @@ -170,7 +170,7 @@ inline int GetStashMethod(bool IsLayoutNHWC, (in_cstride) % ylocalsize == 0 ? ylocalsize : (in_cstride) % ylocalsize; unsigned int last_zlocalsize = n % (zlocalsize * nelements) == 0 ? (zlocalsize * nelements) : n % (zlocalsize * nelements); - if(last_ylocalsize < stash_values && last_zlocalsize >= (size_t)stash_values) + if(last_ylocalsize < stash_values && last_zlocalsize >= static_cast(stash_values)) { stash_method = 1; } @@ -400,8 +400,8 @@ inline bool IsSpatialMultipleApplicable(const miopen::batchnorm::ProblemDescript // be large enough // - if C is not multiple of 2, intermediate results are stored in N dimension splitting // float values in group of 2 bytes. N must be large enough - if((!bfp32parm && (c % 2 != 0 && last_zlocalsize < (size_t)stash_values)) || - ((last_ylocalsize < stash_values) && (last_zlocalsize < (size_t)stash_values))) + if((!bfp32parm && (c % 2 != 0 && last_zlocalsize < static_cast(stash_values))) || + ((last_ylocalsize < stash_values) && (last_zlocalsize < static_cast(stash_values)))) { return false; } @@ -425,7 +425,7 @@ inline bool IsSpatialMultipleApplicable(const miopen::batchnorm::ProblemDescript // - if last block doesn't fit, intermediate results are stored in N dimension which must // be large enough stash_values *= (problem.GetXDesc().GetType() == miopenFloat ? 1 : 2); - if(last_ylocalsize < stash_values && last_zlocalsize < (size_t)stash_values) + if(last_ylocalsize < stash_values && last_zlocalsize < static_cast(stash_values)) { return false; } diff --git a/projects/miopen/src/include/miopen/conv/solvers.hpp b/projects/miopen/src/include/miopen/conv/solvers.hpp index e27ea412219f..bc7079e17e51 100644 --- a/projects/miopen/src/include/miopen/conv/solvers.hpp +++ b/projects/miopen/src/include/miopen/conv/solvers.hpp @@ -4829,18 +4829,18 @@ struct ConvWinogradNHWCTransposingTunableSolver struct TransposedConvBinWinograd3x3U final : ConvWinogradNHWCTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } }; struct TransposedConvBinWinogradRxS final : ConvWinogradNHWCTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } }; struct TransposedConvBinWinogradRxSf2x3g1 final : ConvWinogradNHWCTransposingSolver { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return GetSolverDbId(); } @@ -4850,7 +4850,7 @@ struct TransposedConvMPBidirectWinograd final : ConvWinogradNHWCTransposingSolver< ConvMPBidirectWinograd> { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return this->template GetSolverDbId< TransposedConvMPBidirectWinograd>(); @@ -4862,7 +4862,7 @@ struct TransposedConvWinograd3x3MultipassWrW final : ConvWinogradNHWCTransposingSolver< ConvWinograd3x3MultipassWrW> { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return this->template GetSolverDbId struct TransposedConvWinoFuryRxS final : ConvWinogradNHWCTransposingSolver> { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return this->template GetSolverDbId>(); } @@ -4885,7 +4885,7 @@ template struct TransposedConvWinoRageRxS final : ConvWinogradNHWCTransposingSolver> { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return this->template GetSolverDbId>(); } @@ -4907,7 +4907,7 @@ struct TransposedConvMPBidirectWinograd_xdlops final : ConvWinogradNHWCTransposingTunableSolver< ConvMPBidirectWinograd_xdlops> { - const std::string& SolverDbId() const + const std::string& SolverDbId() const override { return this->template GetSolverDbId #include #else diff --git a/projects/miopen/src/include/miopen/pooling/solvers.hpp b/projects/miopen/src/include/miopen/pooling/solvers.hpp index 38f9019eaa8b..2dd7c084d3c5 100644 --- a/projects/miopen/src/include/miopen/pooling/solvers.hpp +++ b/projects/miopen/src/include/miopen/pooling/solvers.hpp @@ -330,23 +330,23 @@ struct PoolingFwdNCHWTransposingSolver struct TransposedPoolingFwd2d final : PoolingFwdNCHWTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } PoolingForward2d::PerformanceConfigType GetDefaultPerformanceConfig(const ExecutionContext& ctx, - const miopen::pooling::ProblemDescription& problem) const + const miopen::pooling::ProblemDescription& problem) const override { return PoolingForward2d{}.GetDefaultPerformanceConfig(ctx, problem); } bool IsValidPerformanceConfig(const ExecutionContext& ctx, const miopen::pooling::ProblemDescription& problem, - const PoolingForward2d::PerformanceConfigType& config) const + const PoolingForward2d::PerformanceConfigType& config) const override { return PoolingForward2d{}.IsValidPerformanceConfig(ctx, problem, config); } PoolingForward2d::PerformanceConfigType Search(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, - const AnyInvokeParams& invoke_context) const + const AnyInvokeParams& invoke_context) const override { return PoolingForward2d{}.Search(context, problem, invoke_context); } @@ -354,23 +354,23 @@ struct TransposedPoolingFwd2d final : PoolingFwdNCHWTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } PoolingForwardNd::PerformanceConfigType GetDefaultPerformanceConfig(const ExecutionContext& ctx, - const miopen::pooling::ProblemDescription& problem) const + const miopen::pooling::ProblemDescription& problem) const override { return PoolingForwardNd{}.GetDefaultPerformanceConfig(ctx, problem); } bool IsValidPerformanceConfig(const ExecutionContext& ctx, const miopen::pooling::ProblemDescription& problem, - const PoolingForwardNd::PerformanceConfigType& config) const + const PoolingForwardNd::PerformanceConfigType& config) const override { return PoolingForwardNd{}.IsValidPerformanceConfig(ctx, problem, config); } PoolingForwardNd::PerformanceConfigType Search(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, - const AnyInvokeParams& invoke_context) const + const AnyInvokeParams& invoke_context) const override { return PoolingForwardNd{}.Search(context, problem, invoke_context); } @@ -474,23 +474,23 @@ struct PoolingBwdNCHWTransposingSolver struct TransposedPoolingBwd2d final : PoolingBwdNCHWTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } PoolingBackward2d::PerformanceConfigType GetDefaultPerformanceConfig(const ExecutionContext& ctx, - const miopen::pooling::ProblemDescription& problem) const + const miopen::pooling::ProblemDescription& problem) const override { return PoolingBackward2d{}.GetDefaultPerformanceConfig(ctx, problem); } bool IsValidPerformanceConfig(const ExecutionContext& ctx, const miopen::pooling::ProblemDescription& problem, - const PoolingBackward2d::PerformanceConfigType& config) const + const PoolingBackward2d::PerformanceConfigType& config) const override { return PoolingBackward2d{}.IsValidPerformanceConfig(ctx, problem, config); } PoolingBackward2d::PerformanceConfigType Search(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, - const AnyInvokeParams& invoke_context) const + const AnyInvokeParams& invoke_context) const override { return PoolingBackward2d{}.Search(context, problem, invoke_context); } @@ -498,23 +498,23 @@ struct TransposedPoolingBwd2d final : PoolingBwdNCHWTransposingSolver { - const std::string& SolverDbId() const { return GetSolverDbId(); } + const std::string& SolverDbId() const override { return GetSolverDbId(); } PoolingBackwardNd::PerformanceConfigType GetDefaultPerformanceConfig(const ExecutionContext& ctx, - const miopen::pooling::ProblemDescription& problem) const + const miopen::pooling::ProblemDescription& problem) const override { return PoolingBackwardNd{}.GetDefaultPerformanceConfig(ctx, problem); } bool IsValidPerformanceConfig(const ExecutionContext& ctx, const miopen::pooling::ProblemDescription& problem, - const PoolingBackwardNd::PerformanceConfigType& config) const + const PoolingBackwardNd::PerformanceConfigType& config) const override { return PoolingBackwardNd{}.IsValidPerformanceConfig(ctx, problem, config); } PoolingBackwardNd::PerformanceConfigType Search(const ExecutionContext& context, const miopen::pooling::ProblemDescription& problem, - const AnyInvokeParams& invoke_context) const + const AnyInvokeParams& invoke_context) const override { return PoolingBackwardNd{}.Search(context, problem, invoke_context); } diff --git a/projects/miopen/src/include/miopen/solver/implicitgemm_util.hpp b/projects/miopen/src/include/miopen/solver/implicitgemm_util.hpp index 3f4e3ecde5d7..e6d5cb6cb5e8 100644 --- a/projects/miopen/src/include/miopen/solver/implicitgemm_util.hpp +++ b/projects/miopen/src/include/miopen/solver/implicitgemm_util.hpp @@ -177,7 +177,7 @@ inline T igemm_get_max_gks(T gemm_k, T gemm_k_per_block, T max_log2_splits) return 0; T rem = gemm_k / gemm_k_per_block; T rem_pow2 = rem & (~(rem - 1)); - T gks = (T)log2(rem_pow2); + T gks = static_cast(log2(rem_pow2)); if(gks > max_log2_splits) gks = max_log2_splits; diff --git a/projects/miopen/src/kernels/hip_float8.hpp b/projects/miopen/src/kernels/hip_float8.hpp index 352cd97bb9b0..7d837be2a219 100644 --- a/projects/miopen/src/kernels/hip_float8.hpp +++ b/projects/miopen/src/kernels/hip_float8.hpp @@ -46,6 +46,13 @@ #define USE_SIMPLER_HIP_F8x8 0 +// Pull FP8 config from config.h when available (host compilation). +// During kernel compilation config.h is not in the include path; +// the macros are passed via -D flags from the solver instead. +#if __has_include() +#include +#endif + #ifndef MIOPEN_FP8_CLIPPING #define MIOPEN_FP8_CLIPPING 1 #endif diff --git a/projects/miopen/src/kernels/miopen_rocrand.hpp b/projects/miopen/src/kernels/miopen_rocrand.hpp index 3effb18802e2..019e14193be7 100644 --- a/projects/miopen/src/kernels/miopen_rocrand.hpp +++ b/projects/miopen/src/kernels/miopen_rocrand.hpp @@ -44,12 +44,19 @@ #define ROCRAND_DETAIL_XORWOW_BM_NOT_IN_STATE // Use inlined rocrand header for runtime compilation to avoid external dependency -// For host code compilation, use the regular rocrand header +// For host code compilation, use the regular rocrand header. +// Suppress warnings from rocrand's own headers: in TheRock builds rocrand/stage/include +// is injected globally as -I (not -isystem), so rocrand headers are not treated as system +// includes and their internal issues trigger our warnings. +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wold-style-cast" +#pragma clang diagnostic ignored "-Wundef" #ifdef MIOPEN_HIP_RUNTIME_COMPILE #include "rocrand_xorwow_inlined.h" #else #include #endif +#pragma clang diagnostic pop namespace prng { diff --git a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index d44ab5ba59fa..4a30062d10c2 100644 --- a/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/projects/miopen/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -873,7 +873,7 @@ bool PerformanceConfigAsmImplicitGemmGTCBwdXdlopsNHWC::IsValid( n = n | (n >> 16); return n - (n >> 1); }(padded_k_num); - int k_pow2 = (int)log2(prev_pow2); + int k_pow2 = static_cast(log2(prev_pow2)); return std::min(k_pow2, BWD_MAX_GEMM_K_SPLITS); } diff --git a/projects/miopen/test/conv_common.hpp b/projects/miopen/test/conv_common.hpp index 5e141908bab7..87f116690296 100644 --- a/projects/miopen/test/conv_common.hpp +++ b/projects/miopen/test/conv_common.hpp @@ -2272,14 +2272,14 @@ struct conv_driver : test_driver { auto output = get_output_tensor(filter, input, weights, out_layout); - auto gen_positive_value = [=](auto...) { + auto gen_positive_value = [=, this](auto...) { auto data_type = input.desc.GetType(); int v_max = is_int8 ? 16 : (data_type == miopenHalf) ? 4 : 17; return gen_float ? prng::gen_canonical() : static_cast(prng::gen_A_to_B(1, v_max)); }; - auto gen_sign_value = [=](auto... is) { + auto gen_sign_value = [=, this](auto... is) { auto data_type = input.desc.GetType(); int v_max = is_int8 ? 16 : (data_type == miopenHalf) ? 4 : 17; return gen_float ? prng::gen_A_to_B(-1.0, 1.0) diff --git a/projects/miopen/test/driver.hpp b/projects/miopen/test/driver.hpp index 13097aaca388..53b08f24492e 100644 --- a/projects/miopen/test/driver.hpp +++ b/projects/miopen/test/driver.hpp @@ -112,7 +112,7 @@ struct test_driver template void add_source(Source src, T& x) { - data_sources.push_back([=, &x](std::function callback) { + data_sources.push_back([=, this, &x](std::function callback) { for(auto y : src()) // NOLINT { x = T(y); @@ -332,7 +332,7 @@ struct test_driver template generate_tensor_t generate_tensor(std::set dims, X single, G g) { - return {[=]() -> std::set { + return {[=, this]() -> std::set { if(full_set) return dims; else @@ -351,7 +351,7 @@ struct test_driver template auto lazy_generate_tensor(F f, G g) -> generate_tensor_t, G> { - return {[=]() -> decltype(f()) { + return {[=, this]() -> decltype(f()) { if(full_set) return f(); else @@ -363,7 +363,7 @@ struct test_driver template generate_tensor_t lazy_generate_tensor(F f, X single, G g) { - return {[=]() -> std::set { + return {[=, this]() -> std::set { if(full_set) return f(); else @@ -382,7 +382,7 @@ struct test_driver template generate_tensor_t, G> get_tensor(F gen_shapes, G gen_value) { - return lazy_generate_tensor([=] { return gen_shapes(batch_factor); }, gen_value); + return lazy_generate_tensor([=, this] { return gen_shapes(batch_factor); }, gen_value); } template @@ -390,7 +390,7 @@ struct test_driver get_bn_spatial_input_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { return lazy_generate_tensor( - [=] { return get_bn_spatial_inputs(batch_factor); }, {4, 64, 28, 28}, tensor_elem_gen); + [=, this] { return get_bn_spatial_inputs(batch_factor); }, {4, 64, 28, 28}, tensor_elem_gen); } template @@ -398,7 +398,7 @@ struct test_driver get_bn_peract_input_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { return lazy_generate_tensor( - [=] { return get_bn_peract_inputs(batch_factor); }, {16, 32, 8, 8}, tensor_elem_gen); + [=, this] { return get_bn_peract_inputs(batch_factor); }, {16, 32, 8, 8}, tensor_elem_gen); } template @@ -406,14 +406,14 @@ struct test_driver get_input_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { return lazy_generate_tensor( - [=] { return get_inputs(batch_factor); }, {16, 32, 8, 8}, tensor_elem_gen); + [=, this] { return get_inputs(batch_factor); }, {16, 32, 8, 8}, tensor_elem_gen); } template generate_tensor_t, G> get_3d_bn_spatial_input_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { - return lazy_generate_tensor([=] { return get_3d_bn_spatial_inputs(batch_factor); }, + return lazy_generate_tensor([=, this] { return get_3d_bn_spatial_inputs(batch_factor); }, {16, 32, 8, 8, 8}, tensor_elem_gen); } @@ -422,7 +422,7 @@ struct test_driver generate_tensor_t, G> get_3d_bn_peract_input_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { - return lazy_generate_tensor([=] { return get_3d_bn_peract_inputs(batch_factor); }, + return lazy_generate_tensor([=, this] { return get_3d_bn_peract_inputs(batch_factor); }, {16, 32, 8, 8, 8}, tensor_elem_gen); } @@ -432,7 +432,7 @@ struct test_driver get_weights_tensor(G tensor_elem_gen = tensor_elem_gen_integer{}) { return lazy_generate_tensor( - [=] { return get_weights(batch_factor); }, {64, 32, 5, 5}, tensor_elem_gen); + [=, this] { return get_weights(batch_factor); }, {64, 32, 5, 5}, tensor_elem_gen); } template @@ -456,7 +456,7 @@ struct test_driver template generate_data_t> generate_data(std::vector dims, T single) { - return {[=]() -> std::vector { + return {[=, this]() -> std::vector { if(full_set) return dims; else @@ -468,7 +468,7 @@ struct test_driver generate_data_t> generate_data_limited(std::vector dims, int limit_multiplier, T single) { - return {[=]() -> std::vector { + return {[=, this]() -> std::vector { if(full_set) { if(limit_set > 0) @@ -504,7 +504,7 @@ struct test_driver template generate_data_t> generate_data(std::vector dims) { - return {[=]() -> std::vector { + return {[=, this]() -> std::vector { if(full_set) return dims; else @@ -515,13 +515,13 @@ struct test_driver template generate_data_t> generate_multi_data(std::vector> multi_dims) { - return {[=]() -> std::vector { return generate_data(multi_dims.at(dataset_id))(T{}); }}; + return {[=, this]() -> std::vector { return generate_data(multi_dims.at(dataset_id))(T{}); }}; } template generate_data_t> generate_data_limited(std::vector dims, int limit_multiplier) { - return {[=]() -> std::vector { + return {[=, this]() -> std::vector { if(full_set) { if(limit_set > 0) @@ -547,7 +547,7 @@ struct test_driver generate_data_t> generate_multi_data_limited(std::vector> multi_dims, int limit_multiplier) { - return {[=]() -> std::vector { + return {[=, this]() -> std::vector { return generate_data_limited(multi_dims.at(dataset_id), limit_multiplier)(T{}); }}; } @@ -555,7 +555,7 @@ struct test_driver template auto lazy_generate_data(F f, T single) -> generate_data_t { - return {[=]() -> decltype(f()) { + return {[=, this]() -> decltype(f()) { if(full_set) return f(); else @@ -566,7 +566,7 @@ struct test_driver template auto lazy_generate_data(F f) -> generate_data_t { - return {[=]() -> decltype(f()) { + return {[=, this]() -> decltype(f()) { if(full_set) return f(); else @@ -577,7 +577,7 @@ struct test_driver template generate_data_t> generate_single(T single) { - return {[=]() -> std::vector { return {single}; }}; + return {[=, this]() -> std::vector { return {single}; }}; } template @@ -613,7 +613,7 @@ struct test_driver auto verify_reporter() { - return [=](bool pass, + return [=, this](bool pass, std::vector error, const auto& out_cpu, const auto& out_gpu, diff --git a/projects/miopen/test/gtest/kernel_tuning_net.cpp b/projects/miopen/test/gtest/kernel_tuning_net.cpp index 304adb9800d4..ab103db57853 100644 --- a/projects/miopen/test/gtest/kernel_tuning_net.cpp +++ b/projects/miopen/test/gtest/kernel_tuning_net.cpp @@ -197,7 +197,7 @@ template class KernelTuningNetTest : public ::testing::TestWithParam { protected: - void TestParameterPredictionModel(std::string solver_nm) + void TestParameterPredictionModel([[maybe_unused]] std::string solver_nm) { #if MIOPEN_ENABLE_AI_KERNEL_TUNING && MIOPEN_USE_COMPOSABLEKERNEL auto test_case = GetParam(); diff --git a/projects/miopen/test/gtest/kthvalue.hpp b/projects/miopen/test/gtest/kthvalue.hpp index 2aa7e6fd41d1..62d5ebb5905a 100644 --- a/projects/miopen/test/gtest/kthvalue.hpp +++ b/projects/miopen/test/gtest/kthvalue.hpp @@ -156,7 +156,7 @@ struct KthvalueFwdTest : public ::testing::TestWithParam output.desc, output_dev.get(), indicesDesc, - (size_t*)indices_dev.get(), + static_cast(indices_dev.get()), config.k, config.dim, config.keepDim); diff --git a/projects/miopen/test/gtest/na_inference_find2.cpp b/projects/miopen/test/gtest/na_inference_find2.cpp index 0313097c2242..fa45d8eddcc8 100644 --- a/projects/miopen/test/gtest/na_inference_find2.cpp +++ b/projects/miopen/test/gtest/na_inference_find2.cpp @@ -152,7 +152,7 @@ struct verify_inference_batchnorm_activ miopenTensorArgument_t{miopenTensorBatchnormBias, nullptr, bnbias_dev.get()}, miopenTensorArgument_t{miopenTensorBatchnormEstimatedMean, nullptr, estMean_dev.get()}, miopenTensorArgument_t{miopenTensorBatchnormEstimatedVariance, nullptr, estVariance_dev.get()}, - miopenTensorArgument_t{miopenScalarBatchnormEpsilon, nullptr, (void*)&mut_eps}, + miopenTensorArgument_t{miopenScalarBatchnormEpsilon, nullptr, static_cast(&mut_eps)}, // clang-format on }; diff --git a/projects/miopen/test/gtest/rnn_seq_api.hpp b/projects/miopen/test/gtest/rnn_seq_api.hpp index 4f7cded37687..6699a726bd69 100644 --- a/projects/miopen/test/gtest/rnn_seq_api.hpp +++ b/projects/miopen/test/gtest/rnn_seq_api.hpp @@ -1823,7 +1823,7 @@ struct RNNSeqApiCommon : public ::testing::TestWithParam seqTensor dy(output); const auto num_hidden_layers = numLayers * ((dirMode != 0) ? 2 : 1); - tensor hx = [=]() { + tensor hx = [=, this]() { if(pytorchTensorDescriptorFormat) return tensor(std::vector{num_hidden_layers, batchSize, hiddenSize, 1, 1}); else