Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion example/ck_tile/01_fmha/bias.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -106,7 +106,7 @@ struct bias_info
return info;
}

friend std::ostream& operator<<(std::ostream& os, const bias_info& bi)
friend std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os, const bias_info& bi)
{
bi.serialize(os);
return os;
Expand Down
2 changes: 1 addition & 1 deletion example/ck_tile/01_fmha/mask.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -191,7 +191,7 @@ struct mask_info
return area;
}

friend std::ostream& operator<<(std::ostream& os, const mask_info& mi)
friend std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os, const mask_info& mi)
{
mi.serialize(os);
return os;
Expand Down
5 changes: 3 additions & 2 deletions include/ck/host_utility/io.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
namespace ck {

template <typename T>
std::ostream& operator<<(std::ostream& os, const std::vector<T>& v)
std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os, const std::vector<T>& v)
{
std::copy(std::begin(v), std::end(v), std::ostream_iterator<T>(os, " "));
return os;
Expand All @@ -27,7 +27,8 @@ std::ostream& operator<<(std::ostream& os, const std::array<T, N>& v)
}

template <typename... Ts>
std::ostream& operator<<(std::ostream& os, const TensorDescriptor<Ts...>& desc)
std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const TensorDescriptor<Ts...>& desc)
{
constexpr index_t nDim = remove_cvref_t<decltype(desc)>::GetNumOfDimension();

Expand Down
3 changes: 2 additions & 1 deletion include/ck/library/utility/convolution_parameter.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -110,4 +110,5 @@ ConvParam parse_conv_param(int num_dim_spatial, int arg_idx, char* const argv[])
} // namespace utils
} // namespace ck

std::ostream& operator<<(std::ostream& os, const ck::utils::conv::ConvParam& p);
std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const ck::utils::conv::ConvParam& p);
7 changes: 4 additions & 3 deletions include/ck/library/utility/host_tensor.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@
namespace ck {

template <typename Range>
std::ostream& LogRange(std::ostream& os, Range&& range, std::string delim)
std::ostream& LogRange([[clang::lifetimebound]] std::ostream& os, Range&& range, std::string delim)
{
bool first = true;
for(auto&& v : range)
Expand Down Expand Up @@ -580,8 +580,9 @@ struct HostTensorDescriptor
return std::inner_product(iss.begin(), iss.end(), mStrides.begin(), std::size_t{0});
}

friend std::ostream& operator<<(std::ostream& os, const HostTensorDescriptor& desc);
friend std::ostream& operator<<(std::ostream& os, ChosenLayout tag);
friend std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const HostTensorDescriptor& desc);
friend std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os, ChosenLayout tag);

