ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-11 17:00:18 +00:00
Code Issues Packages Projects Releases Wiki Activity

[rocm-libraries] ROCm/rocm-libraries#5921 (commit 032ac1b)

Browse Source 
[CK] fix clang lifetimebound errors with staging compiler
 (#5921)

## Motivation

The ROCm staging compiler (newer Clang) enforces
`[[clang::lifetimebound]]` annotations on methods that return references
or pointers to internal object data. Without these annotations, the
staging compiler emits compilation errors for container accessor methods
across the CK and CK Tile namespaces.

  ## Technical Details

Adds `[[clang::lifetimebound]]` to all reference/pointer-returning
accessors in core container types:

  **`ck::` namespace:**
  - `Array` -- `At()`, `operator[]`, `operator()`, `begin()`, `end()`
  - `index_array` -- `operator[]`
  - `StaticallyIndexedArray_v2` -- `At()`, `operator[]`, `operator()`
  - `IndexLookupTable` -- `operator[]`

  **`ck_tile::` namespace:**
  - `array` -- `get(i)`, `at()`, `operator[]`, `operator()`
  - `static_array` -- `operator[]`
  - `thread_buffer` -- `get(i)`, `at()`, `operator[]`, `operator()`
  - `make_kernel()` -- parameter pack

Also removes the unused `instance_index` variable from
`batched_gemm_reduce_fp16.cpp` and simplifies its argument parsing
  accordingly.

  ## Test Plan

- Compile with the staging compiler to verify all lifetimebound errors
are resolved
- Existing tests pass unchanged -- the attribute is a compile-time
annotation with no runtime effect

 ## Test Result

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

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
This commit is contained in:
Illia Silin
2026-03-30 14:20:20 +00:00
committed by assistant-librarian[bot]
parent 2dcae9d173
commit e6b8094f94
12 changed files with 82 additions and 71 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
include/ck/tensor_description/tensor_space_filling_curve.hpp
View File

@@ -24,6 +24,7 @@ struct IndexLookupTable
MultiIndex<nDim> data[NumAccesses > 0 ? NumAccesses : 1];
__host__ __device__ constexpr const MultiIndex<nDim>& operator[](index_t i) const
[[clang::lifetimebound]]
{
return data[i];
}

31
include/ck/utility/array.hpp
View File

@@ -24,13 +24,22 @@ struct Array
__host__ __device__ static constexpr index_t Size() { return NSize; }
__host__ __device__ constexpr const TData& At(index_t i) const { return mData[i]; }
__host__ __device__ constexpr const TData& At(index_t i) const [[clang::lifetimebound]]
{
return mData[i];
}
__host__ __device__ constexpr TData& At(index_t i) { return mData[i]; }
__host__ __device__ constexpr TData& At(index_t i) [[clang::lifetimebound]] { return mData[i]; }
__host__ __device__ constexpr const TData& operator[](index_t i) const { return At(i); }
__host__ __device__ constexpr const TData& operator[](index_t i) const [[clang::lifetimebound]]
{
return At(i);
}
__host__ __device__ constexpr TData& operator()(index_t i) { return At(i); }
__host__ __device__ constexpr TData& operator()(index_t i) [[clang::lifetimebound]]
{
return At(i);
}
template <typename... Args>
__host__ constexpr auto Emplace(index_t i, Args&&... args)
@@ -50,10 +59,16 @@ struct Array
return *this;
}
__host__ __device__ constexpr const TData* begin() const { return &mData[0]; }
__host__ __device__ constexpr const TData* end() const { return &mData[NSize]; }
__host__ __device__ constexpr TData* begin() { return &mData[0]; }
__host__ __device__ constexpr TData* end() { return &mData[NSize]; }
__host__ __device__ constexpr const TData* begin() const [[clang::lifetimebound]]
{
return &mData[0];
}
__host__ __device__ constexpr const TData* end() const [[clang::lifetimebound]]
{
return &mData[NSize];
}
__host__ __device__ constexpr TData* begin() [[clang::lifetimebound]] { return &mData[0]; }
__host__ __device__ constexpr TData* end() [[clang::lifetimebound]] { return &mData[NSize]; }
};
// empty Array

