From c24e5284814b6f1204961c3826a57b7933ec1b12 Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Wed, 27 May 2026 06:56:58 -0700 Subject: [PATCH] [rocm-libraries] ROCm/rocm-libraries#7760 (commit a61bc76) [CK] suppress compiler warnings while building pytorch. (#7760) ## Motivation Recently added compiler flags that are required to suppress false warnings by latest staging compiler are not recognized by older compiler versions and are triggering an avalanche of warnings. Previous attempt to suppress them by using -Wno-unknown-warning-option flag didn't help, because that flag wasn't recognized either and just added more warnings. I've verified that current approach by checking the clang version actually works as intended and makes the warnings go away. ## Technical Details ## Test Plan ## Test Result ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --- CMakeLists.txt | 5 +++++ codegen/include/ck/host/types.hpp | 5 ++++- codegen/src/headers.cpp | 5 ++++- codegen/test/include/test.hpp | 3 +-- .../grouped_gemm_multiple_d_xdl_fp16.cpp | 6 ++++-- .../grouped_gemm_bias_e_permute_xdl_fp16.cpp | 6 ++++-- ...un_batched_gemm_bias_e_permute_example.inc | 6 ++++-- .../sparse_embedding3_forward_layernorm.cpp | 6 ++++-- example/39_permute/common.hpp | 6 ++++-- .../splitk_gemm_bias_e_permute_xdl_fp16.cpp | 6 ++++-- .../splitk_gemm_bias_e_permute_xdl_fp32.cpp | 6 ++++-- .../01_fmha/codegen/ops/fmha_fwd_splitkv.py | 6 ++++-- .../01_fmha/fmha_fwd_head_grouping.hpp | 5 ++++- example/ck_tile/01_fmha/fmha_fwd_runner.hpp | 5 ++++- example/ck_tile/01_fmha/quant.hpp | 6 ++++-- .../03_gemm/gemm_splitk_two_stage_reduce.cpp | 6 ++++-- .../grouped_convolution_utils.hpp | 6 ++++-- .../builder/test/unit_device_buffer.cpp | 5 ++++- .../ck/library/utility/gpu_verification.hpp | 5 ++++- include/ck/library/utility/host_tensor.hpp | 6 ++++-- include/ck/tensor/static_tensor.hpp | 5 ++++- .../multi_index_transform.hpp | 6 ++++-- .../tensor_description/tensor_descriptor.hpp | 5 ++++- .../tensor_space_filling_curve.hpp | 6 ++++-- .../blockwise_gemm_pipeline_wmmaops_base.hpp | 5 ++++- .../block/blockwise_gemm_pipeline_xdlops.hpp | 5 ++++- .../blockwise_gemm_pipeline_xdlops_base.hpp | 6 ++++-- .../gpu/block/blockwise_gemm_wmma.hpp | 6 ++++-- .../gpu/block/blockwise_gemm_xdlops.hpp | 5 ++++- .../blockwise_gemm_xdlops_skip_b_lds.hpp | 6 ++++-- .../gpu/device/device_grouped_gemm.hpp | 6 ++++-- ...ontraction_multiple_d_wmma_cshuffle_v3.hpp | 5 ++++- ...ed_contraction_multiple_d_xdl_cshuffle.hpp | 5 ++++- ...ched_gemm_gemm_wmma_cshuffle_v3_common.hpp | 11 +++++----- .../device_batched_gemm_gemm_xdl_cshuffle.hpp | 6 ++++-- .../impl/device_batched_gemm_multi_d_xdl.hpp | 5 ++++- ...e_batched_gemm_reduce_wmma_cshuffle_v3.hpp | 6 ++++-- ...evice_batched_gemm_reduce_xdl_cshuffle.hpp | 6 ++++-- ...emm_softmax_gemm_permute_wmma_cshuffle.hpp | 6 ++++-- ...gemm_softmax_gemm_permute_xdl_cshuffle.hpp | 6 ++++-- ...batched_gemm_softmax_gemm_xdl_cshuffle.hpp | 6 ++++-- .../device/impl/device_batched_gemm_xdl.hpp | 6 ++++-- .../impl/device_batchnorm_backward_impl.hpp | 6 ++++-- .../impl/device_batchnorm_forward_impl.hpp | 6 ++++-- ...device_batchnorm_forward_impl_obsolete.hpp | 6 ++++-- .../impl/device_cgemm_4gemm_xdl_cshuffle.hpp | 6 ++++-- ...fle_bias_activation_add_nhwc_kyxc_nhwk.hpp | 6 ++++-- ...shuffle_bias_activation_nhwc_kyxc_nhwk.hpp | 6 ++++-- ...onv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp | 6 ++++-- .../device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp | 6 ++++-- ...ice_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp | 6 ++++-- .../device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp | 6 ++++-- ...device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp | 6 ++++-- ..._gemm_bias_add_reduce_wmma_cshuffle_v3.hpp | 6 ++++-- ...vice_gemm_bias_add_reduce_xdl_cshuffle.hpp | 6 ++++-- ..._multiple_d_layernorm_wmma_cshuffle_v3.hpp | 5 ++++- ...emm_multiple_d_multiple_r_xdl_cshuffle.hpp | 5 ++++- ...evice_gemm_multiple_d_wmma_cshuffle_v3.hpp | 5 ++++- .../device_gemm_multiple_d_xdl_cshuffle.hpp | 5 ++++- .../device_gemm_reduce_wmma_cshuffle_v3.hpp | 6 ++++-- .../impl/device_gemm_reduce_xdl_cshuffle.hpp | 6 ++++-- .../impl/device_gemm_wmma_cshuffle_v3r1.hpp | 6 ++++-- .../device_gemm_xdl_layernorm_cshuffle.hpp | 6 ++++-- .../impl/device_gemm_xdl_skip_b_lds.hpp | 6 ++++-- .../device_gemm_xdl_waveletmodel_cshuffle.hpp | 6 ++++-- .../device_grouped_conv_bwd_weight_dl.hpp | 6 ++++-- ...conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp | 5 ++++- ...ped_conv_fwd_multiple_abd_xdl_cshuffle.hpp | 5 ++++- ..._conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp | 5 ++++- ...fwd_multiple_d_multiple_r_xdl_cshuffle.hpp | 5 ++++- ...uped_conv_fwd_multiple_d_wmma_cshuffle.hpp | 5 ++++- ...e_grouped_query_attention_forward_wmma.hpp | 6 ++++-- .../impl/device_image_to_column_impl.hpp | 6 ++++-- ...ice_multi_query_attention_forward_wmma.hpp | 6 ++++-- .../device_multiple_reduce_multiblock.hpp | 5 ++++- .../device_multiple_reduce_threadwise.hpp | 5 ++++- .../device_normalization_bwd_data_impl.hpp | 6 ++++-- ...vice_normalization_bwd_gamma_beta_impl.hpp | 6 ++++-- .../impl/device_normalization_fwd_impl.hpp | 6 ++++-- .../gpu/device/impl/device_permute_impl.hpp | 5 ++++- .../impl/device_pool2d_fwd_nhwc_nhwc.hpp | 6 ++++-- .../impl/device_pool3d_fwd_ndhwc_ndhwc.hpp | 6 ++++-- .../device/impl/device_put_element_impl.hpp | 6 ++++-- .../device/impl/device_reduce_threadwise.hpp | 6 ++++-- .../impl/device_reduce_threadwise_multi_d.hpp | 6 ++++-- ...ce_sparse_embeddings_forward_layernorm.hpp | 6 ++++-- ...tk_contraction_multiple_d_xdl_cshuffle.hpp | 5 ++++- .../epilogue_cshuffle_v3_welford_wmma.hpp | 6 ++++-- ...ise_batched_gemm_gemm_wmma_cshuffle_v3.hpp | 7 ++++--- .../gpu/grid/gridwise_gemm_dpp.hpp | 6 ++++-- .../grid/gridwise_gemm_wmma_cshuffle_v3.hpp | 6 ++++-- ...ridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp | 6 ++++-- .../gridwise_gemm_xdl_cshuffle_conv_v3.hpp | 6 ++++-- .../gridwise_gemm_xdl_cshuffle_streamk_v3.hpp | 6 ++++-- .../grid/gridwise_gemm_xdl_cshuffle_v1.hpp | 6 ++++-- .../grid/gridwise_gemm_xdl_cshuffle_v2.hpp | 6 ++++-- .../grid/gridwise_gemm_xdl_cshuffle_v3.hpp | 6 ++++-- ...wise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp | 6 ++++-- .../gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp | 6 ++++-- ...ridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp | 5 ++++- .../gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp | 6 ++++-- ..._gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp | 6 ++++-- ...m_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp | 6 ++++-- ...fle_v3_multi_d_blockscale_b_preshuffle.hpp | 6 ++++-- .../grid/gridwise_gemm_xdl_cshuffle_v3_mx.hpp | 6 ++++-- ...se_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp | 6 ++++-- .../grid/gridwise_gemm_xdlops_bwd_weight.hpp | 6 ++++-- ...ise_gemm_xdlops_splitk_lds_direct_load.hpp | 6 ++++-- .../gpu/grid/gridwise_gemm_xdlops_streamk.hpp | 6 ++++-- .../gpu/grid/gridwise_gemm_xdlops_v2r3.hpp | 6 ++++-- .../gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp | 6 ++++-- .../gpu/grid/gridwise_moe_gemm.hpp | 6 ++++-- .../gpu/grid/gridwise_moe_gemm_blockscale.hpp | 6 ++++-- .../gpu/grid/gridwise_moe_mx_gemm.hpp | 6 ++++-- .../gpu/grid/gridwise_moe_mx_gemm_bns.hpp | 6 ++++-- .../grid/gridwise_moe_mx_gemm_bpreshuffle.hpp | 6 ++++-- .../threadwise_tensor_slice_transfer.hpp | 6 ++++-- include/ck/utility/amd_address_space.hpp | 10 ++++++++-- include/ck/utility/amd_cluster_load.hpp | 4 ++++ include/ck/utility/array.hpp | 6 ++++-- include/ck/utility/c_style_pointer_cast.hpp | 5 ++++- include/ck/utility/dtype_vector.hpp | 5 ++++- include/ck/utility/dynamic_buffer.hpp | 6 ++++-- include/ck/utility/env.hpp | 10 ++++++++-- include/ck/utility/functional2.hpp | 6 ++++-- include/ck/utility/sequence.hpp | 6 ++++-- include/ck/utility/span.hpp | 6 ++++-- include/ck/utility/static_buffer.hpp | 5 ++++- .../ck/utility/statically_indexed_array.hpp | 6 ++++-- include/ck/utility/tuple.hpp | 6 ++++-- include/ck/utility/type_convert.hpp | 5 ++++- include/ck/utility/workgroup_barrier.hpp | 6 ++++-- include/ck/wrapper/layout.hpp | 6 ++++-- include/ck/wrapper/tensor.hpp | 6 ++++-- .../core/algorithm/coordinate_transform.hpp | 6 ++++-- .../core/arch/amd_buffer_addressing.hpp | 17 ++++++++++++---- .../arch/amd_buffer_addressing_builtins.hpp | 17 ++++++++++++---- .../ck_tile/core/arch/amd_cluster_load.hpp | 8 ++++++++ include/ck_tile/core/arch/arch.hpp | 10 ++++++++-- include/ck_tile/core/arch/mma/amdgcn_mma.hpp | 6 ++++-- .../ck_tile/core/arch/mma/mma_pipeline.hpp | 6 ++++-- include/ck_tile/core/container/array.hpp | 6 ++++-- include/ck_tile/core/container/map.hpp | 6 ++++-- include/ck_tile/core/container/span.hpp | 6 ++++-- .../ck_tile/core/container/static_array.hpp | 6 ++++-- .../ck_tile/core/container/thread_buffer.hpp | 6 ++++-- include/ck_tile/core/container/tuple.hpp | 6 ++++-- include/ck_tile/core/numeric/e4m3.hpp | 6 ++++-- include/ck_tile/core/numeric/e5m3.hpp | 6 ++++-- include/ck_tile/core/numeric/e8m0.hpp | 6 ++++-- include/ck_tile/core/numeric/pk_fp4.hpp | 6 ++++-- include/ck_tile/core/numeric/type_convert.hpp | 5 ++++- include/ck_tile/core/tensor/buffer_view.hpp | 6 ++++-- .../core/tensor/static_distributed_tensor.hpp | 6 ++++-- .../ck_tile/core/tensor/tensor_adaptor.hpp | 6 ++++-- .../core/tensor/tensor_adaptor_coordinate.hpp | 11 +++++++--- include/ck_tile/core/tensor/tensor_view.hpp | 6 ++++-- .../ck_tile/core/tensor/tile_distribution.hpp | 6 ++++-- .../ck_tile/core/tensor/tile_elementwise.hpp | 5 ++++- include/ck_tile/core/utility/debug.hpp | 5 ++++- include/ck_tile/core/utility/env.hpp | 10 ++++++++-- include/ck_tile/core/utility/functional.hpp | 5 ++++- include/ck_tile/core/utility/type_traits.hpp | 5 ++++- include/ck_tile/host/arg_parser.hpp | 6 ++++-- include/ck_tile/host/host_tensor.hpp | 6 ++++-- include/ck_tile/host/joinable_thread.hpp | 5 ++++- include/ck_tile/host/kernel_launch.hpp | 6 ++++-- .../kernel/batched_contraction_kernel.hpp | 6 ++++-- .../flatmm_32x512x128_1x4x1_16x16x32.hpp | 20 +++++++++++++++---- .../flatmm_sn_32x128x512_1x4x1_16x16x32.hpp | 10 ++++++++-- ...latmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp | 10 ++++++++-- .../ops/flatmm/kernel/flatmm_kernel.hpp | 6 ++++-- .../flatmm/kernel/grouped_flatmm_kernel.hpp | 6 ++++-- .../ops/flatmm/kernel/moe_flatmm_kernel.hpp | 6 ++++-- .../ck_tile/ops/fmha/block/cast_tile_mx.hpp | 10 ++++++++-- .../ops/fmha/block/page_block_navigator.hpp | 6 ++++-- include/ck_tile/ops/fmha/block/variants.hpp | 6 ++++-- .../fused_moe/kernel/moe_sorting_kernel.hpp | 6 ++++-- .../ck_tile/ops/gemm/kernel/gemm_kernel.hpp | 6 ++++-- .../ops/gemm/kernel/gemm_multi_abd_kernel.hpp | 6 ++++-- .../ops/gemm/kernel/gemm_multi_d_kernel.hpp | 6 ++++-- .../ops/gemm/kernel/grouped_gemm_kernel.hpp | 6 ++++-- .../ops/gemm_mx/kernel/scale_pointer.hpp | 6 ++++-- .../block/block_gemm_quant_common.hpp | 6 ++++-- .../gemm_quant/kernel/gemm_quant_kernel.hpp | 6 ++++-- .../kernel/grouped_gemm_quant_kernel.hpp | 6 ++++-- .../utils/grouped_convolution_utils.hpp | 6 ++++-- .../ops/pooling/kernel/pool_kernel.hpp | 6 ++++-- include/ck_tile/ref/naive_attention.hpp | 5 ++++- include/rapidjson/document.h | 2 +- include/rapidjson/encodedstream.h | 4 ++-- include/rapidjson/error/en.h | 4 ++-- include/rapidjson/error/error.h | 4 ++-- include/rapidjson/filereadstream.h | 4 ++-- include/rapidjson/filewritestream.h | 4 ++-- include/rapidjson/internal/diyfp.h | 4 ++-- include/rapidjson/internal/regex.h | 2 +- include/rapidjson/istreamwrapper.h | 2 +- include/rapidjson/memorystream.h | 4 ++-- include/rapidjson/ostreamwrapper.h | 4 ++-- include/rapidjson/pointer.h | 2 +- include/rapidjson/reader.h | 2 +- include/rapidjson/schema.h | 2 +- include/rapidjson/writer.h | 2 +- .../cpu/reference_avgpool_bwd.hpp | 6 ++++-- .../cpu/reference_batched_gemm.hpp | 6 ++++-- .../cpu/reference_cgemm.hpp | 6 ++++-- .../cpu/reference_column_to_image.hpp | 6 ++++-- .../cpu/reference_contraction.hpp | 6 ++++-- .../cpu/reference_conv_bwd_data.hpp | 6 ++++-- .../cpu/reference_conv_bwd_weight.hpp | 6 ++++-- .../cpu/reference_conv_fwd.hpp | 6 ++++-- .../reference_conv_fwd_bias_activation.hpp | 6 ++++-- ...reference_conv_fwd_bias_activation_add.hpp | 6 ++++-- .../cpu/reference_elementwise.hpp | 6 ++++-- .../cpu/reference_fpAintB_gemm.hpp | 6 ++++-- .../cpu/reference_gemm.hpp | 6 ++++-- .../cpu/reference_gemm_layernorm.hpp | 6 ++++-- .../cpu/reference_gemm_multi_abd.hpp | 6 ++++-- .../cpu/reference_gemm_multiple_d.hpp | 6 ++++-- .../cpu/reference_groupnorm.hpp | 6 ++++-- .../cpu/reference_groupnorm_bwd.hpp | 6 ++++-- .../cpu/reference_image_to_column.hpp | 6 ++++-- .../cpu/reference_layernorm.hpp | 6 ++++-- .../cpu/reference_layernorm_bwd.hpp | 6 ++++-- .../cpu/reference_maxpool_bwd.hpp | 6 ++++-- .../cpu/reference_moe_gemm.hpp | 6 ++++-- .../cpu/reference_moe_gemm1_blockscale.hpp | 6 ++++-- .../reference_moe_gemm1_blockscale_splitk.hpp | 6 ++++-- .../cpu/reference_moe_gemm2.hpp | 6 ++++-- .../cpu/reference_moe_gemm2_blockscale.hpp | 6 ++++-- .../cpu/reference_moe_mx_gemm1.hpp | 6 ++++-- .../cpu/reference_moe_mx_gemm2.hpp | 6 ++++-- .../cpu/reference_mx_gemm.hpp | 6 ++++-- .../cpu/reference_pool_fwd.hpp | 6 ++++-- ...ce_sparse_embedding3_forward_layernorm.hpp | 6 ++++-- .../gpu/reference_gemm.hpp | 5 ++++- ...gelu_bf16_i8_bf16_mk_kn_mn_v1_instance.cpp | 6 ++++-- .../profiler/profile_contraction_utils.hpp | 5 ++++- profiler/src/profiler_operation_registry.hpp | 6 ++++-- test/ck_tile/flatmm/test_mx_flatmm_base.hpp | 5 ++++- test/common/csv_test_loader.hpp | 5 ++++- ...st_grouped_convnd_bwd_data_dataset_xdl.cpp | 5 ++++- ..._grouped_convnd_bwd_weight_dataset_xdl.cpp | 5 ++++- .../test_grouped_convnd_fwd_dataset_xdl.cpp | 5 ++++- .../test_grouped_gemm_fixed_nk_bias.cpp | 5 ++++- .../test_grouped_gemm_multi_abd_fixed_nk.cpp | 5 ++++- test/grouped_gemm/test_grouped_gemm_util.hpp | 6 ++++-- .../test_grouped_gemm_tile_loop_util.hpp | 5 ++++- .../position_embedding/position_embedding.cpp | 6 ++++-- .../gemm_multi_d/gemm_multi_d_benchmark.hpp | 6 ++++-- .../gemm_preshuffle_benchmark.hpp | 6 ++++-- .../gemm_universal_benchmark.hpp | 5 ++++- .../grouped_gemm/grouped_gemm_benchmark.hpp | 5 ++++- .../gemm_streamk/gemm_streamk_benchmark.hpp | 6 ++++-- 255 files changed, 1049 insertions(+), 457 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e2f3c46619..2056f23a7f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -135,6 +135,11 @@ add_compile_options(-Wno-bit-int-extension) add_compile_options(-Wno-pass-failed) add_compile_options(-Wno-switch-default) add_compile_options(-Wno-unique-object-duplication) +if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang") + add_compile_options(-Wno-unknown-warning-option) +else() + add_compile_options(-Wno-unknown-pragmas) +endif() # Increase the number of max elements in fold expressions add_compile_options(-fbracket-depth=1024) diff --git a/codegen/include/ck/host/types.hpp b/codegen/include/ck/host/types.hpp index 4db3941414..91dd3013a1 100644 --- a/codegen/include/ck/host/types.hpp +++ b/codegen/include/ck/host/types.hpp @@ -90,12 +90,15 @@ std::string SequenceStr(const std::vector& v); std::string MakeTuple(const std::vector& v); +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wglobal-constructors" +#endif template const std::string S = SequenceStr({xs...}); +#ifdef __clang__ #pragma clang diagnostic pop +#endif constexpr const char* PassThrough = "ck::tensor_operation::element_wise::PassThrough"; constexpr const char* Bilinear = "ck::tensor_operation::element_wise::Bilinear"; diff --git a/codegen/src/headers.cpp b/codegen/src/headers.cpp index 9d59176c92..0929879e6a 100644 --- a/codegen/src/headers.cpp +++ b/codegen/src/headers.cpp @@ -7,11 +7,14 @@ namespace ck { namespace host { +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wglobal-constructors" +#endif const std::string config_header = ""; +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif std::unordered_map GetHeaders() { diff --git a/codegen/test/include/test.hpp b/codegen/test/include/test.hpp index b799df372d..eb9d29187a 100644 --- a/codegen/test/include/test.hpp +++ b/codegen/test/include/test.hpp @@ -843,9 +843,8 @@ inline void run(int argc, const char* argv[]) TEST_CASE_REGISTER(__VA_ARGS__) \ void __VA_ARGS__() -#ifdef __clang__ +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wglobal-constructors" #endif diff --git a/example/15_grouped_gemm/grouped_gemm_multiple_d_xdl_fp16.cpp b/example/15_grouped_gemm/grouped_gemm_multiple_d_xdl_fp16.cpp index ae097865f6..b24e9c089d 100644 --- a/example/15_grouped_gemm/grouped_gemm_multiple_d_xdl_fp16.cpp +++ b/example/15_grouped_gemm/grouped_gemm_multiple_d_xdl_fp16.cpp @@ -23,10 +23,10 @@ #include "ck/library/utility/literals.hpp" #include "ck/library/reference_tensor_operation/cpu/reference_gemm_multiple_d.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::hip_check_error; using ::ck::HostTensorDescriptor; @@ -79,4 +79,6 @@ using DeviceGemmInstance = int main(int argc, char* argv[]) { return !run_grouped_gemm_example(argc, argv); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/28_grouped_gemm_bias_e_permute/grouped_gemm_bias_e_permute_xdl_fp16.cpp b/example/28_grouped_gemm_bias_e_permute/grouped_gemm_bias_e_permute_xdl_fp16.cpp index 63417dfd95..2e0d9c2299 100644 --- a/example/28_grouped_gemm_bias_e_permute/grouped_gemm_bias_e_permute_xdl_fp16.cpp +++ b/example/28_grouped_gemm_bias_e_permute/grouped_gemm_bias_e_permute_xdl_fp16.cpp @@ -18,10 +18,10 @@ #include "ck/library/utility/host_tensor_generator.hpp" #include "ck/library/utility/numeric.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::HostTensorDescriptor; using ::ck::make_ParallelTensorFunctor; @@ -482,4 +482,6 @@ int main(int argc, char* argv[]) return pass ? 0 : 1; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/29_batched_gemm_bias_e_permute/run_batched_gemm_bias_e_permute_example.inc b/example/29_batched_gemm_bias_e_permute/run_batched_gemm_bias_e_permute_example.inc index 081e27c843..f6134fab5a 100644 --- a/example/29_batched_gemm_bias_e_permute/run_batched_gemm_bias_e_permute_example.inc +++ b/example/29_batched_gemm_bias_e_permute/run_batched_gemm_bias_e_permute_example.inc @@ -1,7 +1,7 @@ +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // hardcoded for NumDimM == NumDimN == NumDimK == 2 template = 23 #pragma clang diagnostic pop +#endif diff --git a/example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp b/example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp index 8c6121c402..3f1dafe3a7 100644 --- a/example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp +++ b/example/36_sparse_embedding/sparse_embedding3_forward_layernorm.cpp @@ -19,10 +19,10 @@ #include "ck/library/utility/host_tensor_generator.hpp" #include "ck/library/reference_tensor_operation/cpu/reference_sparse_embedding3_forward_layernorm.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::HostTensorDescriptor; using ::ck::Tensor; @@ -222,4 +222,6 @@ int main(int argc, char* argv[]) return 0; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/39_permute/common.hpp b/example/39_permute/common.hpp index 4478118f34..c6c1087ebe 100644 --- a/example/39_permute/common.hpp +++ b/example/39_permute/common.hpp @@ -26,10 +26,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::HostTensorDescriptor; using ::ck::Tensor; @@ -463,4 +463,6 @@ auto host_permute(const Tensor& src, const Axes& axes, Functor functor, Ten return true; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp16.cpp b/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp16.cpp index d14e4e5c32..b537c59835 100644 --- a/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp16.cpp +++ b/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp16.cpp @@ -16,10 +16,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::HostTensorDescriptor; using ::ck::make_ParallelTensorFunctor; @@ -423,4 +423,6 @@ int main(int argc, char* argv[]) return 0; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp32.cpp b/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp32.cpp index 7181ffebc0..4cad866132 100644 --- a/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp32.cpp +++ b/example/43_splitk_gemm_bias_e_permute/splitk_gemm_bias_e_permute_xdl_fp32.cpp @@ -16,10 +16,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif using ::ck::DeviceMem; using ::ck::HostTensorDescriptor; using ::ck::make_ParallelTensorFunctor; @@ -423,4 +423,6 @@ int main(int argc, char* argv[]) return 0; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py index 875a655976..849f463afa 100644 --- a/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py +++ b/example/ck_tile/01_fmha/codegen/ops/fmha_fwd_splitkv.py @@ -121,10 +121,10 @@ using trait_{F_idx} = fmha_fwd_splitkv_traits_<{F_hdim}, {F_dtype}, {F_mode}, {F {F_pipeline_enum}, {F_logits}, fmha_mask_{F_idx}, {F_bias}, {F_lse}, {F_squant}, {F_pagedkv}, {F_sink}, {F_spad}, {F_skpad}, {F_dpad}, {F_dvpad}>; +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wtautological-compare" - +#endif namespace {{ template void run_instance(const ck_tile::stream_config& s, fmha_fwd_splitkv_args a) {{ @@ -142,7 +142,9 @@ void run_instance(const ck_tile::stream_config& s, fmha_fwd_splitkv_args a) {{ }} }} // anonymous namespace +#ifdef __clang__ #pragma clang diagnostic pop +#endif template<> void fmha_fwd_splitkv_oneshot_(const ck_tile::stream_config& s, fmha_fwd_splitkv_args a) diff --git a/example/ck_tile/01_fmha/fmha_fwd_head_grouping.hpp b/example/ck_tile/01_fmha/fmha_fwd_head_grouping.hpp index 1565fe8566..a4e48ff124 100644 --- a/example/ck_tile/01_fmha/fmha_fwd_head_grouping.hpp +++ b/example/ck_tile/01_fmha/fmha_fwd_head_grouping.hpp @@ -22,9 +22,10 @@ #define CK_TILE_FMHA_ENABLE_HEAD_GROUPING 1 #endif +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif #if CK_TILE_FMHA_ENABLE_HEAD_GROUPING CK_TILE_DECLARE_ENV_VAR_BOOL(CK_TILE_FMHA_HEAD_GROUP_LOG) CK_TILE_DECLARE_ENV_VAR_BOOL(CK_TILE_FMHA_DISABLE_HEAD_GROUPING) @@ -430,4 +431,6 @@ float run_fwd_head_grouped(const ck_tile::stream_config& sc, } // namespace fmha_fwd_head_grouping #endif // CK_TILE_FMHA_ENABLE_HEAD_GROUPING +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/ck_tile/01_fmha/fmha_fwd_runner.hpp b/example/ck_tile/01_fmha/fmha_fwd_runner.hpp index 2b0033524b..0b51dffa46 100644 --- a/example/ck_tile/01_fmha/fmha_fwd_runner.hpp +++ b/example/ck_tile/01_fmha/fmha_fwd_runner.hpp @@ -27,9 +27,10 @@ #error "we should enable fmha_fwd_splitkv() api in order to cooperate with fmha_fwd_appendkv()" #endif +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif enum class fwd_result { success, @@ -2492,4 +2493,6 @@ fwd_result fmha_fwd_run(mode_enum mode, return pass ? fwd_result::success : fwd_result::failure; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/ck_tile/01_fmha/quant.hpp b/example/ck_tile/01_fmha/quant.hpp index 885d014695..4f854fde1d 100644 --- a/example/ck_tile/01_fmha/quant.hpp +++ b/example/ck_tile/01_fmha/quant.hpp @@ -8,10 +8,10 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/fmha.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // keep sync with BlockAttentionQuantScaleEnum enum class quant_scale_enum { @@ -76,4 +76,6 @@ struct quant_scale_info return os; } }; +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/ck_tile/03_gemm/gemm_splitk_two_stage_reduce.cpp b/example/ck_tile/03_gemm/gemm_splitk_two_stage_reduce.cpp index 1b339c39bb..6faca63c17 100644 --- a/example/ck_tile/03_gemm/gemm_splitk_two_stage_reduce.cpp +++ b/example/ck_tile/03_gemm/gemm_splitk_two_stage_reduce.cpp @@ -16,10 +16,10 @@ #include "gemm_utils.hpp" #include "run_gemm_example.inc" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif /** * @brief Tile partitioner with output offset support. * @@ -979,4 +979,6 @@ int main(int argc, char* argv[]) return EXIT_SUCCESS; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_utils.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_utils.hpp index bb403e6062..fcd7004f35 100644 --- a/example/ck_tile/20_grouped_convolution/grouped_convolution_utils.hpp +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_utils.hpp @@ -15,10 +15,10 @@ #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" #include "conv_configs.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" - +#endif struct GemmWarpConfig_Mfma { static constexpr ck_tile::index_t M_Warp_Tile = 32; @@ -163,4 +163,6 @@ struct InvokerResult float ave_time; ck_tile::index_t split_k; }; +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/experimental/builder/test/unit_device_buffer.cpp b/experimental/builder/test/unit_device_buffer.cpp index 1691088e1d..1f512720b3 100644 --- a/experimental/builder/test/unit_device_buffer.cpp +++ b/experimental/builder/test/unit_device_buffer.cpp @@ -50,9 +50,10 @@ TEST(DeviceBuffer, AutoFree) // In this test we are explicitly testing a pointer that is out of scope, so // we have to disable the clang compiler's lifestime safety checks. +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-permissive" +#endif { auto buffer = ckt::alloc_buffer(size); ptr = buffer.get(); @@ -60,7 +61,9 @@ TEST(DeviceBuffer, AutoFree) // Trying to use a pointer after freeing should return en error in HIP. EXPECT_THAT(hipMemset(ptr, 0xFF, size), HipError(hipErrorInvalidValue)); +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif // Reset internal HIP error state. // Otherwise, the error may leak into other tests, triggering anything that diff --git a/include/ck/library/utility/gpu_verification.hpp b/include/ck/library/utility/gpu_verification.hpp index b69fb33691..cd16354541 100644 --- a/include/ck/library/utility/gpu_verification.hpp +++ b/include/ck/library/utility/gpu_verification.hpp @@ -18,9 +18,10 @@ #include "ck/host_utility/hip_check_error.hpp" #include "ck/library/utility/check_err.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace profiler { @@ -520,4 +521,6 @@ float gpu_reduce_max(Iterator device_buffer, std::size_t size, hipStream_t strea } // namespace profiler } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/library/utility/host_tensor.hpp b/include/ck/library/utility/host_tensor.hpp index c3865820d1..c6028e5351 100644 --- a/include/ck/library/utility/host_tensor.hpp +++ b/include/ck/library/utility/host_tensor.hpp @@ -23,11 +23,11 @@ #include "ck/tensor_operation/gpu/device/tensor_layout.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" #pragma clang diagnostic ignored "-Wlifetime-safety-cross-tu-suggestions" - +#endif namespace ck { template @@ -1202,4 +1202,6 @@ struct Tensor }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor/static_tensor.hpp b/include/ck/tensor/static_tensor.hpp index c29c1a4a66..5543ab6271 100644 --- a/include/ck/tensor/static_tensor.hpp +++ b/include/ck/tensor/static_tensor.hpp @@ -4,9 +4,10 @@ #ifndef CK_STATIC_TENSOR_HPP #define CK_STATIC_TENSOR_HPP +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { // StaticTensor for Scalar @@ -273,5 +274,7 @@ __host__ __device__ constexpr auto make_static_tensor(TensorDesc, X invalid_elem } } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop #endif +#endif diff --git a/include/ck/tensor_description/multi_index_transform.hpp b/include/ck/tensor_description/multi_index_transform.hpp index ef1d6177fa..ab806b1bb2 100644 --- a/include/ck/tensor_description/multi_index_transform.hpp +++ b/include/ck/tensor_description/multi_index_transform.hpp @@ -6,10 +6,10 @@ #include "ck/utility/common_header.hpp" #include "ck/utility/multi_index.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -2281,4 +2281,6 @@ struct Xor } }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_description/tensor_descriptor.hpp b/include/ck/tensor_description/tensor_descriptor.hpp index 1958c2b981..3a5d258fae 100644 --- a/include/ck/tensor_description/tensor_descriptor.hpp +++ b/include/ck/tensor_description/tensor_descriptor.hpp @@ -7,9 +7,10 @@ #include "ck/utility/sequence_helper.hpp" #include "ck/tensor_description/multi_index_transform.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template @@ -642,4 +643,6 @@ using TensorCoordinateStep_t = decltype(make_tensor_coordinate_step( TensorDesc{}, MultiIndex::GetNumOfDimension()>{})); } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_description/tensor_space_filling_curve.hpp b/include/ck/tensor_description/tensor_space_filling_curve.hpp index 353ce65c61..673485cba1 100644 --- a/include/ck/tensor_description/tensor_space_filling_curve.hpp +++ b/include/ck/tensor_description/tensor_space_filling_curve.hpp @@ -10,10 +10,10 @@ #include "ck/utility/tuple_helper.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace detail { @@ -197,4 +197,6 @@ struct SpaceFillingCurve } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp index 6d684696b4..9b30fd5580 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp @@ -10,9 +10,10 @@ #include "ck/tensor_operation/gpu/warp/wmma_gemm.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp index 4611aa6d6d..c2e9e2edac 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp @@ -13,9 +13,10 @@ // Prefetech 2 stage // Local prefetch 1 stage +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp index 40a2aa28b6..ba710781b5 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp @@ -9,10 +9,10 @@ #include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_wmma.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_wmma.hpp index bb3377930d..827ea29c07 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_wmma.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_wmma.hpp @@ -11,10 +11,10 @@ #define CK_MNK_LOOP +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { #ifdef __gfx12__ @@ -1014,4 +1014,6 @@ struct BlockwiseGemmWMMA #endif } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp index 25a29a277b..4f73be7773 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp @@ -9,9 +9,10 @@ #include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template @@ -1033,4 +1034,6 @@ struct BlockwiseGemmXdlops_v2 }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops_skip_b_lds.hpp b/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops_skip_b_lds.hpp index 4fe0b77f27..e80562ed8e 100644 --- a/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops_skip_b_lds.hpp +++ b/include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops_skip_b_lds.hpp @@ -8,10 +8,10 @@ #include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp" #include "ck/tensor_description/tensor_adaptor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/device_grouped_gemm.hpp b/include/ck/tensor_operation/gpu/device/device_grouped_gemm.hpp index 426b008257..7b6dbd6f94 100644 --- a/include/ck/tensor_operation/gpu/device/device_grouped_gemm.hpp +++ b/include/ck/tensor_operation/gpu/device/device_grouped_gemm.hpp @@ -12,10 +12,10 @@ #include "device_base.hpp" #include "ck/utility/ignore.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -183,4 +183,6 @@ struct DeviceGroupedGemm : public BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle_v3.hpp index 2f4926d452..14fc62f4a9 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_wmma_cshuffle_v3.hpp @@ -20,9 +20,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/utility/scheduler_enum.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_xdl_cshuffle.hpp index 419b8b3a1f..af6d12e5ad 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_contraction_multiple_d_xdl_cshuffle.hpp @@ -18,9 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_wmma_cshuffle_v3_common.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_wmma_cshuffle_v3_common.hpp index 0346f90b72..9055f9a02a 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_wmma_cshuffle_v3_common.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_wmma_cshuffle_v3_common.hpp @@ -18,10 +18,10 @@ #include "ck/utility/scheduler_enum.hpp" #include "ck/utility/integral_constant.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -705,16 +705,15 @@ struct DeviceGemmGemm_Wmma_CShuffleV3_Common_Invoker_Arg auto print = [&curFunc](const char* format, ...) -> void { if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING))) { -#if defined(__clang__) +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wformat-nonliteral" #endif va_list args; va_start(args, format); std::vfprintf(stdout, format, args); va_end(args); -#if defined(__clang__) +#ifdef __clang__ #pragma clang diagnostic pop #endif std::cout << "In file: " << __FILE__ << ", function: " << curFunc << "\n"; @@ -905,4 +904,6 @@ struct DeviceGemmGemm_Wmma_CShuffleV3_Common_Invoker_Arg } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_xdl_cshuffle.hpp index 97af5fd88d..5392055556 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_gemm_xdl_cshuffle.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/host_utility/io.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -809,4 +809,6 @@ struct DeviceBatchedGemmGemm_Xdl_CShuffle : public DeviceBatchedGemmGemm= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multi_d_xdl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multi_d_xdl.hpp index 7e414eeed7..c6d18d57d5 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multi_d_xdl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_multi_d_xdl.hpp @@ -18,9 +18,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/host_utility/io.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -794,4 +795,6 @@ struct DeviceBatchedGemmMultiD_Xdl : public DeviceBatchedGemmMultiD= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_wmma_cshuffle_v3.hpp index 26891b4367..febc1b079e 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_wmma_cshuffle_v3.hpp @@ -19,10 +19,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_xdl_cshuffle.hpp index f197ee3157..3748c6fa7d 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_reduce_xdl_cshuffle.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1048,4 +1048,6 @@ struct DeviceBatchedGemmReduce_Xdl_CShuffle : public DeviceGemmReduce<0, ReduceO } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp index adcfe9c5b6..069d77dd1b 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp @@ -22,10 +22,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1599,4 +1599,6 @@ struct DeviceBatchedGemmSoftmaxGemmPermute_Wmma_CShuffle } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp index 187ed9ad72..c840c1217b 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp @@ -19,10 +19,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1010,4 +1010,6 @@ struct DeviceBatchedGemmSoftmaxGemmPermute_Xdl_CShuffle } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp index 63d1125fc9..9b84f665fc 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp @@ -22,10 +22,10 @@ #include "ck/tensor_operation/gpu/device/matrix_padder.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_batched_gemm_softmax_gemm_xdl_cshuffle_v1.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1209,4 +1209,6 @@ struct DeviceBatchedGemmSoftmaxGemm_Xdl_CShuffle } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp index 14f89a49bb..f87c41bff8 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batched_gemm_xdl.hpp @@ -16,10 +16,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -493,4 +493,6 @@ struct DeviceBatchedGemmXdl : public DeviceBatchedGemm= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_backward_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_backward_impl.hpp index 202ed1088d..634cfec213 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_backward_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_backward_impl.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -879,4 +879,6 @@ struct DeviceBatchNormBwdImpl : public DeviceBatchNormBwd= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl.hpp index a4995ccc62..8b7afbce00 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl.hpp @@ -19,10 +19,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/host_utility/hip_check_error.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -827,4 +827,6 @@ struct DeviceBatchNormFwdImpl : public DeviceBatchNormFwd= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl_obsolete.hpp b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl_obsolete.hpp index 01f2b60cc8..2c2f494467 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl_obsolete.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_batchnorm_forward_impl_obsolete.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -719,4 +719,6 @@ struct DeviceBatchNormFwdImpl : public DeviceBatchNormFwd= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_cgemm_4gemm_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_cgemm_4gemm_xdl_cshuffle.hpp index e43babb783..46d940f17a 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_cgemm_4gemm_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_cgemm_4gemm_xdl_cshuffle.hpp @@ -19,10 +19,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -679,4 +679,6 @@ struct DeviceCGemm_4Gemm_Xdl_CShuffle } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp index 3aae690e61..43fa9fb8ad 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1014,4 +1014,6 @@ struct } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp index b5b0fa7674..d76e1505fd 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -974,4 +974,6 @@ struct DeviceConv2dFwdXdl_C_Shuffle_Bias_Activation_Input_N_Hi_Wi_C_Weight_K_Y_X } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp index e8995de746..dc1fe4da0e 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -941,4 +941,6 @@ struct DeviceConv2dFwdXdl_C_Shuffle_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_W } // namespace tensor_operation } // namespace ck // +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp index 894dd7c3ac..dc84f2aaca 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -728,4 +728,6 @@ struct DeviceConv2dFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp b/include/ck/tensor_operation/gpu/device/impl/device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp index 19fb1afbe8..d44458620a 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp @@ -14,10 +14,10 @@ #include "ck/library/utility/convolution_parameter.hpp" #include "ck/library/reference_tensor_operation/gpu/naive_conv_fwd_gpu.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -266,4 +266,6 @@ struct DeviceConv3dFwdNaive_Input_N_Di_Hi_Wi_C_Weight_K_Z_Y_X_C_Output_N_Do_Ho_W } // namespace ck #endif +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp index 3980c055cf..cbd97bfb33 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1589,4 +1589,6 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Dl } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp index f037f00b28..a21d6c633c 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/library/utility/numeric.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1589,5 +1589,7 @@ struct DeviceConvNdBwdDataNwcKxcNwk_Xdl } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_wmma_cshuffle_v3.hpp index c84fa51a40..6f3715378c 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_wmma_cshuffle_v3.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_xdl_cshuffle.hpp index ac70ef0dbc..6caa070543 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_bias_add_reduce_xdl_cshuffle.hpp @@ -16,10 +16,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -919,4 +919,6 @@ struct DeviceGemmBiasAddReduce_Xdl_CShuffle : public DeviceGemmReduce<1, ReduceO } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_wmma_cshuffle_v3.hpp index c1e9048d1b..4e0eaec3b2 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_layernorm_wmma_cshuffle_v3.hpp @@ -19,9 +19,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp index bf22814a07..d077dc5438 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp @@ -17,9 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp index 53e8964bfb..18a46b3bb8 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_wmma_cshuffle_v3.hpp @@ -18,9 +18,10 @@ #include "ck/host_utility/flush_cache.hpp" #include "ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -411,4 +412,6 @@ struct DeviceGemmMultipleD_Wmma_CShuffleV3 } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp index 42ed344702..b21aa08f69 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_multiple_d_xdl_cshuffle.hpp @@ -19,9 +19,10 @@ #include "ck/tensor_operation/gpu/device/matrix_padder.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_multiple_d_xdl_cshuffle.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_wmma_cshuffle_v3.hpp index e223cb5fbe..618fd8abb7 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_wmma_cshuffle_v3.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_xdl_cshuffle.hpp index 4cf8456117..f4e97a4bd9 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_reduce_xdl_cshuffle.hpp @@ -16,10 +16,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -883,4 +883,6 @@ struct DeviceGemmReduce_Xdl_CShuffle : public DeviceGemmReduce<0, ReduceOperatio } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3r1.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3r1.hpp index 2e8d32ddb4..3112dea36d 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3r1.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_wmma_cshuffle_v3r1.hpp @@ -26,10 +26,10 @@ #include "ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp" #include "ck/tensor_operation/gpu/device/impl/device_reduce_threadwise_multi_d.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -584,4 +584,6 @@ struct DeviceGemm_Wmma_CShuffleV3R1 : public DeviceGemmV2R1= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_layernorm_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_layernorm_cshuffle.hpp index 1c0d00b8f7..02ad9079ea 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_layernorm_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_layernorm_cshuffle.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -816,4 +816,6 @@ struct DeviceGemmLayerNorm_Xdl_CShuffle : public BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_skip_b_lds.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_skip_b_lds.hpp index 92f6d4370f..fb5fea27f1 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_skip_b_lds.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_skip_b_lds.hpp @@ -18,10 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -560,4 +560,6 @@ struct DeviceGemmXdlSkipBLds : public DeviceGemm= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_waveletmodel_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_waveletmodel_cshuffle.hpp index 9c1ada3b68..075f3db489 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_waveletmodel_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_gemm_xdl_waveletmodel_cshuffle.hpp @@ -17,10 +17,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp index 35bc2d1b11..136d59a160 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_dl.hpp @@ -27,10 +27,10 @@ #endif #include "ck/tensor_operation/gpu/device/tensor_size_check.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1396,4 +1396,6 @@ struct DeviceGroupedConvBwdWeight_Dl : public DeviceGroupedConvBwdWeight= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp index 862ddb50d5..9b0358e5d2 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.hpp @@ -36,9 +36,10 @@ #endif #include "ck/tensor_operation/gpu/device/tensor_size_check.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -2410,4 +2411,6 @@ struct DeviceGroupedConvFwdMultipleABD_Wmma_CShuffle_V3 #ifdef CK_EXPERIMENTAL_BUILDER #include "ck_tile/builder/reflect/reflect_device_grouped_conv_fwd_multiple_abd_wmma_cshuffle_v3.inc" #endif +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp index 4d2351a348..21eebd4481 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp @@ -34,9 +34,10 @@ #endif #include "ck/tensor_operation/gpu/device/tensor_size_check.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -2264,4 +2265,6 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp index ee578971e2..10785a17d0 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp @@ -33,9 +33,10 @@ #endif #include "ck/tensor_operation/gpu/device/tensor_size_check.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -2424,4 +2425,6 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp index 572314c023..371e59988d 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp @@ -24,9 +24,10 @@ #include "ck/host_utility/io.hpp" #include "ck/library/utility/numeric.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1160,4 +1161,6 @@ struct DeviceGroupedConvFwdMultipleDMultipleR_Xdl_CShuffle } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp index 68845f7fa9..737c019edd 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp @@ -30,9 +30,10 @@ #endif #include "ck/tensor_operation/gpu/device/tensor_size_check.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1059,4 +1060,6 @@ struct DeviceGroupedConvFwdMultipleD_Wmma_CShuffle } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp index 72b1fb6b5f..d878cfc1b5 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_query_attention_forward_wmma.hpp @@ -22,10 +22,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1124,4 +1124,6 @@ struct DeviceGroupedQueryAttentionForward_Wmma } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp index 7f7f669602..1fac0160c5 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_image_to_column_impl.hpp @@ -19,10 +19,10 @@ #include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp" #include "ck/host_utility/io.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -411,4 +411,6 @@ struct DeviceImageToColumnImpl } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp b/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp index 0a8922de8a..91dfd504b2 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_multi_query_attention_forward_wmma.hpp @@ -22,10 +22,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -1114,4 +1114,6 @@ struct DeviceMultiQueryAttentionForward_Wmma } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_multiblock.hpp b/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_multiblock.hpp index 0e31488487..671894546b 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_multiblock.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_multiblock.hpp @@ -17,9 +17,10 @@ #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -596,4 +597,6 @@ struct DeviceMultipleReduceMultiBlock : public DeviceMultipleReduce= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_threadwise.hpp b/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_threadwise.hpp index 389754f533..62198843e6 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_threadwise.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_multiple_reduce_threadwise.hpp @@ -16,9 +16,10 @@ #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -423,4 +424,6 @@ struct DeviceMultipleReduceThreadWise : public DeviceMultipleReduce= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_data_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_data_impl.hpp index 887322379c..8ae6a305c6 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_data_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_data_impl.hpp @@ -14,10 +14,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // M is Invariant dimension, K is reduced dimension namespace ck { namespace tensor_operation { @@ -468,4 +468,6 @@ struct DeviceNormalizationBwdDataImpl : public DeviceNormalizationBwdData= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_gamma_beta_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_gamma_beta_impl.hpp index 153685af13..b2d50c4c41 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_gamma_beta_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_normalization_bwd_gamma_beta_impl.hpp @@ -14,10 +14,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // M is Invariant dimension, K is reduced dimension namespace ck { namespace tensor_operation { @@ -481,4 +481,6 @@ struct DeviceNormalizationBwdGammaBetaImpl } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_normalization_fwd_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_normalization_fwd_impl.hpp index 429e9cff79..cbf670b0eb 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_normalization_fwd_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_normalization_fwd_impl.hpp @@ -14,10 +14,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -478,4 +478,6 @@ struct DeviceNormalizationFwdImpl : public DeviceNormalizationFwd= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp index ea52a442ae..12ca511a5e 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_permute_impl.hpp @@ -17,9 +17,10 @@ #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { namespace tensor_operation { namespace device { @@ -283,4 +284,6 @@ struct DevicePermuteImpl : DevicePermute= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp b/include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp index 3c0d110540..9ee5e9b07c 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_pool2d_fwd_nhwc_nhwc.hpp @@ -16,10 +16,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -391,4 +391,6 @@ struct DevicePool2dFwd_NHWC_NHWC : public DevicePoolFwd<4, } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp b/include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp index d53ab6cadb..30c6b51103 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_pool3d_fwd_ndhwc_ndhwc.hpp @@ -16,10 +16,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -414,4 +414,6 @@ struct DevicePool3dFwd_NDHWC_NDHWC : public DevicePoolFwd<5, } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_put_element_impl.hpp b/include/ck/tensor_operation/gpu/device/impl/device_put_element_impl.hpp index cef984d581..10bdd780fb 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_put_element_impl.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_put_element_impl.hpp @@ -14,10 +14,10 @@ #include "ck/host_utility/kernel_launch.hpp" #include "ck/host_utility/stream_utility.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -158,4 +158,6 @@ struct DevicePutElementImpl } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise.hpp b/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise.hpp index d711e4705b..13c16a4bcb 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise.hpp @@ -13,10 +13,10 @@ #include "ck/tensor_operation/gpu/device/impl/device_reduce_common.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_2d_reduction_threadwise.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -395,4 +395,6 @@ struct DeviceReduceThreadWise : public DeviceReduce= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise_multi_d.hpp b/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise_multi_d.hpp index d6c7afd8b2..dc9db637aa 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise_multi_d.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_reduce_threadwise_multi_d.hpp @@ -14,10 +14,10 @@ #include "ck/tensor_operation/gpu/grid/gridwise_2d_reduction_threadwise_multi_d.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -415,4 +415,6 @@ struct DeviceReduceThreadWiseMultiD : public DeviceReduceMultiD= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_sparse_embeddings_forward_layernorm.hpp b/include/ck/tensor_operation/gpu/device/impl/device_sparse_embeddings_forward_layernorm.hpp index 10a4d74ddc..1b5e401319 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_sparse_embeddings_forward_layernorm.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_sparse_embeddings_forward_layernorm.hpp @@ -18,10 +18,10 @@ #include "ck/tensor_operation/gpu/grid/gridwise_sparse_embeddings_forward_layernorm.hpp" #endif +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -200,4 +200,6 @@ struct DeviceSparseEmbeddingsForwardLayernorm : public BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/device/impl/device_splitk_contraction_multiple_d_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_splitk_contraction_multiple_d_xdl_cshuffle.hpp index 6213b44db4..29db59fe32 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_splitk_contraction_multiple_d_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_splitk_contraction_multiple_d_xdl_cshuffle.hpp @@ -18,9 +18,10 @@ #include "ck/host_utility/device_prop.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_welford_wmma.hpp b/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_welford_wmma.hpp index 4d92dbdc0c..aaeaf96338 100644 --- a/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_welford_wmma.hpp +++ b/include/ck/tensor_operation/gpu/grid/epilogue_cshuffle_v3_welford_wmma.hpp @@ -8,10 +8,10 @@ #include "ck/tensor_operation/gpu/block/blockwise_welford.hpp" #include "ck/tensor_operation/gpu/thread/threadwise_welford.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_gemm_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_gemm_wmma_cshuffle_v3.hpp index 36ea80f472..c276be5dbe 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_gemm_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_gemm_wmma_cshuffle_v3.hpp @@ -510,16 +510,17 @@ struct GridwiseBatchedGemmGemm_wmma_cshuffle_v3 auto print = [&curFunc](const char* format, ...) -> void { if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING))) { -#if defined(__clang__) + +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wformat-nonliteral" #endif va_list args; va_start(args, format); std::vfprintf(stdout, format, args); va_end(args); -#if defined(__clang__) + +#ifdef __clang__ #pragma clang diagnostic pop #endif std::cout << "In file: " << __FILE__ << ", function: " << curFunc << "\n"; diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp index 787a957219..52bd7e33fb 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_dpp.hpp @@ -16,10 +16,10 @@ #include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -697,4 +697,6 @@ struct GridwiseGemm_ak0mak1_bk0nbk1_mn_dpp } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3.hpp index 844b1cc7a2..897891ea4c 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3.hpp @@ -20,10 +20,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { /// @brief \"Universal\" GEMM kernel with SplitK support. @@ -1189,4 +1189,6 @@ struct GridwiseGemm_wmma_cshuffle_v3 }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp index 2c7a8e2eaf..1d0eb4f2d0 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp @@ -16,10 +16,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_wmma_cshuffle_v3_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_conv_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_conv_v3.hpp index 1673be0dcd..c9434ebfa2 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_conv_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_conv_v3.hpp @@ -15,10 +15,10 @@ #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_direct_load.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp index f2abdefff0..ad3c31e467 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_streamk_v3.hpp @@ -17,10 +17,10 @@ #include "ck/utility/reduction_functions_accumulate.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -2164,4 +2164,6 @@ struct GridwiseGemm_xdl_cshuffle_streamk_v3 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp index 2ff3f04821..ebace17416 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v1.hpp @@ -17,10 +17,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -1034,4 +1034,6 @@ struct GridwiseGemm_k0mk1_k0nk1_mn_xdl_cshuffle_v1 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v2.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v2.hpp index e49fb5be40..adaf3209a3 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v2.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v2.hpp @@ -15,10 +15,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -1028,4 +1028,6 @@ struct GridwiseGemm_xdl_cshuffle_v2 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp index 3bdcbc336c..0c8707476e 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3.hpp @@ -17,10 +17,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1739,4 +1739,6 @@ struct GridwiseGemm_xdl_cshuffle_v3 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp index 1c39f62e4d..a8eb85cf7b 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp @@ -15,10 +15,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1476,4 +1476,6 @@ struct GridwiseGemm_xdl_cshuffle_v3_b_preshuffle }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif 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 7a6de68579..19e98469d0 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 @@ -15,10 +15,10 @@ #include "ck/utility/env.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1607,4 +1607,6 @@ struct GridwiseGemm_xdl_cshuffle_v3 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp index 73f54ea93b..753d801a98 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp @@ -15,9 +15,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1696,4 +1697,6 @@ struct GridwiseGemm_xdl_cshuffle_v3 }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp index 38385711ba..99e5828661 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp @@ -20,10 +20,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1961,4 +1961,6 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp index 7d2c6d8f28..27d5c4be61 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp @@ -17,10 +17,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1288,4 +1288,6 @@ struct GridwiseGemmMultiD_ABScale_xdl_cshuffle_v3 }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp index e2a2c0145e..f0b750aef2 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp @@ -17,10 +17,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1474,4 +1474,6 @@ struct GridwiseGemmMultiD_xdl_cshuffle_v3_b_preshuffle } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_blockscale_b_preshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_blockscale_b_preshuffle.hpp index 6c4c279fa4..c0abfb107c 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_blockscale_b_preshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_multi_d_blockscale_b_preshuffle.hpp @@ -17,10 +17,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1549,4 +1549,6 @@ struct GridwiseGemmMultiD_blockscale_xdl_cshuffle_v3_b_preshuffle }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx.hpp index 25e68bcb10..061aaf8c39 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx.hpp @@ -17,10 +17,10 @@ #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_direct_load.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1766,4 +1766,6 @@ struct GridwiseGemmMX_xdl_cshuffle_v3 }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp index 7206a5766b..bad66d58a6 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp @@ -16,10 +16,10 @@ #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_direct_load.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -1502,4 +1502,6 @@ struct GridwiseGemmMX_xdl_cshuffle_v3_bpreshuffle }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_bwd_weight.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_bwd_weight.hpp index 0539f5a2af..181572d3f9 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_bwd_weight.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_bwd_weight.hpp @@ -17,10 +17,10 @@ #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" #include "ck/tensor_operation/gpu/device/device_base.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Implementation of "Merge" transformation primitive that uses division and mod. It is supposed to @@ -1147,4 +1147,6 @@ struct GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_bwd_weight }; // namespace ck } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_splitk_lds_direct_load.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_splitk_lds_direct_load.hpp index 3bdc593ca8..1270220dbe 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_splitk_lds_direct_load.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_splitk_lds_direct_load.hpp @@ -21,10 +21,10 @@ #include "ck/tensor_operation/gpu/device/matrix_padder.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_streamk.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_streamk.hpp index 66188babee..84809dbeee 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_streamk.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_streamk.hpp @@ -20,10 +20,10 @@ #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" #include "ck/host_utility/device_prop.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -1212,4 +1212,6 @@ struct GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_streamk } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp index 10020c166b..ed84774cae 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r3.hpp @@ -17,10 +17,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp index 7b97c0f5b7..f99190ee45 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_v2r4r2.hpp @@ -18,10 +18,10 @@ #include "ck/tensor_operation/gpu/device/matrix_padder.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp index 7ea7439f22..049388766c 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm.hpp @@ -18,10 +18,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -2126,4 +2126,6 @@ struct GridwiseMoeGemm : public GridwiseGemm_xdl_cshuffle_base< }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm_blockscale.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm_blockscale.hpp index ce89c94c04..5b791d3668 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm_blockscale.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_moe_gemm_blockscale.hpp @@ -15,10 +15,10 @@ #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -2224,4 +2224,6 @@ struct GridwiseMoeGemmBlockScale }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm.hpp index 75b4b22465..ab71fb2f77 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm.hpp @@ -19,10 +19,10 @@ #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_gather_direct_load.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #define DEBUG_LOG 0 namespace ck { @@ -2512,4 +2512,6 @@ struct GridwiseMoeGemmMX }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bns.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bns.hpp index 3dd1b6557b..5c7c37272f 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bns.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bns.hpp @@ -21,10 +21,10 @@ #define DEBUG_LOG 0 +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Currently we do not have a elegant way to put single lds buffer & double lds buffer pipe in same @@ -2394,4 +2394,6 @@ struct GridwiseMoeGemmMXBNS }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bpreshuffle.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bpreshuffle.hpp index 29bac1d4a9..04042530a2 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bpreshuffle.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_moe_mx_gemm_bpreshuffle.hpp @@ -19,10 +19,10 @@ #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_gather_direct_load.hpp" #include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #define DEBUG_LOG 0 namespace ck { @@ -2222,4 +2222,6 @@ struct GridwiseMoeGemmMX_BPreshuffle }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp index 142a90ae3a..765b395cfe 100644 --- a/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp +++ b/include/ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer.hpp @@ -12,10 +12,10 @@ #include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_util.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // Assume: // 1. src: @@ -2100,4 +2100,6 @@ struct ThreadwiseTensorSliceTransfer_StaticToStatic_IntraRow } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/amd_address_space.hpp b/include/ck/utility/amd_address_space.hpp index 46557b3d41..cfafed05fe 100644 --- a/include/ck/utility/amd_address_space.hpp +++ b/include/ck/utility/amd_address_space.hpp @@ -25,11 +25,14 @@ __device__ T* cast_pointer_to_generic_address_space(T CK_CONSTANT_ADDRESS_SPACE* { // cast a pointer in "Constant" address space (4) to "Generic" address space (0) // only c-style pointer cast seems be able to be compiled +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif return (T*)p; // NOLINT(old-style-cast) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template @@ -37,11 +40,14 @@ __host__ __device__ T CK_CONSTANT_ADDRESS_SPACE* cast_pointer_to_constant_addres { // cast a pointer in "Generic" address space (0) to "Constant" address space (4) // only c-style pointer cast seems be able to be compiled +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif return (T CK_CONSTANT_ADDRESS_SPACE*)p; // NOLINT(old-style-cast) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } } // namespace ck diff --git a/include/ck/utility/amd_cluster_load.hpp b/include/ck/utility/amd_cluster_load.hpp index dcd811e8aa..6a70281dbd 100644 --- a/include/ck/utility/amd_cluster_load.hpp +++ b/include/ck/utility/amd_cluster_load.hpp @@ -19,11 +19,15 @@ namespace detail { template __device__ __attribute__((address_space(1))) T* to_global(const T* ptr) { +#ifdef __clang__ #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wold-style-cast" #pragma clang diagnostic ignored "-Wcast-qual" +#endif return (__attribute__((address_space(1))) T*)(ptr); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } } // namespace detail diff --git a/include/ck/utility/array.hpp b/include/ck/utility/array.hpp index b0e90b2cc0..4f1147feae 100644 --- a/include/ck/utility/array.hpp +++ b/include/ck/utility/array.hpp @@ -9,10 +9,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -98,6 +98,8 @@ __host__ __device__ constexpr auto make_array() } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif #endif diff --git a/include/ck/utility/c_style_pointer_cast.hpp b/include/ck/utility/c_style_pointer_cast.hpp index 023e0c38ad..dc973ca0fb 100644 --- a/include/ck/utility/c_style_pointer_cast.hpp +++ b/include/ck/utility/c_style_pointer_cast.hpp @@ -14,12 +14,15 @@ template && is_pointer_v, bool>::type = false> __host__ __device__ PY c_style_pointer_cast([[clang::lifetimebound]] PX p_x) { +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" #pragma clang diagnostic ignored "-Wcast-align" +#endif return (PY)p_x; // NOLINT(old-style-cast, cast-align) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } } // namespace ck diff --git a/include/ck/utility/dtype_vector.hpp b/include/ck/utility/dtype_vector.hpp index 1e44d3bb28..017bd814b6 100644 --- a/include/ck/utility/dtype_vector.hpp +++ b/include/ck/utility/dtype_vector.hpp @@ -5,9 +5,10 @@ #include "ck/utility/data_type.hpp" #include "ck/utility/math.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { __device__ int static err = 0; @@ -855,4 +856,6 @@ using pk_i4x4_t = typename vector_type::type; using pk_i4x8_t = typename vector_type::type; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/dynamic_buffer.hpp b/include/ck/utility/dynamic_buffer.hpp index 2f1ed3b121..07fb93e9e2 100644 --- a/include/ck/utility/dynamic_buffer.hpp +++ b/include/ck/utility/dynamic_buffer.hpp @@ -16,10 +16,10 @@ #include "generic_memory_space_atomic.hpp" #include "data_cache_prefetch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { // T may be scalar or vector @@ -542,4 +542,6 @@ make_dynamic_buffer(T* p, ElementSpaceSize element_space_size, X invalid_element } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/env.hpp b/include/ck/utility/env.hpp index 1460f02b62..812b17a807 100644 --- a/include/ck/utility/env.hpp +++ b/include/ck/utility/env.hpp @@ -9,10 +9,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace internal { template @@ -87,8 +87,10 @@ struct EnvVar is_unset = false; value = val; } +#if __clang_major__ >= 23 #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wdeprecated-declarations" +#endif explicit EnvVar(const char* const name, const T& def_val) { // NOLINTNEXTLINE (concurrency-mt-unsafe) @@ -103,7 +105,9 @@ struct EnvVar value = def_val; } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif }; } // end namespace internal @@ -194,5 +198,7 @@ void UpdateEnvVar(EnvVar, const std::string_view& val) // environment variable to enable logging: // export CK_LOGGING=ON or CK_LOGGING=1 or CK_LOGGING=ENABLED CK_DECLARE_ENV_VAR_BOOL(CK_LOGGING) +#if __clang_major__ >= 23 #pragma clang diagnostic pop #endif +#endif diff --git a/include/ck/utility/functional2.hpp b/include/ck/utility/functional2.hpp index b704f2881f..827b26c7cf 100644 --- a/include/ck/utility/functional2.hpp +++ b/include/ck/utility/functional2.hpp @@ -8,10 +8,10 @@ #include "ck/utility/tuple.hpp" #include "ck/utility/type.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace detail { @@ -117,4 +117,6 @@ struct identity }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/sequence.hpp b/include/ck/utility/sequence.hpp index 8987f9fbf8..7f15d5bb3b 100644 --- a/include/ck/utility/sequence.hpp +++ b/include/ck/utility/sequence.hpp @@ -12,10 +12,10 @@ #include "ck/utility/functional.hpp" #include "ck/utility/math.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -1027,7 +1027,9 @@ using uniform_sequence_gen_t = typename uniform_sequence_gen::type; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif #if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC) template diff --git a/include/ck/utility/span.hpp b/include/ck/utility/span.hpp index 72efd79711..d6f40a2a5e 100644 --- a/include/ck/utility/span.hpp +++ b/include/ck/utility/span.hpp @@ -7,10 +7,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { template @@ -70,4 +70,6 @@ class span } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/static_buffer.hpp b/include/ck/utility/static_buffer.hpp index bde35b8a0c..ecfc2b8ebe 100644 --- a/include/ck/utility/static_buffer.hpp +++ b/include/ck/utility/static_buffer.hpp @@ -5,9 +5,10 @@ #include "statically_indexed_array.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { // static buffer for scalar @@ -198,4 +199,6 @@ __host__ __device__ constexpr auto make_static_buffer(LongNumber) } } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/statically_indexed_array.hpp b/include/ck/utility/statically_indexed_array.hpp index 0d3be06c6b..b7660b7178 100644 --- a/include/ck/utility/statically_indexed_array.hpp +++ b/include/ck/utility/statically_indexed_array.hpp @@ -7,10 +7,10 @@ #include "functional2.hpp" #include "tuple.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace detail { @@ -107,6 +107,8 @@ struct StaticallyIndexedArray_v2 } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif #endif diff --git a/include/ck/utility/tuple.hpp b/include/ck/utility/tuple.hpp index 0067433839..685e4aa160 100644 --- a/include/ck/utility/tuple.hpp +++ b/include/ck/utility/tuple.hpp @@ -9,11 +9,11 @@ #include "ck/utility/enable_if.hpp" #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" #pragma clang diagnostic ignored "-Wlifetime-safety-lifetimebound-violation" - +#endif namespace ck { namespace detail { @@ -275,4 +275,6 @@ using tuple_element_or_t = typename detail::tuple_element_or_impl= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/type_convert.hpp b/include/ck/utility/type_convert.hpp index 2e7317a378..5d035e10a0 100644 --- a/include/ck/utility/type_convert.hpp +++ b/include/ck/utility/type_convert.hpp @@ -13,9 +13,10 @@ #include "ck/utility/amd_inline_asm.hpp" #include "ck/utility/type.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { // Define the common macro for MI300 models #if defined(__gfx942__) || defined(__gfx950__) @@ -2948,4 +2949,6 @@ inline __host__ __device__ void array_convert(Array& y, const Array } } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/utility/workgroup_barrier.hpp b/include/ck/utility/workgroup_barrier.hpp index 791357b8fa..75127075f6 100644 --- a/include/ck/utility/workgroup_barrier.hpp +++ b/include/ck/utility/workgroup_barrier.hpp @@ -5,10 +5,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { struct workgroup_barrier { @@ -65,4 +65,6 @@ struct workgroup_barrier }; } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/wrapper/layout.hpp b/include/ck/wrapper/layout.hpp index 11d16ab144..1d289f93f1 100644 --- a/include/ck/wrapper/layout.hpp +++ b/include/ck/wrapper/layout.hpp @@ -5,10 +5,10 @@ #include "ck/wrapper/utils/layout_utils.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // Disable from doxygen docs generation /// @cond INTERNAL namespace ck { @@ -485,4 +485,6 @@ struct Layout } // namespace wrapper } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck/wrapper/tensor.hpp b/include/ck/wrapper/tensor.hpp index 3b88228db4..38cd450f76 100644 --- a/include/ck/wrapper/tensor.hpp +++ b/include/ck/wrapper/tensor.hpp @@ -7,10 +7,10 @@ #include "utils/tensor_partition.hpp" #include "utils/layout_utils.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // Disable from doxygen docs generation /// @cond INTERNAL namespace ck { @@ -444,4 +444,6 @@ struct Tensor } // namespace wrapper } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/algorithm/coordinate_transform.hpp b/include/ck_tile/core/algorithm/coordinate_transform.hpp index 71a1a14b05..5969440dd5 100644 --- a/include/ck_tile/core/algorithm/coordinate_transform.hpp +++ b/include/ck_tile/core/algorithm/coordinate_transform.hpp @@ -11,10 +11,10 @@ #include "ck_tile/core/utility/magic_div.hpp" #include "ck_tile/core/utility/print.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { enum struct coord_transform_enum @@ -1808,4 +1808,6 @@ make_indexing_transform_with_adaptor(const UpLength& up_lengths, const IndexingA } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/arch/amd_buffer_addressing.hpp b/include/ck_tile/core/arch/amd_buffer_addressing.hpp index 86c52078f4..4134b077db 100644 --- a/include/ck_tile/core/arch/amd_buffer_addressing.hpp +++ b/include/ck_tile/core/arch/amd_buffer_addressing.hpp @@ -106,9 +106,10 @@ struct buffer_store; template struct buffer_store_if; +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wundefined-reinterpret-cast" +#endif // TODO: strict aliasing rule seems fail when reinterpret_cast between vector type // (exp_vector_type(xxx)) @@ -487,7 +488,9 @@ struct buffer_load_if<1, pre_nop> }; #endif -#pragma clang diagnostic pop // "-Wundefined-reinterpret-cast" +#ifdef __clang__ +#pragma clang diagnostic pop +#endif // "-Wundefined-reinterpret-cast" template <> struct buffer_store<16> @@ -1917,9 +1920,10 @@ CK_TILE_DEVICE void amd_async_buffer_load(CK_TILE_LDS_ADDR T* smem, if constexpr(oob_conditional_check) v_offset = flag ? v_offset : src_wave_buffer_resource[2]; +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif // Use C-style cast to change address space without dropping llvm noalias attribute llvm_amdgcn_raw_buffer_load_lds(src_wave_buffer_resource, (as3_uint32_ptr)(smem), @@ -1928,7 +1932,9 @@ CK_TILE_DEVICE void amd_async_buffer_load(CK_TILE_LDS_ADDR T* smem, src_wave_addr_offset, /*src_immediate_addr_offset*/ 0, static_cast(coherence)); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template (in_ptr)); +#ifdef __clang__ #pragma clang diagnostic pop +#endif if constexpr(std::is_same_v, ck_tile::half_t>) { typedef __attribute__((__vector_size__(4 * sizeof(__fp16)))) __fp16 llvm_fp16x4_t; diff --git a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp index 7f1c264bbf..ac2b81cbfb 100644 --- a/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp +++ b/include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp @@ -113,9 +113,10 @@ struct buffer_store; template struct buffer_store_if; +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wundefined-reinterpret-cast" +#endif // TODO: strict aliasing rule seems fail when reinterpret_cast between vector type // (exp_vector_type(xxx)) @@ -500,7 +501,9 @@ struct buffer_load_if<1, pre_nop> } }; #endif -#pragma clang diagnostic pop // "-Wundefined-reinterpret-cast" +#ifdef __clang__ +#pragma clang diagnostic pop +#endif // "-Wundefined-reinterpret-cast" template struct buffer_store; @@ -1984,9 +1987,10 @@ CK_TILE_DEVICE void amd_async_buffer_load(CK_TILE_LDS_ADDR T* smem, if constexpr(oob_conditional_check) v_offset = flag ? v_offset : 0x7fffffff; // large offset to cause OOB access +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif // Use C-style cast to change address space without dropping llvm noalias attribute __builtin_amdgcn_raw_ptr_buffer_load_lds(rsrc, smem, @@ -1995,7 +1999,9 @@ CK_TILE_DEVICE void amd_async_buffer_load(CK_TILE_LDS_ADDR T* smem, src_wave_addr_offset, /*imm*/ IMM, static_cast(coherence)); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template (in_ptr)); +#ifdef __clang__ #pragma clang diagnostic pop +#endif if constexpr(std::is_same_v, ck_tile::half_t>) { #if defined(__gfx950__) diff --git a/include/ck_tile/core/arch/amd_cluster_load.hpp b/include/ck_tile/core/arch/amd_cluster_load.hpp index 19ca7690a6..893d2d7334 100644 --- a/include/ck_tile/core/arch/amd_cluster_load.hpp +++ b/include/ck_tile/core/arch/amd_cluster_load.hpp @@ -14,20 +14,28 @@ namespace ck_tile { template CK_TILE_DEVICE __attribute__((address_space(1))) T* to_global(const T* ptr) { +#ifdef __clang__ #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wold-style-cast" #pragma clang diagnostic ignored "-Wcast-qual" +#endif return (__attribute__((address_space(1))) T*)(ptr); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template CK_TILE_DEVICE __attribute__((address_space(3))) T* to_lds(T* ptr) { +#ifdef __clang__ #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wold-style-cast" +#endif return (__attribute__((address_space(3))) T*)(ptr); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } #endif // __gfx1250__ diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index d51a02df0b..6c8dcd7290 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -1299,11 +1299,14 @@ __device__ T* cast_pointer_to_generic_address_space(T CK_TILE_CONSTANT_ADDRESS_S { // cast a pointer in "Constant" address space (4) to "Generic" address space (0) // only c-style pointer cast seems be able to be compiled +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif return (T*)(p); // NOLINT(old-style-cast) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template @@ -1311,11 +1314,14 @@ __host__ __device__ T CK_TILE_CONSTANT_ADDRESS_SPACE* cast_pointer_to_constant_a { // cast a pointer in "Generic" address space (0) to "Constant" address space (4) // only c-style pointer cast seems be able to be compiled; +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif return (T CK_TILE_CONSTANT_ADDRESS_SPACE*)p; // NOLINT(old-style-cast) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } CK_TILE_HOST_DEVICE constexpr index_t get_smem_capacity() diff --git a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp index c2a31e92c8..5985f63440 100644 --- a/include/ck_tile/core/arch/mma/amdgcn_mma.hpp +++ b/include/ck_tile/core/arch/mma/amdgcn_mma.hpp @@ -23,10 +23,10 @@ #include #endif +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile::core::arch::mma { /**--------------------------------------------------- @@ -390,7 +390,9 @@ CK_TILE_HOST_DEVICE void print(amdgcn_mma= 23 #pragma clang diagnostic pop +#endif // Include the implementations #include "wmma/wmma.hpp" // should be included before the below headers diff --git a/include/ck_tile/core/arch/mma/mma_pipeline.hpp b/include/ck_tile/core/arch/mma/mma_pipeline.hpp index 418cad282d..de01760620 100644 --- a/include/ck_tile/core/arch/mma/mma_pipeline.hpp +++ b/include/ck_tile/core/arch/mma/mma_pipeline.hpp @@ -9,10 +9,10 @@ #include "mma_traits.hpp" #include "mma_transforms.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile::core::arch::mma { /*! @enum MmaPipelineOptionFlag @@ -346,4 +346,6 @@ concept MmaPipelineInterface = std::derived_from= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/array.hpp b/include/ck_tile/core/container/array.hpp index ab1d64c0ea..69f460bfe9 100644 --- a/include/ck_tile/core/container/array.hpp +++ b/include/ck_tile/core/container/array.hpp @@ -12,10 +12,10 @@ #include "ck_tile/core/utility/type_traits.hpp" #include "ck_tile/core/utility/functional.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // use aggregate initialization for this type @@ -289,4 +289,6 @@ CK_TILE_HOST_DEVICE constexpr auto to_array(const X& x) } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/map.hpp b/include/ck_tile/core/container/map.hpp index ef91fb6af9..5b7ea9c406 100644 --- a/include/ck_tile/core/container/map.hpp +++ b/include/ck_tile/core/container/map.hpp @@ -8,10 +8,10 @@ #include "ck_tile/core/container/sequence.hpp" #include "ck_tile/core/container/tuple.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // naive map @@ -161,4 +161,6 @@ CK_TILE_HOST_DEVICE static void print(const map& m) } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/span.hpp b/include/ck_tile/core/container/span.hpp index 9839cb533f..7835bfd90f 100644 --- a/include/ck_tile/core/container/span.hpp +++ b/include/ck_tile/core/container/span.hpp @@ -8,10 +8,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // implement the c++20 std::span, lightweight, non-owning reference to a sequence @@ -80,4 +80,6 @@ class span }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/static_array.hpp b/include/ck_tile/core/container/static_array.hpp index 00c352c6f6..2d24ae8bfb 100644 --- a/include/ck_tile/core/container/static_array.hpp +++ b/include/ck_tile/core/container/static_array.hpp @@ -6,10 +6,10 @@ #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/integer.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /** @@ -41,4 +41,6 @@ struct static_array }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/thread_buffer.hpp b/include/ck_tile/core/container/thread_buffer.hpp index 37a9b0c3e0..2976b878ee 100644 --- a/include/ck_tile/core/container/thread_buffer.hpp +++ b/include/ck_tile/core/container/thread_buffer.hpp @@ -7,10 +7,10 @@ #include "ck_tile/core/container/array.hpp" #include "ck_tile/core/container/tuple.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { #if CK_TILE_THREAD_BUFFER_DEFAULT == CK_TILE_THREAD_BUFFER_USE_TUPLE @@ -148,4 +148,6 @@ struct vector_traits, std::enable_if_t>> } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/container/tuple.hpp b/include/ck_tile/core/container/tuple.hpp index c147208f78..996ea3f52f 100644 --- a/include/ck_tile/core/container/tuple.hpp +++ b/include/ck_tile/core/container/tuple.hpp @@ -13,11 +13,11 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" #pragma clang diagnostic ignored "-Wlifetime-safety-lifetimebound-violation" - +#endif #ifndef CK_TILE_TUPLE_IMPL #define CK_TILE_TUPLE_IMPL 1 #endif @@ -861,4 +861,6 @@ struct tuple_element> } \ }() #endif +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/numeric/e4m3.hpp b/include/ck_tile/core/numeric/e4m3.hpp index 1010c01517..715cb80f26 100644 --- a/include/ck_tile/core/numeric/e4m3.hpp +++ b/include/ck_tile/core/numeric/e4m3.hpp @@ -5,10 +5,10 @@ #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/scale_util.hpp" - +#if __clang_major__ >= 23 #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { struct e4m3_bexp_t @@ -89,4 +89,6 @@ CK_TILE_HOST_DEVICE e4m3_bexp_t::operator float() const } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/numeric/e5m3.hpp b/include/ck_tile/core/numeric/e5m3.hpp index 07f0868dec..daf799daf6 100644 --- a/include/ck_tile/core/numeric/e5m3.hpp +++ b/include/ck_tile/core/numeric/e5m3.hpp @@ -5,10 +5,10 @@ #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/scale_util.hpp" - +#if __clang_major__ >= 23 #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { struct e5m3_bexp_t @@ -95,4 +95,6 @@ CK_TILE_HOST_DEVICE e5m3_bexp_t::operator float() const } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/numeric/e8m0.hpp b/include/ck_tile/core/numeric/e8m0.hpp index c3f85bf1f2..79b5e485db 100644 --- a/include/ck_tile/core/numeric/e8m0.hpp +++ b/include/ck_tile/core/numeric/e8m0.hpp @@ -6,10 +6,10 @@ #include "ck_tile/core/config.hpp" #include "ck_tile/core/numeric/mxfp_convert.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /** @@ -104,4 +104,6 @@ CK_TILE_HOST_DEVICE constexpr e8m0_bexp_t::operator float() const } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/numeric/pk_fp4.hpp b/include/ck_tile/core/numeric/pk_fp4.hpp index 3a63a374b3..1ee8a5557c 100644 --- a/include/ck_tile/core/numeric/pk_fp4.hpp +++ b/include/ck_tile/core/numeric/pk_fp4.hpp @@ -12,10 +12,10 @@ #include "ck_tile/core/numeric/mxfp_scale.hpp" #include "ck_tile/core/numeric/vector_type.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #if defined(__gfx950__) || defined(__gfx125__) #define CK_TILE_FP4_CVT_DEVICE 1 #else @@ -736,4 +736,6 @@ CK_TILE_HOST_DEVICE constexpr bf8x2_t pk_fp4_t::to_bf8x2(float scale) const } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/numeric/type_convert.hpp b/include/ck_tile/core/numeric/type_convert.hpp index 2c582f67fd..af21d6e655 100644 --- a/include/ck_tile/core/numeric/type_convert.hpp +++ b/include/ck_tile/core/numeric/type_convert.hpp @@ -13,9 +13,10 @@ #include "ck_tile/core/numeric/int8.hpp" #include "ck_tile/core/numeric/mxfp_convert.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck_tile { #if CK_TILE_USE_CUSTOM_DATA_TYPE @@ -318,4 +319,6 @@ struct pk4scaled_type_convert_impl #endif } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/buffer_view.hpp b/include/ck_tile/core/tensor/buffer_view.hpp index bd69cedbb6..029c390157 100644 --- a/include/ck_tile/core/tensor/buffer_view.hpp +++ b/include/ck_tile/core/tensor/buffer_view.hpp @@ -19,11 +19,11 @@ #include "ck_tile/core/utility/type_traits.hpp" #include "ck_tile/core/utility/ignore.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" #pragma clang diagnostic ignored "-Wlifetime-safety-lifetimebound-violation" - +#endif namespace ck_tile { // T may be scalar or vector @@ -1588,4 +1588,6 @@ CK_TILE_HOST_DEVICE void print(const buffer_view= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/static_distributed_tensor.hpp b/include/ck_tile/core/tensor/static_distributed_tensor.hpp index e9eb51f1a5..9207fd2be1 100644 --- a/include/ck_tile/core/tensor/static_distributed_tensor.hpp +++ b/include/ck_tile/core/tensor/static_distributed_tensor.hpp @@ -14,10 +14,10 @@ #include "ck_tile/core/tensor/tile_distribution.hpp" #include "ck_tile/core/container/thread_buffer.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -271,4 +271,6 @@ inline constexpr bool is_similiar_distributed_tensor_v = } // namespace detail } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/tensor_adaptor.hpp b/include/ck_tile/core/tensor/tensor_adaptor.hpp index 5ca7f468b3..43937f9b8f 100644 --- a/include/ck_tile/core/tensor/tensor_adaptor.hpp +++ b/include/ck_tile/core/tensor/tensor_adaptor.hpp @@ -12,10 +12,10 @@ #include "ck_tile/core/utility/type_traits.hpp" #include "ck_tile/core/numeric/numeric.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // Transforms: Tuple @@ -955,4 +955,6 @@ CK_TILE_HOST_DEVICE constexpr auto chain_tensor_adaptors(const X& x, const Xs&.. remove_cvref_t, \ remove_cvref_t>{trans}; \ }() +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp b/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp index a517a99dbd..fb0b627314 100644 --- a/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp +++ b/include/ck_tile/core/tensor/tensor_adaptor_coordinate.hpp @@ -14,10 +14,10 @@ #include "ck_tile/core/utility/type_traits.hpp" #include "ck_tile/core/utility/print.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -348,11 +348,14 @@ struct CK_PRINT_X_, str_literal> Args&&... args) const { constexpr auto fmt_wrap_v = get_prefix() + str_literal{} + get_suffix(); +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wformat-nonliteral" +#endif printf(fmt_wrap_v.data, get_thread_id(), args..., targs.at(number())...); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template CK_TILE_HOST_DEVICE void operator()(T&& x, Args&&... args) const @@ -372,4 +375,6 @@ CK_TILE_HOST_DEVICE void print(const tensor_adaptor_coordinate& coord) detail::CK_PRINT_X_<>{}(coord); } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/tensor_view.hpp b/include/ck_tile/core/tensor/tensor_view.hpp index 6096b47e00..dd7ca94559 100644 --- a/include/ck_tile/core/tensor/tensor_view.hpp +++ b/include/ck_tile/core/tensor/tensor_view.hpp @@ -14,10 +14,10 @@ #include "ck_tile/core/utility/functional.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // placeholder type if we want to opt-out a tensor view parameter @@ -704,4 +704,6 @@ pad_tensor_view(const TensorView& tensor_view, const TileLengths& tile_lengths, } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/tile_distribution.hpp b/include/ck_tile/core/tensor/tile_distribution.hpp index 8b63b26964..606b346e19 100644 --- a/include/ck_tile/core/tensor/tile_distribution.hpp +++ b/include/ck_tile/core/tensor/tile_distribution.hpp @@ -15,10 +15,10 @@ #include "ck_tile/core/utility/functional.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -708,4 +708,6 @@ CK_TILE_HOST_DEVICE void print(const tile_distribution= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/tensor/tile_elementwise.hpp b/include/ck_tile/core/tensor/tile_elementwise.hpp index 12437095bd..914306a9ab 100644 --- a/include/ck_tile/core/tensor/tile_elementwise.hpp +++ b/include/ck_tile/core/tensor/tile_elementwise.hpp @@ -193,9 +193,10 @@ CK_TILE_DEVICE auto cast_tile_pk_fp8_fp32(const InTensor& in_dstr_tensors) constexpr index_t thread_buffer_size_pk = thread_buffer_size / 4; auto out_dstr_tensor = make_static_distributed_tensor(in_tile_dstr); +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wuninitialized" +#endif // __builtin_amdgcn_cvt_pk_fp8_f32() this builtin requires the old value, and // will generate a v_mov_b32 vxxx [old] before cvt, which result in unwanted ISA // so we prepare an uninitialized variable purposely, and turn off the warning @@ -218,7 +219,9 @@ CK_TILE_DEVICE auto cast_tile_pk_fp8_fp32(const InTensor& in_dstr_tensors) vec_t d = bit_cast(y); out_dstr_tensor.get_thread_buffer().template set_as(number{}, d); }); +#ifdef __clang__ #pragma clang diagnostic pop +#endif return out_dstr_tensor; #else diff --git a/include/ck_tile/core/utility/debug.hpp b/include/ck_tile/core/utility/debug.hpp index 23c6215a5e..111bca09bc 100644 --- a/include/ck_tile/core/utility/debug.hpp +++ b/include/ck_tile/core/utility/debug.hpp @@ -90,15 +90,18 @@ struct CK_PRINTF(make_str_literal(" ")); constexpr auto fmt_wrap_v = get_prefix() + fmt_v + get_suffix(); +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wformat-nonliteral" +#endif printf(fmt_wrap_v.data, get_thread_id(), N, args..., bit_cast>(type_convert(buf[Is]))...); +#ifdef __clang__ #pragma clang diagnostic pop +#endif } template diff --git a/include/ck_tile/core/utility/env.hpp b/include/ck_tile/core/utility/env.hpp index 01a6b39a74..1800467caa 100644 --- a/include/ck_tile/core/utility/env.hpp +++ b/include/ck_tile/core/utility/env.hpp @@ -6,10 +6,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -112,8 +112,10 @@ struct EnvVar is_unset = false; value = val; } +#if __clang_major__ >= 23 #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wdeprecated-declarations" +#endif explicit EnvVar(const char* const name, const T& def_val) { // NOLINTNEXTLINE (concurrency-mt-unsafe) @@ -128,7 +130,9 @@ struct EnvVar value = def_val; } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif }; } // end namespace internal @@ -220,4 +224,6 @@ void UpdateEnvVar(EnvVar, const std::string_view& val) // environment variable to enable logging: // export CK_TILE_LOGGING=ON or CK_TILE_LOGGING=1 or CK_TILE_LOGGING=ENABLED CK_TILE_DECLARE_ENV_VAR_BOOL(CK_TILE_LOGGING) +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/utility/functional.hpp b/include/ck_tile/core/utility/functional.hpp index cceb87fb98..1125b2098d 100644 --- a/include/ck_tile/core/utility/functional.hpp +++ b/include/ck_tile/core/utility/functional.hpp @@ -10,9 +10,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck_tile { namespace detail { @@ -355,4 +356,6 @@ constexpr auto conditional_expr(X&& x, Y&& y) } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/core/utility/type_traits.hpp b/include/ck_tile/core/utility/type_traits.hpp index 1273062612..a8bedd964f 100644 --- a/include/ck_tile/core/utility/type_traits.hpp +++ b/include/ck_tile/core/utility/type_traits.hpp @@ -119,12 +119,15 @@ template < typename std::enable_if && std::is_pointer_v, bool>::type = false> CK_TILE_HOST_DEVICE PY c_style_pointer_cast([[clang::lifetimebound]] PX p_x) { +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" #pragma clang diagnostic ignored "-Wcast-align" +#endif return (PY)p_x; // NOLINT(old-style-cast, cast-align) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } // Template ternary: if Cond == Match, use TrueType, else FalseType diff --git a/include/ck_tile/host/arg_parser.hpp b/include/ck_tile/host/arg_parser.hpp index dee28f9263..1a8f220bc4 100644 --- a/include/ck_tile/host/arg_parser.hpp +++ b/include/ck_tile/host/arg_parser.hpp @@ -13,10 +13,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /* * a host side utility, arg parser for, either @@ -238,4 +238,6 @@ class ArgParser std::vector keys; }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/host/host_tensor.hpp b/include/ck_tile/host/host_tensor.hpp index ca7ae452f8..e6cd9f6e77 100644 --- a/include/ck_tile/host/host_tensor.hpp +++ b/include/ck_tile/host/host_tensor.hpp @@ -17,10 +17,10 @@ #include "ck_tile/host/joinable_thread.hpp" #include "ck_tile/host/ranges.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -785,4 +785,6 @@ auto get_default_stride(std::size_t row, return stride; } } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/host/joinable_thread.hpp b/include/ck_tile/host/joinable_thread.hpp index 586524221e..474c7ba5f4 100644 --- a/include/ck_tile/host/joinable_thread.hpp +++ b/include/ck_tile/host/joinable_thread.hpp @@ -59,11 +59,14 @@ class cpu_core_guard CPU_ZERO(&new_cpu_set); for(unsigned int i = 0; i < num_cores; ++i) { +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wold-style-cast" +#endif CPU_SET(i, &new_cpu_set); // NOLINT(old-style-cast) +#ifdef __clang__ #pragma clang diagnostic pop +#endif } sched_setaffinity(0, sizeof(cpu_set_t), &new_cpu_set); } diff --git a/include/ck_tile/host/kernel_launch.hpp b/include/ck_tile/host/kernel_launch.hpp index 6213ccfdb9..06cd45c1ea 100644 --- a/include/ck_tile/host/kernel_launch.hpp +++ b/include/ck_tile/host/kernel_launch.hpp @@ -15,11 +15,11 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" #pragma clang diagnostic ignored "-Wlifetime-safety-lifetimebound-violation" - +#endif namespace ck_tile { template @@ -375,4 +375,6 @@ CK_TILE_HOST float launch_kernel_time_mask_flush_cache(const stream_config& s, } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp b/include/ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp index bb0d1f7bb3..47a358600f 100644 --- a/include/ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp +++ b/include/ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp @@ -8,10 +8,10 @@ #include "ck_tile/ops/batched_contraction/utils/tensor_descriptor_utils.hpp" #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif /** * @file batched_contraction_kernel.hpp * @brief Batched Tensor Contraction Operations @@ -689,4 +689,6 @@ struct BatchedContractionKernel } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp b/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp index c3e105e3b2..922393c861 100644 --- a/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp +++ b/include/ck_tile/ops/flatmm/block/flatmm_32x512x128_1x4x1_16x16x32.hpp @@ -467,9 +467,10 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16 fp32x4_t v_acc[32]{.0f}; // B nr->kr +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif // clang-format off asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16 @@ -484,7 +485,9 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16 : _EXPAND_ASM_ARGS_CLOBBER, "s24", "s25", "s26", "s27" ); // clang-format on +#ifdef __clang__ #pragma clang diagnostic pop +#endif // return local scratch auto c = make_tuple(MakeCBlockTile(), MakeCBlockTile()); @@ -510,9 +513,10 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16 fp32x4_t v_acc[16]{.0f}; // B nr->kr +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif // clang-format off asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16 @@ -522,7 +526,9 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_BF16 : public Flatmm_32x512x128_1x4x1_16 : _EXPAND_ASM_ARGS_CLOBBER ); // clang-format on +#ifdef __clang__ #pragma clang diagnostic pop +#endif // return local scratch auto c = MakeCBlockTile(); @@ -606,9 +612,10 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16 fp32x4_t v_acc[32]{.0f}; // B nr->kr +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif // clang-format off asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16 @@ -623,7 +630,9 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16 : _EXPAND_ASM_ARGS_CLOBBER, "s24", "s25", "s26", "s27" ); // clang-format on +#ifdef __clang__ #pragma clang diagnostic pop +#endif // return local scratch auto c = make_tuple(MakeCBlockTile(), MakeCBlockTile()); @@ -649,9 +658,10 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16 fp32x4_t v_acc[16]{.0f}; // B nr->kr +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif // clang-format off asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16 @@ -661,7 +671,9 @@ struct Flatmm_32x512x128_1x4x1_16x16x32_FP16 : public Flatmm_32x512x128_1x4x1_16 : _EXPAND_ASM_ARGS_CLOBBER ); // clang-format on +#ifdef __clang__ #pragma clang diagnostic pop +#endif // return local scratch auto c = MakeCBlockTile(); diff --git a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp index 968c79471e..15886fcf85 100644 --- a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp +++ b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32.hpp @@ -169,9 +169,10 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 : public FlatmmSn_32x128x512_1x4x // B nr->kr // clang-format off +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc" @@ -310,7 +311,9 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_BF16 : public FlatmmSn_32x128x512_1x4x "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251", "v252", "v253", "v254", "v255" ); +#ifdef __clang__ #pragma clang diagnostic pop +#endif // clang-format on } }; @@ -414,9 +417,10 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16 : public FlatmmSn_32x128x512_1x4x // B nr->kr // clang-format off +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc" @@ -555,7 +559,9 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16 : public FlatmmSn_32x128x512_1x4x "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251", "v252", "v253", "v254", "v255" ); +#ifdef __clang__ #pragma clang diagnostic pop +#endif // clang-format on } }; diff --git a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp index 6384db685a..066ccc4eaf 100644 --- a/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp +++ b/include/ck_tile/ops/flatmm/block/flatmm_sn_32x128x512_1x4x1_16x16x32_itl.hpp @@ -113,9 +113,10 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl : public FlatmmSn_32x128x512_ // B nr->kr // clang-format off +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_BF16 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16_itl.inc" @@ -257,7 +258,9 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_BF16_itl : public FlatmmSn_32x128x512_ "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251", "v252", "v253", "v254", "v255" ); +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif // clang-format on } }; @@ -361,9 +364,10 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16_itl : public FlatmmSn_32x128x512_ // B nr->kr // clang-format off +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Winline-asm" +#endif asm volatile( #define CK_TILE_FLATMM_UK_MFMA CK_TILE_FLATMM_UK_MFMA_FP16 #include "uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16_itl.inc" @@ -504,7 +508,9 @@ struct FlatmmSn_32x128x512_1x4x1_16x16x32_FP16_itl : public FlatmmSn_32x128x512_ "v244", "v245", "v246", "v247", "v248", "v249", "v250", "v251", "v252", "v253", "v254", "v255" ); +#ifdef __clang__ #pragma clang diagnostic pop +#endif // clang-format on } }; diff --git a/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp b/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp index 8ad4a24ff5..5a4b194edb 100644 --- a/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp +++ b/include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp @@ -10,10 +10,10 @@ #include "ck_tile/ops/common.hpp" #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { struct FlatmmProblem { @@ -982,4 +982,6 @@ struct FlatmmKernel } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/flatmm/kernel/grouped_flatmm_kernel.hpp b/include/ck_tile/ops/flatmm/kernel/grouped_flatmm_kernel.hpp index d1feda700f..161ff8e0d0 100644 --- a/include/ck_tile/ops/flatmm/kernel/grouped_flatmm_kernel.hpp +++ b/include/ck_tile/ops/flatmm/kernel/grouped_flatmm_kernel.hpp @@ -10,10 +10,10 @@ #include "ck_tile/ops/common.hpp" #include "ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template , @@ -473,4 +473,6 @@ struct GroupedFlatmmKernel : FlatmmKernel= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp b/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp index 284ecb22f5..cc6385865c 100644 --- a/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp +++ b/include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp @@ -10,10 +10,10 @@ #include "ck_tile/ops/gemm/kernel/gemm_tile_partitioner.hpp" #include "ck_tile/host.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // #define disable_tile_gs namespace ck_tile { @@ -1499,4 +1499,6 @@ struct MoeFlatmmKernel } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/fmha/block/cast_tile_mx.hpp b/include/ck_tile/ops/fmha/block/cast_tile_mx.hpp index 4ed82f4ab4..a6925efa9a 100644 --- a/include/ck_tile/ops/fmha/block/cast_tile_mx.hpp +++ b/include/ck_tile/ops/fmha/block/cast_tile_mx.hpp @@ -56,9 +56,10 @@ cast_tile_mx(DstTensor& dst_tensor, DstScaleTensor& dst_scale_tensor, const SrcT // These builtins require the old value, and will generate a v_mov_b32 // vxxx [old] before cvt, which result in unwanted ISA so we prepare an // uninitialized variable x purposely, and turn off the warning +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wuninitialized" +#endif vec_t x; x = __builtin_amdgcn_cvt_scalef32_pk_fp4_f32( x, @@ -86,7 +87,9 @@ cast_tile_mx(DstTensor& dst_tensor, DstScaleTensor& dst_scale_tensor, const SrcT 3); // byte 3 dst_tensor.get_thread_buffer().template set_as( number{}, x); +#ifdef __clang__ #pragma clang diagnostic pop +#endif }); // Save scale for the corresponding lane @@ -184,9 +187,10 @@ cast_tile_mx(DstTensor& dst_tensor, DstScaleTensor& dst_scale_tensor, const SrcT // These builtins require the old value, and will generate a v_mov_b32 // vxxx [old] before cvt, which result in unwanted ISA so we prepare an // uninitialized variable x purposely, and turn off the warning +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wuninitialized" +#endif vec_t x; if constexpr(std::is_same_v) { @@ -220,7 +224,9 @@ cast_tile_mx(DstTensor& dst_tensor, DstScaleTensor& dst_scale_tensor, const SrcT } dst_tensor.get_thread_buffer().template set_as( number{}, x); +#ifdef __clang__ #pragma clang diagnostic pop +#endif }); // Save scale for the corresponding lane diff --git a/include/ck_tile/ops/fmha/block/page_block_navigator.hpp b/include/ck_tile/ops/fmha/block/page_block_navigator.hpp index 28d5796f9a..412d8bbd04 100644 --- a/include/ck_tile/ops/fmha/block/page_block_navigator.hpp +++ b/include/ck_tile/ops/fmha/block/page_block_navigator.hpp @@ -6,10 +6,10 @@ #include "ck_tile/core.hpp" #include "ck_tile/core/tensor/tile_window.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // assume that we have only 1 page-block/tensor view @@ -361,4 +361,6 @@ CK_TILE_HOST_DEVICE auto make_page_block_navigator(copy_const_t* } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/fmha/block/variants.hpp b/include/ck_tile/ops/fmha/block/variants.hpp index f287ec8382..0eb4ae28b4 100644 --- a/include/ck_tile/ops/fmha/block/variants.hpp +++ b/include/ck_tile/ops/fmha/block/variants.hpp @@ -20,10 +20,10 @@ #define CK_TILE_ATTENTION_USE_SOFTSIGN_ASM 0 #endif +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { namespace internal { __device__ inline float @@ -339,4 +339,6 @@ struct ComposedAttention } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp b/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp index 9e610d803f..2223b6bb32 100644 --- a/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp +++ b/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp @@ -10,10 +10,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #if !defined(CK_TILE_HAS_ROW_NEWBCAST) // row_newbcast (DPP modifier 0x157) support by architecture: // - Not supported: gfx908 (MI100) and older @@ -3130,4 +3130,6 @@ struct MoeSortingMultiPhaseKernel_P23 } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp index df77349ace..ba8188047b 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp @@ -15,10 +15,10 @@ #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief The GEMM kernel host arguments. @@ -175,4 +175,6 @@ struct GemmKernel }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm/kernel/gemm_multi_abd_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_multi_abd_kernel.hpp index 356377fba6..02a27968ec 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_multi_abd_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_multi_abd_kernel.hpp @@ -15,10 +15,10 @@ #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief The MultiABD GEMM kernel host arguments. @@ -200,4 +200,6 @@ struct GemmKernelMultiABD }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm/kernel/gemm_multi_d_kernel.hpp b/include/ck_tile/ops/gemm/kernel/gemm_multi_d_kernel.hpp index 1a6bcba7e8..9a03428d53 100644 --- a/include/ck_tile/ops/gemm/kernel/gemm_multi_d_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/gemm_multi_d_kernel.hpp @@ -15,10 +15,10 @@ #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp" #include "ck_tile/core/utility/type_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief The MultiD GEMM kernel host arguments. @@ -195,4 +195,6 @@ struct GemmKernelMultiD }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp b/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp index 1e33ba0bcb..d2c5e4598c 100644 --- a/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp +++ b/include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp @@ -14,10 +14,10 @@ #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief The Grouped GEMM kernel host arguments. @@ -583,4 +583,6 @@ struct GroupedGemmKernel } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm_mx/kernel/scale_pointer.hpp b/include/ck_tile/ops/gemm_mx/kernel/scale_pointer.hpp index c0a95fca05..0214f26a9c 100644 --- a/include/ck_tile/ops/gemm_mx/kernel/scale_pointer.hpp +++ b/include/ck_tile/ops/gemm_mx/kernel/scale_pointer.hpp @@ -5,10 +5,10 @@ #include "ck_tile/core.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { template @@ -115,4 +115,6 @@ struct MXScalePointer }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm_quant/block/block_gemm_quant_common.hpp b/include/ck_tile/ops/gemm_quant/block/block_gemm_quant_common.hpp index 4730439501..d4ff15a50e 100644 --- a/include/ck_tile/ops/gemm_quant/block/block_gemm_quant_common.hpp +++ b/include/ck_tile/ops/gemm_quant/block/block_gemm_quant_common.hpp @@ -5,10 +5,10 @@ #include "ck_tile/core.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { // Common utilities for quantized GEMM block operations @@ -228,4 +228,6 @@ struct AQPickerCommon : public BlockGemmQuantBase float scale_reg_f = 0.0f; }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp b/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp index b277780d6a..4c81718468 100644 --- a/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp +++ b/include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp @@ -14,10 +14,10 @@ #include "ck_tile/host/concat.hpp" #include "ck_tile/ops/gemm_quant/pipeline/tile_gemm_quant_traits.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { namespace detail { @@ -1578,4 +1578,6 @@ struct QuantGemmKernel }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp b/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp index 1e43836d57..46284c62ae 100644 --- a/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp +++ b/include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp @@ -15,10 +15,10 @@ #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief The Grouped GEMM kernel host arguments. @@ -650,4 +650,6 @@ struct QuantGroupedGemmKernel }; } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp index e526918ff7..b513f4fd75 100644 --- a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp +++ b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp @@ -7,10 +7,10 @@ #include "ck_tile/host/convolution_parameter.hpp" #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { enum class GroupedConvDirection @@ -376,4 +376,6 @@ CK_TILE_HOST SplitImagePieceInfo calculate_spatial_piece(ck_tile::index_t piece_ } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ops/pooling/kernel/pool_kernel.hpp b/include/ck_tile/ops/pooling/kernel/pool_kernel.hpp index 297e397b11..8f6ee14663 100644 --- a/include/ck_tile/ops/pooling/kernel/pool_kernel.hpp +++ b/include/ck_tile/ops/pooling/kernel/pool_kernel.hpp @@ -8,10 +8,10 @@ #include "ck_tile/ops/common.hpp" #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck_tile { /// @brief Host arguments for pooling operations @@ -580,4 +580,6 @@ struct PoolKernel } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/ck_tile/ref/naive_attention.hpp b/include/ck_tile/ref/naive_attention.hpp index f7d1da36e2..6e401b128f 100644 --- a/include/ck_tile/ref/naive_attention.hpp +++ b/include/ck_tile/ref/naive_attention.hpp @@ -9,9 +9,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck_tile { enum class naive_attention_layout_enum @@ -827,4 +828,6 @@ CK_TILE_HOST float naive_attention_fwd(naive_attention_fwd_traits t, #undef CK_TILE_DISPATCH_NAIVE_ATTEN_FWD_INTERNAL_ } // namespace ck_tile +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/include/rapidjson/document.h b/include/rapidjson/document.h index 0b12550a00..e937a3ca99 100644 --- a/include/rapidjson/document.h +++ b/include/rapidjson/document.h @@ -29,7 +29,7 @@ #endif RAPIDJSON_DIAG_PUSH -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_OFF(padded) RAPIDJSON_DIAG_OFF(switch - enum) RAPIDJSON_DIAG_OFF(c++ 98 - compat) diff --git a/include/rapidjson/encodedstream.h b/include/rapidjson/encodedstream.h index 4b96e79b7b..075dfc8691 100644 --- a/include/rapidjson/encodedstream.h +++ b/include/rapidjson/encodedstream.h @@ -23,7 +23,7 @@ RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(effc++) #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) #endif @@ -416,7 +416,7 @@ class AutoUTFOutputStream RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/error/en.h b/include/rapidjson/error/en.h index 9ea11e40ec..f11abe5cd7 100644 --- a/include/rapidjson/error/en.h +++ b/include/rapidjson/error/en.h @@ -17,7 +17,7 @@ #include "error.h" -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(switch - enum) RAPIDJSON_DIAG_OFF(covered - switch - default) @@ -266,7 +266,7 @@ GetPointerParseError_En(PointerParseErrorCode pointerParseErrorCode) RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/error/error.h b/include/rapidjson/error/error.h index 12ca3e085c..fcd308dc99 100644 --- a/include/rapidjson/error/error.h +++ b/include/rapidjson/error/error.h @@ -17,7 +17,7 @@ #include "../rapidjson.h" -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) #endif @@ -298,7 +298,7 @@ typedef const RAPIDJSON_ERROR_CHARTYPE* (*GetPointerParseErrorFunc)(PointerParse RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/filereadstream.h b/include/rapidjson/filereadstream.h index 8cdd792f44..a12d408ffb 100644 --- a/include/rapidjson/filereadstream.h +++ b/include/rapidjson/filereadstream.h @@ -18,7 +18,7 @@ #include "stream.h" #include -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) RAPIDJSON_DIAG_OFF(unreachable - code) @@ -116,7 +116,7 @@ class FileReadStream RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/filewritestream.h b/include/rapidjson/filewritestream.h index fd805d38d0..5ab74e646d 100644 --- a/include/rapidjson/filewritestream.h +++ b/include/rapidjson/filewritestream.h @@ -18,7 +18,7 @@ #include "stream.h" #include -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(unreachable - code) #endif @@ -128,7 +128,7 @@ inline void PutN(FileWriteStream& stream, char c, size_t n) RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/internal/diyfp.h b/include/rapidjson/internal/diyfp.h index 7dfe4aa550..354495088e 100644 --- a/include/rapidjson/internal/diyfp.h +++ b/include/rapidjson/internal/diyfp.h @@ -40,7 +40,7 @@ RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(effc++) #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) #endif @@ -264,7 +264,7 @@ inline DiyFp GetCachedPower10(int exp, int* outExp) RAPIDJSON_DIAG_POP #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP RAPIDJSON_DIAG_OFF(padded) #endif diff --git a/include/rapidjson/internal/regex.h b/include/rapidjson/internal/regex.h index 1a078f1ba0..6f86d673cc 100644 --- a/include/rapidjson/internal/regex.h +++ b/include/rapidjson/internal/regex.h @@ -19,7 +19,7 @@ #include "../stream.h" #include "stack.h" -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) RAPIDJSON_DIAG_OFF(switch - enum) diff --git a/include/rapidjson/istreamwrapper.h b/include/rapidjson/istreamwrapper.h index ad07e5ca3c..a2fe18a48b 100644 --- a/include/rapidjson/istreamwrapper.h +++ b/include/rapidjson/istreamwrapper.h @@ -19,7 +19,7 @@ #include #include -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) #elif defined(_MSC_VER) diff --git a/include/rapidjson/memorystream.h b/include/rapidjson/memorystream.h index 221b756d7f..7a6471ac3e 100644 --- a/include/rapidjson/memorystream.h +++ b/include/rapidjson/memorystream.h @@ -17,7 +17,7 @@ #include "stream.h" -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(unreachable - code) RAPIDJSON_DIAG_OFF(missing - noreturn) @@ -74,7 +74,7 @@ struct MemoryStream RAPIDJSON_NAMESPACE_END -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/ostreamwrapper.h b/include/rapidjson/ostreamwrapper.h index 547bd09fc7..f3dc362b11 100644 --- a/include/rapidjson/ostreamwrapper.h +++ b/include/rapidjson/ostreamwrapper.h @@ -18,7 +18,7 @@ #include "stream.h" #include -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) #endif @@ -89,7 +89,7 @@ class BasicOStreamWrapper typedef BasicOStreamWrapper OStreamWrapper; typedef BasicOStreamWrapper WOStreamWrapper; -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_POP #endif diff --git a/include/rapidjson/pointer.h b/include/rapidjson/pointer.h index bba4519774..387837c88b 100644 --- a/include/rapidjson/pointer.h +++ b/include/rapidjson/pointer.h @@ -20,7 +20,7 @@ #include "internal/itoa.h" #include "error/error.h" // PointerParseErrorCode -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(switch - enum) #elif defined(_MSC_VER) diff --git a/include/rapidjson/reader.h b/include/rapidjson/reader.h index fe4d6e3ec9..f580dfaf18 100644 --- a/include/rapidjson/reader.h +++ b/include/rapidjson/reader.h @@ -38,7 +38,7 @@ #include #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(old - style - cast) RAPIDJSON_DIAG_OFF(padded) diff --git a/include/rapidjson/schema.h b/include/rapidjson/schema.h index 8a542afc66..dca3f2d910 100644 --- a/include/rapidjson/schema.h +++ b/include/rapidjson/schema.h @@ -53,7 +53,7 @@ RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(effc++) #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_OFF(weak - vtables) RAPIDJSON_DIAG_OFF(exit - time - destructors) RAPIDJSON_DIAG_OFF(c++ 98 - compat - pedantic) diff --git a/include/rapidjson/writer.h b/include/rapidjson/writer.h index 634060020d..4a325d5bdb 100644 --- a/include/rapidjson/writer.h +++ b/include/rapidjson/writer.h @@ -37,7 +37,7 @@ #include #endif -#ifdef __clang__ +#if defined(__clang__) RAPIDJSON_DIAG_PUSH RAPIDJSON_DIAG_OFF(padded) RAPIDJSON_DIAG_OFF(unreachable - code) diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_avgpool_bwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_avgpool_bwd.hpp index c61ab5eae7..76fbf9d922 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_avgpool_bwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_avgpool_bwd.hpp @@ -10,10 +10,10 @@ #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -357,4 +357,6 @@ struct ReferenceAvgPoolBwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_batched_gemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_batched_gemm.hpp index 7e219f12d5..de23d18d15 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_batched_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_batched_gemm.hpp @@ -10,10 +10,10 @@ #include "ck/library/utility/host_tensor.hpp" #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -426,4 +426,6 @@ struct ReferenceBatchedGemm_GQA : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_cgemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_cgemm.hpp index 23b8d69ba7..1897271c5b 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_cgemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_cgemm.hpp @@ -10,10 +10,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -187,4 +187,6 @@ struct ReferenceCGemm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_column_to_image.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_column_to_image.hpp index d193f08a99..19f1f43e21 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_column_to_image.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_column_to_image.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/tensor_layout.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -371,4 +371,6 @@ struct ReferenceColumnToImage : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp index c4164a3e18..3d2eb46649 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_contraction.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/element/element_wise_operation.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -512,4 +512,6 @@ struct ReferenceBatchedContraction_G1_M3_N2_K1 : public ck::tensor_operation::de } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp index 572855d4d4..da26103ada 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_data.hpp @@ -10,10 +10,10 @@ #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -511,4 +511,6 @@ struct ReferenceConvBwdData : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp index cc5ca29c6d..d96e7b16d3 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_bwd_weight.hpp @@ -10,10 +10,10 @@ #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -456,4 +456,6 @@ struct ReferenceConvBwdWeight : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp index 6f6d0f464e..c42676b209 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd.hpp @@ -23,10 +23,10 @@ #include "ck/library/utility/convolution_parameter.hpp" #include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -572,4 +572,6 @@ struct ReferenceConvFwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation.hpp index 6daa0d0ba1..475d6e1fc2 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation.hpp @@ -9,10 +9,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -195,4 +195,6 @@ struct ReferenceConvFwd_Bias_Activation : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation_add.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation_add.hpp index 61a44acc28..b6c86abf7c 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation_add.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_conv_fwd_bias_activation_add.hpp @@ -9,10 +9,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -203,4 +203,6 @@ struct ReferenceConvFwd_Bias_Activation_Add : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_elementwise.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_elementwise.hpp index fe2cb4be4a..d6be12e20f 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_elementwise.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_elementwise.hpp @@ -10,10 +10,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -113,4 +113,6 @@ struct ReferenceElementwise : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_fpAintB_gemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_fpAintB_gemm.hpp index cf8fb45179..26cecfc0fa 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_fpAintB_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_fpAintB_gemm.hpp @@ -10,10 +10,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -149,4 +149,6 @@ struct ReferencefpAintBGemm : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm.hpp index b9dc834e80..8073b41bd2 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -251,4 +251,6 @@ struct ReferenceGemm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_layernorm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_layernorm.hpp index 33ef0b97cc..304a13644e 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_layernorm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_layernorm.hpp @@ -7,10 +7,10 @@ #include #include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -239,4 +239,6 @@ struct ReferenceGemmLayernorm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multi_abd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multi_abd.hpp index b55b1dbf99..490b808399 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multi_abd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multi_abd.hpp @@ -14,10 +14,10 @@ #include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -197,4 +197,6 @@ struct ReferenceGemmMultiABD : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multiple_d.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multiple_d.hpp index 8aeb0069b7..1ec311c7df 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multiple_d.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_gemm_multiple_d.hpp @@ -10,10 +10,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -160,4 +160,6 @@ struct ReferenceGemmMultipleD : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm.hpp index 2e1e4e0261..3b4764721e 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -215,4 +215,6 @@ struct ReferenceGroupnorm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm_bwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm_bwd.hpp index 57c9c16c77..7d89ea6fd8 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm_bwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_groupnorm_bwd.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -235,4 +235,6 @@ struct ReferenceGroupnormBwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_image_to_column.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_image_to_column.hpp index d935c766e7..33910f5c91 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_image_to_column.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_image_to_column.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/numeric.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -335,4 +335,6 @@ struct ReferenceImageToColumn : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm.hpp index a9c372510b..aa4837aac6 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -263,4 +263,6 @@ struct ReferenceLayernorm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm_bwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm_bwd.hpp index 5564bb7249..428b944d15 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm_bwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_layernorm_bwd.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -204,4 +204,6 @@ struct ReferenceLayernormBwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_maxpool_bwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_maxpool_bwd.hpp index cd3fb7922e..157b6e56e5 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_maxpool_bwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_maxpool_bwd.hpp @@ -11,10 +11,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -115,4 +115,6 @@ struct ReferenceMaxPoolBwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp index 843128ce97..9ebadcd665 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -297,4 +297,6 @@ struct ReferenceMoeGemm : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale.hpp index aa5a22b305..0b2c9152ee 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -282,4 +282,6 @@ struct ReferenceMoeGemm1BlockScale : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale_splitk.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale_splitk.hpp index 1f152c6aae..0755cf3cd7 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale_splitk.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm1_blockscale_splitk.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -234,4 +234,6 @@ struct ReferenceMoeGemm1BlockScaleSplitK : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2.hpp index 451528fb50..074a0b6f82 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2.hpp @@ -12,10 +12,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -270,4 +270,6 @@ struct ReferenceMoeGemm2 : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2_blockscale.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2_blockscale.hpp index 8f2a35bd26..f672546a6b 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2_blockscale.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_gemm2_blockscale.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -250,4 +250,6 @@ struct ReferenceMoeGemm2BlockScale : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm1.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm1.hpp index 5911f539b1..4d50cb6af0 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm1.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm1.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -266,4 +266,6 @@ struct ReferenceMoeMXGemm1 : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm2.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm2.hpp index 88e5c6bc3e..0d3022e517 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm2.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_moe_mx_gemm2.hpp @@ -11,10 +11,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/library/utility/host_tensor.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -243,4 +243,6 @@ struct ReferenceMoeMXGemm2 : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_mx_gemm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_mx_gemm.hpp index 1732cef098..caa962e0e1 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_mx_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_mx_gemm.hpp @@ -11,10 +11,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/reference_tensor_operation/cpu/reference_gemm.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -239,4 +239,6 @@ struct ReferenceMXGemm : public device::BaseOperator } // namespace host } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp index 8bb217ff83..d14da641ab 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_pool_fwd.hpp @@ -14,10 +14,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -363,4 +363,6 @@ struct ReferencePoolingFwd : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/cpu/reference_sparse_embedding3_forward_layernorm.hpp b/library/include/ck/library/reference_tensor_operation/cpu/reference_sparse_embedding3_forward_layernorm.hpp index 9d4e34ed4e..b050a35337 100644 --- a/library/include/ck/library/reference_tensor_operation/cpu/reference_sparse_embedding3_forward_layernorm.hpp +++ b/library/include/ck/library/reference_tensor_operation/cpu/reference_sparse_embedding3_forward_layernorm.hpp @@ -12,10 +12,10 @@ #include "ck/library/utility/host_tensor.hpp" #include "ck/library/utility/host_tensor_generator.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace host { @@ -208,4 +208,6 @@ struct ReferenceSparseEmbedding3ForwardLayernorm : public device::BaseOperator } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/library/include/ck/library/reference_tensor_operation/gpu/reference_gemm.hpp b/library/include/ck/library/reference_tensor_operation/gpu/reference_gemm.hpp index 3155428e57..b366659e7f 100644 --- a/library/include/ck/library/reference_tensor_operation/gpu/reference_gemm.hpp +++ b/library/include/ck/library/reference_tensor_operation/gpu/reference_gemm.hpp @@ -10,9 +10,10 @@ #include "ck/tensor_operation/gpu/device/device_base.hpp" #include "ck/host_utility/kernel_launch.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif namespace ck { template = 23 #pragma clang diagnostic pop +#endif diff --git a/library/src/tensor_operation_instance/gpu/gemm_multi_abd/device_gemm_wmma_multi_abd_gelu_bf16_i8_bf16_mk_kn_mn_v1_instance.cpp b/library/src/tensor_operation_instance/gpu/gemm_multi_abd/device_gemm_wmma_multi_abd_gelu_bf16_i8_bf16_mk_kn_mn_v1_instance.cpp index d93eb1f89e..5893116f23 100644 --- a/library/src/tensor_operation_instance/gpu/gemm_multi_abd/device_gemm_wmma_multi_abd_gelu_bf16_i8_bf16_mk_kn_mn_v1_instance.cpp +++ b/library/src/tensor_operation_instance/gpu/gemm_multi_abd/device_gemm_wmma_multi_abd_gelu_bf16_i8_bf16_mk_kn_mn_v1_instance.cpp @@ -12,10 +12,10 @@ #include "ck/library/tensor_operation_instance/add_device_operation_instance.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif namespace ck { namespace tensor_operation { namespace device { @@ -61,4 +61,6 @@ void add_device_gemm_wmma_multi_abd_bf16_i8_bf16_mk_kn_mn_gelu_v1_instances( } // namespace device } // namespace tensor_operation } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/profiler/include/profiler/profile_contraction_utils.hpp b/profiler/include/profiler/profile_contraction_utils.hpp index e5cc521178..88d4af2f0b 100644 --- a/profiler/include/profiler/profile_contraction_utils.hpp +++ b/profiler/include/profiler/profile_contraction_utils.hpp @@ -7,9 +7,10 @@ #include "ck/ck.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif using Row = ck::tensor_layout::gemm::RowMajor; using Col = ck::tensor_layout::gemm::ColumnMajor; @@ -84,4 +85,6 @@ assign_default_strides(Col, std::vector& strides, std::vector= 23 #pragma clang diagnostic pop +#endif diff --git a/profiler/src/profiler_operation_registry.hpp b/profiler/src/profiler_operation_registry.hpp index 262793d09e..49e27f36e9 100644 --- a/profiler/src/profiler_operation_registry.hpp +++ b/profiler/src/profiler_operation_registry.hpp @@ -9,10 +9,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif class ProfilerOperationRegistry final { ProfilerOperationRegistry() = default; @@ -87,4 +87,6 @@ class ProfilerOperationRegistry final ::ProfilerOperationRegistry::GetInstance().Add(name, description, operation) \ _Pragma("clang diagnostic pop") // clang-format on +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/ck_tile/flatmm/test_mx_flatmm_base.hpp b/test/ck_tile/flatmm/test_mx_flatmm_base.hpp index 919a295d52..8ce3fb5a86 100644 --- a/test/ck_tile/flatmm/test_mx_flatmm_base.hpp +++ b/test/ck_tile/flatmm/test_mx_flatmm_base.hpp @@ -203,9 +203,10 @@ class TestMXFlatmmBase : public ::testing::Test // mx_flatmm_calc is explicitly instantiated in the linked object library; // suppress the -Wundefined-func-template warning that fires when the // compiler sees only the forward declaration in mx_flatmm.hpp. +#ifdef __clang__ #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wundefined-func-template" +#endif BaseFlatmmPipeline::template TailHandler( [&](auto has_hot_loop_, auto tail_num_) { constexpr auto has_hot_loop_v = has_hot_loop_.value; @@ -232,7 +233,9 @@ class TestMXFlatmmBase : public ::testing::Test }, has_hot_loop, tail_num); +#ifdef __clang__ #pragma clang diagnostic pop +#endif c_dev_buf.FromDevice(c_rslt_host.data()); diff --git a/test/common/csv_test_loader.hpp b/test/common/csv_test_loader.hpp index ac3ac5f97e..43d94049c0 100644 --- a/test/common/csv_test_loader.hpp +++ b/test/common/csv_test_loader.hpp @@ -12,9 +12,10 @@ #include "ck/library/utility/convolution_parameter.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif namespace ck { namespace test { @@ -247,4 +248,6 @@ inline bool load_and_populate_test_cases(const std::vector& csv_pat } // namespace test } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_dataset_xdl.cpp b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_dataset_xdl.cpp index 28795d1415..a779ae1107 100644 --- a/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_dataset_xdl.cpp +++ b/test/grouped_convnd_bwd_data/test_grouped_convnd_bwd_data_dataset_xdl.cpp @@ -15,9 +15,10 @@ using namespace ck::tensor_layout::convolution; // Import tensor layout names (G static ck::index_t param_mask = 0xffff; static ck::index_t instance_index = -1; +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif // Load CSV data for 2D tests static std::vector Get2DTestCases() { @@ -343,4 +344,6 @@ int main(int argc, char** argv) } return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_dataset_xdl.cpp b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_dataset_xdl.cpp index e78660325a..7dadcbd4e2 100644 --- a/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_dataset_xdl.cpp +++ b/test/grouped_convnd_bwd_weight/test_grouped_convnd_bwd_weight_dataset_xdl.cpp @@ -15,9 +15,10 @@ #include "profiler/profile_grouped_conv_bwd_weight_impl.hpp" // The actual GPU profiler that does convolution work #include "../common/csv_test_loader.hpp" // Shared CSV test case loader +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif using namespace ck::tensor_layout::convolution; static ck::index_t param_mask = 0xffff; static ck::index_t instance_index = -1; @@ -282,4 +283,6 @@ int main(int argc, char** argv) } return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp index ffc9067c12..3c4a77f4bd 100644 --- a/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp +++ b/test/grouped_convnd_fwd/test_grouped_convnd_fwd_dataset_xdl.cpp @@ -11,9 +11,10 @@ #include "profiler/profile_grouped_conv_fwd_impl.hpp" // The actual GPU profiler that does convolution work #include "../common/csv_test_loader.hpp" // Shared CSV test case loader +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif using namespace ck::tensor_layout::convolution; // Import tensor layout names (NHWGC, GKYXC, etc.) static ck::index_t param_mask = 0xffff; static ck::index_t instance_index = -1; @@ -196,4 +197,6 @@ int main(int argc, char** argv) } return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_gemm/test_grouped_gemm_fixed_nk_bias.cpp b/test/grouped_gemm/test_grouped_gemm_fixed_nk_bias.cpp index 07c107d6c8..265f6e33e6 100644 --- a/test/grouped_gemm/test_grouped_gemm_fixed_nk_bias.cpp +++ b/test/grouped_gemm/test_grouped_gemm_fixed_nk_bias.cpp @@ -18,9 +18,10 @@ #include #include +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif static ck::index_t param_mask = 0xffffff; using FP32 = float; @@ -305,4 +306,6 @@ int main(int argc, char** argv) } return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_gemm/test_grouped_gemm_multi_abd_fixed_nk.cpp b/test/grouped_gemm/test_grouped_gemm_multi_abd_fixed_nk.cpp index 70ca4b7705..b6be6e6f69 100644 --- a/test/grouped_gemm/test_grouped_gemm_multi_abd_fixed_nk.cpp +++ b/test/grouped_gemm/test_grouped_gemm_multi_abd_fixed_nk.cpp @@ -15,9 +15,10 @@ #include "gtest/gtest.h" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif using FP32 = float; using FP16 = ck::half_t; using BF16 = ck::bhalf_t; @@ -240,4 +241,6 @@ int main(int argc, char** argv) testing::InitGoogleTest(&argc, argv); return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_gemm/test_grouped_gemm_util.hpp b/test/grouped_gemm/test_grouped_gemm_util.hpp index fd835bb5aa..ca5932093e 100644 --- a/test/grouped_gemm/test_grouped_gemm_util.hpp +++ b/test/grouped_gemm/test_grouped_gemm_util.hpp @@ -18,10 +18,10 @@ #include "profiler/profile_grouped_gemm_impl.hpp" #include "profiler/profile_grouped_gemm_fixed_nk_impl.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" - +#endif static ck::index_t param_mask = 0xffffff; static ck::index_t instance_index = -1; @@ -346,4 +346,6 @@ int main(int argc, char** argv) } return RUN_ALL_TESTS(); } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/grouped_gemm_tile_loop/test_grouped_gemm_tile_loop_util.hpp b/test/grouped_gemm_tile_loop/test_grouped_gemm_tile_loop_util.hpp index dd31a02cec..fabd74daf6 100644 --- a/test/grouped_gemm_tile_loop/test_grouped_gemm_tile_loop_util.hpp +++ b/test/grouped_gemm_tile_loop/test_grouped_gemm_tile_loop_util.hpp @@ -19,9 +19,10 @@ #include "example/ck_tile/17_grouped_gemm/grouped_gemm_multi_d.hpp" #include "profiler/profile_grouped_gemm_tile_loop_generic_impl.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-invalidation" +#endif extern ck::index_t param_mask; extern ck::index_t instance_index; @@ -174,4 +175,6 @@ class TestGroupedGemmTileLoop : public testing::Test } // namespace test } // namespace ck +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/test/position_embedding/position_embedding.cpp b/test/position_embedding/position_embedding.cpp index 04b77be9cd..d72c0fff03 100644 --- a/test/position_embedding/position_embedding.cpp +++ b/test/position_embedding/position_embedding.cpp @@ -9,10 +9,10 @@ #include "ck_tile/core.hpp" #include "ck_tile/ops/fmha.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #ifndef TEST_ALIBI_VERBOSE #define TEST_ALIBI_VERBOSE 0 #endif @@ -217,4 +217,6 @@ int main() // clang-format on return rtn ? 0 : -1; } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp index 24cfebdf65..d1345c527e 100644 --- a/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_multi_d/gemm_multi_d_benchmark.hpp @@ -13,10 +13,10 @@ #include "ck_tile/host.hpp" #include "gemm/gemm_benchmark.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts struct GemmMultiDProblem : GemmProblem @@ -74,4 +74,6 @@ void gemm_multi_d_host_reference(int verify, a_m_k, b_k_n, {d0_m_n, d1_m_n}, c_m_n_host_result); } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp index 8b9d95a087..beebc0ffc6 100644 --- a/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_preshuffle/gemm_preshuffle_benchmark.hpp @@ -2,10 +2,10 @@ // SPDX-License-Identifier: MIT #pragma once +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif #include "ck_tile/core.hpp" #include "ck_tile/host.hpp" #include "gemm_preshuffle_common.hpp" @@ -70,4 +70,6 @@ void gemm_host_reference(int verify, c_m_n_gpu_buf_ref.FromDevice(c_m_n_ref.data()); } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp index 64e8b3e226..23338a6cd0 100644 --- a/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp +++ b/tile_engine/ops/gemm/gemm_universal/gemm_universal_benchmark.hpp @@ -13,9 +13,10 @@ #include "ck_tile/host.hpp" #include "gemm/gemm_benchmark.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts @@ -67,4 +68,6 @@ void gemm_host_reference(int verify, c_m_n_gpu_buf_ref.FromDevice(c_m_n_host_result.data()); } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/tile_engine/ops/gemm/grouped_gemm/grouped_gemm_benchmark.hpp b/tile_engine/ops/gemm/grouped_gemm/grouped_gemm_benchmark.hpp index 521e0c1e21..3cbeb6caba 100644 --- a/tile_engine/ops/gemm/grouped_gemm/grouped_gemm_benchmark.hpp +++ b/tile_engine/ops/gemm/grouped_gemm/grouped_gemm_benchmark.hpp @@ -15,9 +15,10 @@ #include "ck_tile/host.hpp" #include "grouped_gemm_common.hpp" +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" +#endif // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts @@ -264,4 +265,6 @@ void gemm_host_reference_grouped(int verify, } } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif diff --git a/tile_engine/ops/gemm_streamk/gemm_streamk_benchmark.hpp b/tile_engine/ops/gemm_streamk/gemm_streamk_benchmark.hpp index 9097dbac78..f92910b40a 100644 --- a/tile_engine/ops/gemm_streamk/gemm_streamk_benchmark.hpp +++ b/tile_engine/ops/gemm_streamk/gemm_streamk_benchmark.hpp @@ -17,10 +17,10 @@ // Data types and Layouts are defined by the generated kernel headers // No hardcoded type definitions here to avoid conflicts +#if __clang_major__ >= 23 #pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wno-unknown-warning-option" #pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions" - +#endif enum class Metric { LATENCY = 0, @@ -203,4 +203,6 @@ void gemm_host_reference(int verify, c_m_n_gpu_buf_ref.FromDevice(c_m_n_host_result.data()); } } +#if __clang_major__ >= 23 #pragma clang diagnostic pop +#endif