Skip to content

[Compiler] Addressing new compiler warnings #3640

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
jplehr wants to merge 7 commits into
base: develop
Choose a base branch
from
+220 −43
Open

[Compiler] Addressing new compiler warnings #3640

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
a34b09c
[Compiler] Addressing new compiler warnings
jplehr Jan 23, 2026
dc34993
Update some more instances
jplehr Jan 23, 2026
d1dece6
Adds file-level ignores via clang diagnostic pragma
jplehr Jan 29, 2026
1a9b028
This adds the remaining instances
jplehr Jan 29, 2026
5a0b74d
Merge branch 'develop' into janplehr/clang-lifetime-warning-annotation
illsilin Jan 29, 2026
673ff17
fix clang format
illsilin Jan 29, 2026
d677cde
Adding couple more instances from gfx1200 build
jplehr Jan 29, 2026
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: 2 additions & 0 deletions cmake/gtest.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,6 +68,8 @@ set(GTEST_CXX_FLAGS
-Wno-deprecated
-Wno-unsafe-buffer-usage
-Wno-float-equal
-Wno-lifetime-safety-intra-tu-suggestions
-Wno-lifetime-safety-cross-tu-suggestions
)

if(WIN32)
Expand Down
2 changes: 1 addition & 1 deletion example/ck_tile/01_fmha/bias.hpp
View file
Open in desktop
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
View file
Open in desktop
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
4 changes: 4 additions & 0 deletions example/ck_tile/01_fmha/quant.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,9 @@
#include "ck_tile/core.hpp"
#include "ck_tile/ops/fmha.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

// keep sync with BlockAttentionQuantScaleEnum
enum class quant_scale_enum
{
Expand Down Expand Up @@ -58,3 +61,4 @@ struct quant_scale_info
return os;
}
};
#pragma clang diagnostic pop
5 changes: 3 additions & 2 deletions include/ck/host_utility/io.hpp
View file
Open in desktop
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
View file
Open in desktop
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);
12 changes: 9 additions & 3 deletions include/ck/library/utility/host_tensor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,10 +23,14 @@

#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
#pragma clang diagnostic ignored "-Wlifetime-safety-cross-tu-suggestions"

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 +584,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 Expand Up @@ -1171,3 +1176,4 @@ struct Tensor
};

} // namespace ck
#pragma clang diagnostic pop
3 changes: 3 additions & 0 deletions include/ck/tensor/static_tensor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,8 @@
#ifndef CK_STATIC_TENSOR_HPP
#define CK_STATIC_TENSOR_HPP

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
namespace ck {

// StaticTensor for Scalar
Expand Down Expand Up @@ -270,4 +272,5 @@ __host__ __device__ constexpr auto make_static_tensor(TensorDesc, X invalid_elem
}

} // namespace ck
#pragma clang diagnostic pop
#endif
39 changes: 32 additions & 7 deletions include/ck/tensor_description/multi_index_transform.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,9 @@
#include "ck/utility/common_header.hpp"
#include "ck/utility/multi_index.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