11
include/ck/utility/sequence.hpp
View File

@@ -403,8 +403,15 @@ struct index_array
{
index_t data[N > 0 ? N : 1];
__host__ __device__ constexpr index_t& operator[](index_t i) { return data[i]; }
__host__ __device__ constexpr const index_t& operator[](index_t i) const { return data[i]; }
__host__ __device__ constexpr index_t& operator[](index_t i) [[clang::lifetimebound]]
{
return data[i];
}
__host__ __device__ constexpr const index_t& operator[](index_t i) const
[[clang::lifetimebound]]
{
return data[i];
}
};
/**

8
include/ck/utility/statically_indexed_array.hpp
View File

@@ -69,7 +69,7 @@ struct StaticallyIndexedArray_v2
// read access
template <index_t I>
__host__ __device__ constexpr const auto& At(Number<I>) const
__host__ __device__ constexpr const auto& At(Number<I>) const [[clang::lifetimebound]]
{
static_assert(I < N, "wrong! out of range");
@@ -78,7 +78,7 @@ struct StaticallyIndexedArray_v2
// write access
template <index_t I>
__host__ __device__ constexpr auto& At(Number<I>)
__host__ __device__ constexpr auto& At(Number<I>) [[clang::lifetimebound]]
{
static_assert(I < N, "wrong! out of range");
@@ -87,14 +87,14 @@ struct StaticallyIndexedArray_v2
// read access
template <index_t I>
__host__ __device__ constexpr const auto& operator[](Number<I> i) const
__host__ __device__ constexpr const auto& operator[](Number<I> i) const [[clang::lifetimebound]]
{
return At(i);
}
// write access
template <index_t I>
__host__ __device__ constexpr auto& operator()(Number<I> i)
__host__ __device__ constexpr auto& operator()(Number<I> i) [[clang::lifetimebound]]
{
return At(i);
}

22
include/ck_tile/core/container/array.hpp
View File

@@ -103,23 +103,23 @@ struct array
// clang-format off
CK_TILE_HOST_DEVICE constexpr auto& get() { return data; }
CK_TILE_HOST_DEVICE constexpr const auto& get() const { return data; }
CK_TILE_HOST_DEVICE constexpr auto& get(index_t i) { return data[i]; }
CK_TILE_HOST_DEVICE constexpr const auto& get(index_t i) const { return data[i]; }
CK_TILE_HOST_DEVICE constexpr auto& get(index_t i) [[clang::lifetimebound]] { return data[i]; }
CK_TILE_HOST_DEVICE constexpr const auto& get(index_t i) const [[clang::lifetimebound]] { return data[i]; }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& get() { return data[I]; }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& get() const { return data[I]; }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& get(number<I>) { return data[I]; }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& get(number<I>)[[clang::lifetimebound]] { return data[I]; }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& get(number<I>) const { return data[I]; }
CK_TILE_HOST_DEVICE constexpr auto& at(index_t i) { return get(i); }
CK_TILE_HOST_DEVICE constexpr const auto& at(index_t i) const { return get(i); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at() { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at() const { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at(number<I>) { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at(number<I>) const { return get(I); }
CK_TILE_HOST_DEVICE constexpr const auto& at(index_t i) const [[clang::lifetimebound]] { return get(i); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at() [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at() const [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at(number<I>) [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at(number<I>) const [[clang::lifetimebound]] { return get(I); }
CK_TILE_HOST_DEVICE constexpr const value_type& operator[](index_t i) const { return get(i); }
CK_TILE_HOST_DEVICE constexpr value_type& operator[](index_t i) { return get(i); }
CK_TILE_HOST_DEVICE constexpr value_type& operator()(index_t i) { return get(i); } // TODO: compatible
CK_TILE_HOST_DEVICE constexpr const value_type& operator[](index_t i) const [[clang::lifetimebound]] { return get(i); }
CK_TILE_HOST_DEVICE constexpr value_type& operator[](index_t i) [[clang::lifetimebound]] { return get(i); }
CK_TILE_HOST_DEVICE constexpr value_type& operator()(index_t i) [[clang::lifetimebound]] { return get(i); } // TODO: compatible
#if 0
template <typename ArrayLike>
CK_TILE_HOST_DEVICE constexpr auto operator=(const ArrayLike& arr)

10
include/ck_tile/core/container/static_array.hpp
View File

@@ -27,8 +27,14 @@ struct static_array
T elems[N > 0 ? N : 1];
// Basic constexpr accessors
CK_TILE_HOST_DEVICE constexpr const T& operator[](index_t i) const { return elems[i]; }
CK_TILE_HOST_DEVICE constexpr T& operator[](index_t i) { return elems[i]; }
CK_TILE_HOST_DEVICE constexpr const T& operator[](index_t i) const [[clang::lifetimebound]]
{
return elems[i];
}
CK_TILE_HOST_DEVICE constexpr T& operator[](index_t i) [[clang::lifetimebound]]
{
return elems[i];
}
CK_TILE_HOST_DEVICE static constexpr index_t size() { return N; }
};

22
include/ck_tile/core/container/thread_buffer.hpp
View File

@@ -54,17 +54,17 @@ struct thread_buffer {
CK_TILE_HOST_DEVICE static constexpr auto size() { return N; }
CK_TILE_HOST_DEVICE auto & get() {return data; }
CK_TILE_HOST_DEVICE const auto & get() const {return data; }
CK_TILE_HOST_DEVICE auto & get(index_t i) {return data[i]; }
CK_TILE_HOST_DEVICE const auto & get(index_t i) const {return data[i]; }
CK_TILE_HOST_DEVICE constexpr const auto& operator[](index_t i) const { return get(i); }
CK_TILE_HOST_DEVICE constexpr auto& operator[](index_t i) { return get(i); }
CK_TILE_HOST_DEVICE constexpr auto& operator()(index_t i) { return get(i); } // TODO: compatible
CK_TILE_HOST_DEVICE constexpr auto& at(index_t i) { return get(i); }
CK_TILE_HOST_DEVICE constexpr const auto& at(index_t i) const { return get(i); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at() { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at() const { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at(number<I>) { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at(number<I>) const { return get(I); }
CK_TILE_HOST_DEVICE auto & get(index_t i) [[clang::lifetimebound]] {return data[i]; }
CK_TILE_HOST_DEVICE const auto & get(index_t i) const [[clang::lifetimebound]] {return data[i]; }
CK_TILE_HOST_DEVICE constexpr const auto& operator[](index_t i) const [[clang::lifetimebound]] {return get(i); }
CK_TILE_HOST_DEVICE constexpr auto& operator[](index_t i) [[clang::lifetimebound]] { return get(i); }
CK_TILE_HOST_DEVICE constexpr auto& operator()(index_t i) [[clang::lifetimebound]] { return get(i); } // TODO: compatible
CK_TILE_HOST_DEVICE constexpr auto& at(index_t i) [[clang::lifetimebound]] { return get(i); }
CK_TILE_HOST_DEVICE constexpr const auto& at(index_t i) const [[clang::lifetimebound]] { return get(i); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at() [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at() const [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr auto& at(number<I>) [[clang::lifetimebound]] { return get(I); }
template <index_t I> CK_TILE_HOST_DEVICE constexpr const auto& at(number<I>) const [[clang::lifetimebound]] { return get(I); }
template <typename X_,
typename std::enable_if<has_same_scalar_type<value_type, X_>::value, bool>::type = false>

7
include/ck_tile/host/kernel_launch.hpp
View File

@@ -87,8 +87,11 @@ template <int MinBlockPerCu = CK_TILE_MIN_BLOCK_PER_CU,
typename Attr = void,
typename KernelImpl,
typename... Args>
CK_TILE_HOST auto
make_kernel(KernelImpl /*f*/, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)
CK_TILE_HOST auto make_kernel(KernelImpl /*f*/,
dim3 grid_dim,
dim3 block_dim,
std::size_t lds_byte,
[[clang::lifetimebound]] Args... args)
{
const auto kernel = []() {
if constexpr(std::is_void_v<Attr>)