[CK] Fix/suppress clang lifetimebound warnings with staging compiler. (#6550)

## Motivation

New changes from upstream llvm-project cause an avalanche of warnings in
CK. Gonna disable them by ignoring the
lifetime-safety-intra-tu-suggestions flag until a better permanent
solution is found.

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

## Test Result

<!-- Briefly summarize test outcomes. -->

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

include/ck_tile/core/arch/mma/mma_pipeline.hpp

@@ -9,6 +9,9 @@
 #include "mma_traits.hpp"
 #include "mma_transforms.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile::core::arch::mma {
 
 /*! @enum MmaPipelineOptionFlag
@@ -341,3 +344,5 @@ concept MmaPipelineInterface = std::derived_from<Derived, MmaPipelineBase<Flags,
 #endif // CK_TILE_CONCEPTS && CK_TILE_CONCEPTS_HEADER
 
 } // namespace ck_tile::core::arch::mma
+
+#pragma clang diagnostic pop

include/ck_tile/core/arch/mma/mma_transforms.hpp

@@ -12,7 +12,7 @@ namespace ck_tile::core::arch::mma {
 struct PassThroughTransform
 {
     template <typename VecType>
-    CK_TILE_DEVICE static decltype(auto) exec(VecType&& v)
+    CK_TILE_DEVICE static decltype(auto) exec([[clang::lifetimebound]] VecType&& v)
     {
         return std::forward<VecType>(v);
     }

include/ck_tile/core/arch/workgroup_barrier.hpp

@@ -10,7 +10,7 @@ namespace ck_tile {
 
 struct workgroup_barrier
 {
-    CK_TILE_DEVICE workgroup_barrier(uint32_t* ptr) : base_ptr(ptr) {}
+    CK_TILE_DEVICE workgroup_barrier([[clang::lifetimebound]] uint32_t* ptr) : base_ptr(ptr) {}
 
     CK_TILE_DEVICE uint32_t ld(uint32_t offset = 0)
     {

include/ck_tile/core/container/span.hpp

@@ -8,6 +8,9 @@
 #include <array>
 #include <type_traits>
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 // implement the c++20 std::span, lightweight, non-owning reference to a sequence
@@ -76,3 +79,4 @@ class span
 };
 
 } // namespace ck_tile
+#pragma clang diagnostic pop

include/ck_tile/core/tensor/buffer_view.hpp

@@ -17,6 +17,9 @@
 #include "ck_tile/core/utility/type_traits.hpp"
 #include "ck_tile/core/utility/ignore.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 // T may be scalar or vector
@@ -240,7 +243,8 @@ struct buffer_view<address_space_enum::global,
     {
     }
 
-    CK_TILE_HOST_DEVICE constexpr buffer_view(T* __restrict__ p_data, BufferSizeType buffer_size)
+    CK_TILE_HOST_DEVICE constexpr buffer_view([[clang::lifetimebound]] T* __restrict__ p_data,
+                                              BufferSizeType buffer_size)
         : p_data_{p_data},
           buffer_size_{buffer_size / PackedSize},
           cached_buf_res_{0},
@@ -770,12 +774,13 @@ struct buffer_view<address_space_enum::lds,
     {
     }
 
-    CK_TILE_HOST_DEVICE constexpr buffer_view(T* __restrict__ p_data, BufferSizeType buffer_size)
+    CK_TILE_HOST_DEVICE constexpr buffer_view([[clang::lifetimebound]] T* __restrict__ p_data,
+                                              BufferSizeType buffer_size)
         : p_data_{p_data}, buffer_size_{buffer_size}, invalid_element_value_{0}
     {
     }
 
-    CK_TILE_HOST_DEVICE constexpr buffer_view(T* __restrict__ p_data,
+    CK_TILE_HOST_DEVICE constexpr buffer_view([[clang::lifetimebound]] T* __restrict__ p_data,
                                               BufferSizeType buffer_size,
                                               T invalid_element_value)
         : p_data_{p_data}, buffer_size_{buffer_size}, invalid_element_value_{invalid_element_value}
@@ -1284,7 +1289,8 @@ template <address_space_enum BufferAddressSpace,
           amd_buffer_coherence_enum Coherence = amd_buffer_coherence_enum::coherence_default,
           typename T,
           typename BufferSizeType>
-CK_TILE_HOST_DEVICE constexpr auto make_buffer_view(T* __restrict__ p, BufferSizeType buffer_size)
+CK_TILE_HOST_DEVICE constexpr auto make_buffer_view([[clang::lifetimebound]] T* __restrict__ p,
+                                                    BufferSizeType buffer_size)
 {
     return buffer_view<BufferAddressSpace, T, BufferSizeType, true, Coherence>{p, buffer_size};
 }
@@ -1325,3 +1331,5 @@ CK_TILE_HOST_DEVICE void print(const buffer_view<BufferAddressSpace,
 }
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/core/utility/type_traits.hpp

@@ -103,7 +103,7 @@ template <
     typename PY,
     typename PX,
     typename std::enable_if<std::is_pointer_v<PY> && std::is_pointer_v<PX>, bool>::type = false>
-CK_TILE_HOST_DEVICE PY c_style_pointer_cast(PX p_x)
+CK_TILE_HOST_DEVICE PY c_style_pointer_cast([[clang::lifetimebound]] PX p_x)
 {
 #pragma clang diagnostic push
 #pragma clang diagnostic ignored "-Wold-style-cast"

include/ck_tile/ops/batched_contraction/kernel/batched_contraction_kernel.hpp

@@ -8,6 +8,9 @@
 #include "ck_tile/ops/batched_contraction/utils/tensor_descriptor_utils.hpp"
 #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 /**
  * @file batched_contraction_kernel.hpp
  * @brief Batched Tensor Contraction Operations
@@ -687,3 +690,5 @@ struct BatchedContractionKernel
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp

@@ -10,6 +10,9 @@
 #include "ck_tile/ops/common.hpp"
 #include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 struct FlatmmProblem
 {
@@ -970,3 +973,5 @@ struct FlatmmKernel
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/flatmm/kernel/grouped_flatmm_kernel.hpp

@@ -10,6 +10,9 @@
 #include "ck_tile/ops/common.hpp"
 #include "ck_tile/ops/flatmm/kernel/flatmm_kernel.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 template <class ScaleM       = FlatmmScalePointer<-1>,
@@ -468,3 +471,5 @@ struct GroupedFlatmmKernel : FlatmmKernel<TilePartitioner_, FlatmmPipeline_, Epi
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/flatmm/kernel/moe_flatmm_kernel.hpp

@@ -10,6 +10,9 @@
 #include "ck_tile/ops/gemm/kernel/gemm_tile_partitioner.hpp"
 #include "ck_tile/host.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 // #define disable_tile_gs
 
 namespace ck_tile {
@@ -1493,3 +1496,5 @@ struct MoeFlatmmKernel
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/fmha/block/page_block_navigator.hpp

@@ -6,6 +6,9 @@
 #include "ck_tile/core.hpp"
 #include "ck_tile/core/tensor/tile_window.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 // assume that we have only 1 page-block/tensor view
@@ -356,3 +359,5 @@ CK_TILE_HOST_DEVICE auto make_page_block_navigator(copy_const_t<DataType, void>*
 }
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/fmha/block/variants.hpp

@@ -19,6 +19,9 @@
 #define CK_TILE_ATTENTION_USE_SOFTSIGN_ASM 0
 #endif
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 namespace internal {
 __device__ inline float
@@ -333,3 +336,5 @@ struct ComposedAttention
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp

@@ -10,6 +10,9 @@
 #include <string>
 #include <type_traits>
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 #if !defined(CK_TILE_HAS_ROW_NEWBCAST)
 // row_newbcast (DPP modifier 0x157) support by architecture:
 // - Not supported: gfx908 (MI100) and older
@@ -3125,3 +3128,5 @@ struct MoeSortingMultiPhaseKernel_P23
 #undef MOE_SORTING_MOCK_ID
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp

@@ -15,6 +15,9 @@
 #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief The GEMM kernel host arguments.
@@ -168,3 +171,5 @@ struct GemmKernel
     }
 };
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm/kernel/gemm_multi_abd_kernel.hpp

@@ -15,6 +15,9 @@
 #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief The MultiABD GEMM kernel host arguments.
@@ -195,3 +198,5 @@ struct GemmKernelMultiABD
     }
 };
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm/kernel/gemm_multi_d_kernel.hpp

@@ -15,6 +15,9 @@
 #include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp"
 #include "ck_tile/core/utility/type_traits.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief The MultiD GEMM kernel host arguments.
@@ -190,3 +193,5 @@ struct GemmKernelMultiD
     }
 };
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm/kernel/grouped_gemm_kernel.hpp

@@ -14,6 +14,9 @@
 
 #include <hip/hip_runtime.h>
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief The Grouped GEMM kernel host arguments.
@@ -575,3 +578,5 @@ struct GroupedGemmKernel
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp

@@ -36,7 +36,7 @@ struct UniversalGemmHostArgs
         const std::array<const void*, NumATensor>& as_ptr_,
         const std::array<const void*, NumBTensor>& bs_ptr_,
         const std::array<const void*, NumDTensor>& ds_ptr_,
-        void* e_ptr_,
+        [[clang::lifetimebound]] void* e_ptr_,
         index_t k_batch_,
         index_t M_,
         index_t N_,

include/ck_tile/ops/gemm_mx/kernel/scale_pointer.hpp

@@ -5,6 +5,9 @@
 
 #include "ck_tile/core.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 template <typename ScaleType, int SharedGranularityMN, int SharedGranularityK = 0>
@@ -111,3 +114,4 @@ struct MXScalePointer<ScaleType, -1, 0>
 };
 
 } // namespace ck_tile
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm_quant/block/block_gemm_quant_common.hpp

@@ -5,6 +5,9 @@
 
 #include "ck_tile/core.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 // Common utilities for quantized GEMM block operations
@@ -224,3 +227,4 @@ struct AQPickerCommon : public BlockGemmQuantBase
     float scale_reg_f = 0.0f;
 };
 } // namespace ck_tile
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm_quant/kernel/gemm_quant_kernel.hpp

@@ -14,6 +14,9 @@
 #include "ck_tile/host/concat.hpp"
 #include "ck_tile/ops/gemm_quant/pipeline/tile_gemm_quant_traits.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 namespace detail {
@@ -1574,3 +1577,4 @@ struct QuantGemmKernel
 };
 
 } // namespace ck_tile
+#pragma clang diagnostic pop

include/ck_tile/ops/gemm_quant/kernel/grouped_gemm_quant_kernel.hpp

@@ -15,6 +15,9 @@
 
 #include <hip/hip_runtime.h>
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief The Grouped GEMM kernel host arguments.
@@ -646,3 +649,4 @@ struct QuantGroupedGemmKernel
 };
 
 } // namespace ck_tile
+#pragma clang diagnostic pop

include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp

@@ -7,6 +7,9 @@
 #include "ck_tile/host/convolution_parameter.hpp"
 #include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp"
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 enum class GroupedConvDirection
@@ -261,3 +264,5 @@ CK_TILE_HOST SplitImagePieceInfo calculate_spatial_piece(ck_tile::index_t piece_
 }
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop

include/ck_tile/ops/pooling/kernel/pool_kernel.hpp

@@ -8,6 +8,9 @@
 #include "ck_tile/ops/common.hpp"
 #include <type_traits>
 
+#pragma clang diagnostic push
+#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
+
 namespace ck_tile {
 
 /// @brief Host arguments for pooling operations
@@ -575,3 +578,5 @@ struct PoolKernel
 };
 
 } // namespace ck_tile
+
+#pragma clang diagnostic pop