namespace ck {

template <typename LowLength>
Expand All @@ -29,7 +32,10 @@ struct PassThrough

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return 1; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ static constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -305,7 +311,10 @@ struct RightPad

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return 1; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ static constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -403,7 +412,10 @@ struct Embed

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return NDimUp; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -1074,7 +1086,10 @@ struct Merge_v2_magic_division

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return 1; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -1366,7 +1381,10 @@ struct Merge_v3_division_mod

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return 1; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -1480,7 +1498,10 @@ struct UnMerge

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return NDimUp; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename LowIdx, typename UpIdx>
__host__ __device__ constexpr void CalculateLowerIndex(LowIdx& idx_low,
Expand Down Expand Up @@ -1640,7 +1661,10 @@ struct ConvBwdDataImplicitGemmOutTransform

__host__ __device__ static constexpr index_t GetNumOfUpperDimension() { return 3; }

__host__ __device__ constexpr const auto& GetUpperLengths() const { return up_lengths_; }
__host__ __device__ constexpr const auto& GetUpperLengths() const [[clang::lifetimebound]]
{
return up_lengths_;
}

template <typename UpIdx>
__host__ __device__ constexpr auto CalculateLowerIndexN(const UpIdx& idx_up) const
Expand Down Expand Up @@ -2236,3 +2260,4 @@ struct Xor
}
};
} // namespace ck
#pragma clang diagnostic pop
5 changes: 4 additions & 1 deletion include/ck/tensor_description/tensor_adaptor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,10 @@ struct TensorAdaptor
{
__host__ __device__ static constexpr index_t GetNumOfTransform() { return Transforms::Size(); }

__host__ __device__ constexpr const auto& GetTransforms() const { return transforms_; }
__host__ __device__ constexpr const auto& GetTransforms() const [[clang::lifetimebound]]
{
return transforms_;
}

__host__ __device__ static constexpr auto GetLowerDimensionHiddenIdss()
{
Expand Down
17 changes: 13 additions & 4 deletions include/ck/tensor_description/tensor_descriptor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,8 @@
#include "ck/utility/sequence_helper.hpp"
#include "ck/tensor_description/multi_index_transform.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
namespace ck {

template <index_t NDimHidden, typename VisibleDimensionIds>
Expand Down Expand Up @@ -179,7 +181,10 @@ struct TensorDescriptor
}

// TODO make these private
__host__ __device__ constexpr const auto& GetTransforms() const { return transforms_; }
__host__ __device__ constexpr const auto& GetTransforms() const [[clang::lifetimebound]]
{
return transforms_;
}

__host__ __device__ static constexpr auto GetLowerDimensionIdss()
{
Expand Down Expand Up @@ -253,9 +258,12 @@ struct TensorCoordinate
__host__ __device__ constexpr index_t GetOffset() const { return idx_hidden_[Number<0>{}]; }

// TODO make these private
__host__ __device__ constexpr const auto& GetHiddenIndex() const { return idx_hidden_; }
__host__ __device__ constexpr const auto& GetHiddenIndex() const [[clang::lifetimebound]]
{
return idx_hidden_;
}

__host__ __device__ auto& GetHiddenIndex() { return idx_hidden_; }
__host__ __device__ auto& GetHiddenIndex() [[clang::lifetimebound]] { return idx_hidden_; }

__host__ __device__ constexpr auto GetVisibleIndex() const
{
Expand Down Expand Up @@ -284,7 +292,7 @@ struct TensorCoordinateStep
__host__ __device__ constexpr const auto& GetIndexDiff() const { return GetVisibleIndexDiff(); }

// TODO make these private
__host__ __device__ constexpr const auto& GetVisibleIndexDiff() const
__host__ __device__ constexpr const auto& GetVisibleIndexDiff() const [[clang::lifetimebound]]
{
return idx_diff_visible_;
}
Expand Down Expand Up @@ -613,3 +621,4 @@ using TensorCoordinateStep_t = decltype(make_tensor_coordinate_step(
TensorDesc{}, MultiIndex<remove_cvref_t<TensorDesc>::GetNumOfDimension()>{}));

} // namespace ck
#pragma clang diagnostic pop
3 changes: 3 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_wmmaops_base.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,8 @@
#include "ck/tensor_operation/gpu/warp/wmma_gemm.hpp"
#include "ck/tensor_description/tensor_adaptor.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
namespace ck {

template <index_t BlockSize,
Expand Down Expand Up @@ -485,3 +487,4 @@ struct BlockwiseGemmWmmaops_pipeline_base
};

} // namespace ck
#pragma clang diagnostic pop
3 changes: 3 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,8 @@
// Prefetech 2 stage
// Local prefetch 1 stage

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
namespace ck {

template <index_t BlockSize,
Expand Down Expand Up @@ -992,3 +994,4 @@ struct BlockwiseGemmXdlops_pipeline_v4
};

} // namespace ck
#pragma clang diagnostic pop
4 changes: 4 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_pipeline_xdlops_base.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,9 @@
#include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp"
#include "ck/tensor_description/tensor_adaptor.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

namespace ck {

template <index_t BlockSize,
Expand Down Expand Up @@ -404,3 +407,4 @@ struct BlockwiseGemmXdlops_pipeline_base
};

} // namespace ck
#pragma clang diagnostic pop
4 changes: 4 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_wmma.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,9 @@

#define CK_MNK_LOOP

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

namespace ck {

#ifdef __gfx12__
Expand Down Expand Up @@ -1028,3 +1031,4 @@ struct BlockwiseGemmWMMA
#endif

} // namespace ck
#pragma clang diagnostic pop
3 changes: 3 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,8 @@
#include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp"
#include "ck/tensor_description/tensor_adaptor.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"
namespace ck {

template <index_t MNXdlPerWave, index_t MNWaves, index_t MNPerXdl, typename TileDesc_K0_MN_K1>
Expand Down Expand Up @@ -1031,3 +1033,4 @@ struct BlockwiseGemmXdlops_v2
};

} // namespace ck
#pragma clang diagnostic pop
4 changes: 4 additions & 0 deletions include/ck/tensor_operation/gpu/block/blockwise_gemm_xdlops_skip_b_lds.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,9 @@
#include "ck/tensor_operation/gpu/warp/xdlops_gemm.hpp"
#include "ck/tensor_description/tensor_adaptor.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

namespace ck {

template <index_t BlockSize,
Expand Down Expand Up @@ -317,3 +320,4 @@ struct BlockwiseGemmXdlops_k0mk1_k0nk1_m0n0m1n1m2m3m4n2_v1r1
};

} // namespace ck
#pragma clang diagnostic pop
2 changes: 1 addition & 1 deletion include/ck/tensor_operation/gpu/device/tensor_layout.hpp
View file
Open in desktop
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
4 changes: 4 additions & 0 deletions include/ck/tensor_operation/gpu/grid/gridwise_gemm_xdlops_bwd_weight.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,9 @@
#include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_common.hpp"
#include "ck/tensor_operation/gpu/device/device_base.hpp"

#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wlifetime-safety-intra-tu-suggestions"

namespace ck {

// Implementation of "Merge" transformation primitive that uses division and mod. It is supposed to
Expand Down Expand Up @@ -1132,3 +1135,4 @@ struct GridwiseGemm_bk0mk1_bk0nk1_mn_xdlops_bwd_weight
}; // namespace ck

} // namespace ck
#pragma clang diagnostic pop
3 changes: 2 additions & 1 deletion include/ck/utility/amd_wave_read_first_lane.hpp
View file
Open in desktop
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
Loading