private:
std::vector<std::size_t> mLens;
Expand Down
2 changes: 1 addition & 1 deletion include/ck/tensor_operation/gpu/device/tensor_layout.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -455,7 +455,7 @@ struct G_NDHW : public BaseConvolutionLayout
template <
typename Layout,
typename std::enable_if<std::is_base_of<BaseTensorLayout, Layout>::value, bool>::type = false>
std::ostream& operator<<(std::ostream& os, const Layout&)
std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os, const Layout&)
{
os << Layout::name;
return os;
Expand Down
3 changes: 2 additions & 1 deletion include/ck/utility/amd_wave_read_first_lane.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,8 @@ struct get_carrier<3>

// replacement of host std::copy_n()
template <typename InputIterator, typename Size, typename OutputIterator>
__device__ static OutputIterator copy_n(InputIterator from, Size size, OutputIterator to)
__device__ static OutputIterator
copy_n(InputIterator from, Size size, [[clang::lifetimebound]] OutputIterator to)
Copy link

Copilot AI Jan 23, 2026

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The [[clang::lifetimebound]] attribute on the OutputIterator parameter is incorrect. This function returns the iterator by value after advancing it, not a reference bound to the input parameter's lifetime. The lifetimebound attribute should only be applied when the return value's lifetime is directly tied to the parameter (e.g., returning a reference to data owned by the parameter). Remove this attribute as it doesn't apply to this pattern.

Suggested change
copy_n(InputIterator from, Size size, [[clang::lifetimebound]] OutputIterator to)
copy_n(InputIterator from, Size size, OutputIterator to)

Copilot uses AI. Check for mistakes.
{
if(0 < size)
{
Expand Down
3 changes: 2 additions & 1 deletion include/ck/utility/pipeline_enum.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,8 @@ enum struct PipelineVersion
} // namespace ck

#if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC)
inline std::ostream& operator<<(std::ostream& os, const ck::PipelineVersion& p)
inline std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const ck::PipelineVersion& p)
{
switch(p)
{
Expand Down
3 changes: 2 additions & 1 deletion include/ck/utility/scheduler_enum.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -70,7 +70,8 @@ enum struct TailNumber
} // namespace ck

#if !defined(__HIPCC_RTC__) || !defined(CK_CODE_GEN_RTC)
inline std::ostream& operator<<(std::ostream& os, const ck::LoopScheduler& s)
inline std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const ck::LoopScheduler& s)
{
switch(s)
{
Expand Down
2 changes: 1 addition & 1 deletion include/ck/utility/tuple.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -51,7 +51,7 @@ get_tuple_element_data_reference(const TupleElementKeyData<Key, Data>& x)
// for write access of tuple element
template <typename Key, typename Data>
__host__ __device__ constexpr Data&
get_tuple_element_data_reference(TupleElementKeyData<Key, Data>& x)
get_tuple_element_data_reference([[clang::lifetimebound]] TupleElementKeyData<Key, Data>& x)
{
return x.mData;
}
Expand Down
7 changes: 4 additions & 3 deletions include/ck_tile/core/container/tuple.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -98,13 +98,14 @@ CK_TILE_HOST_DEVICE constexpr T getv(const tuple_object<I, T, true>&)
}

template <index_t I, class T>
CK_TILE_HOST_DEVICE constexpr const T& getv(const tuple_object<I, T, false>& x)
CK_TILE_HOST_DEVICE constexpr const T&
getv([[clang::lifetimebound]] const tuple_object<I, T, false>& x)
{
return x.element;
}

template <index_t I, class T>
CK_TILE_HOST_DEVICE constexpr T& getv(tuple_object<I, T, false>& x)
CK_TILE_HOST_DEVICE constexpr T& getv([[clang::lifetimebound]] tuple_object<I, T, false>& x)
{
return x.element;
}
Expand Down Expand Up @@ -292,7 +293,7 @@ struct tuple : impl::tuple_base<make_index_sequence<sizeof...(T)>, T...>
//template <typename Tx> CK_TILE_HOST_DEVICE constexpr decltype(auto) get_as(index_t i) const { TP_COM_(); return reinterpret_cast<const tuple_array<Tx, size()>&>(*this).at(i); }
template <typename Tx, index_t I> CK_TILE_HOST_DEVICE constexpr decltype(auto) get_as(number<I>) { TP_COM_(); return reinterpret_cast<tuple_array<Tx, size()>&>(*this).at(number<I>{}); }
template <typename Tx, index_t I> CK_TILE_HOST_DEVICE constexpr decltype(auto) get_as(number<I>) const { TP_COM_(); return reinterpret_cast<const tuple_array<Tx, size()>&>(*this).at(number<I>{}); }

// template <typename Tx> CK_TILE_HOST_DEVICE constexpr void set_as(index_t i, const Tx & x) { TP_COM_(); reinterpret_cast<tuple_array<Tx, size()>&>(*this).at(i) = x; }
template <typename Tx, index_t I> CK_TILE_HOST_DEVICE constexpr void set_as(number<I>, const Tx & x) { TP_COM_(); reinterpret_cast<tuple_array<Tx, size()>&>(*this).at(number<I>{}) = x; }

Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -41,7 +41,8 @@ enum struct TailNumber

} // namespace ck_tile

inline std::ostream& operator<<(std::ostream& os, const ck_tile::GemmPipelineScheduler& s)
inline std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const ck_tile::GemmPipelineScheduler& s)
{
switch(s)
{
Expand All @@ -53,7 +54,8 @@ inline std::ostream& operator<<(std::ostream& os, const ck_tile::GemmPipelineSch
return os;
}

inline std::ostream& operator<<(std::ostream& os, const ck_tile::TailNumber& s)
inline std::ostream& operator<<([[clang::lifetimebound]] std::ostream& os,
const ck_tile::TailNumber& s)
{
switch(s)
{
Expand Down