Fixing warnings issues by clang-19, both host and device (#825)

Author: cliffburdick

CMakeLists.txt

@@ -134,8 +134,10 @@ rapids_cpm_cccl(
 
 target_link_libraries(matx INTERFACE CCCL::CCCL)
 
-# Set flags for compiling tests faster
-set(MATX_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} --threads 0 -ftemplate-backtrace-limit=0)
+# Set flags for compiling tests faster (only for nvcc)
+if (NOT CMAKE_CUDA_COMPILER_ID STREQUAL "Clang")
+    set(MATX_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} --threads 0 -ftemplate-backtrace-limit=0)
+endif()
 
 # Hack because CMake doesn't have short circult evaluation
 if (NOT CMAKE_BUILD_TYPE OR "${CMAKE_BUILD_TYPE}" STREQUAL "Debug")
@@ -165,9 +167,13 @@ if (NOT ${IS_NVCPP} GREATER -1)
     endif()
 endif()
 
-
-
-set(WARN_FLAGS ${WARN_FLAGS} $<$<COMPILE_LANGUAGE:CUDA>:-Werror all-warnings>)
+if (CMAKE_CUDA_COMPILER_ID STREQUAL "Clang")
+message((STATUS "Using Clang compiler"))
+    # Workaround for clang bug: https://github.com/llvm/llvm-project/issues/58491
+    set(WARN_FLAGS ${WARN_FLAGS} $<$<COMPILE_LANGUAGE:CUDA>:-Wno-unused-command-line-argument>)
+else()
+    set(WARN_FLAGS ${WARN_FLAGS} $<$<COMPILE_LANGUAGE:CUDA>:-Werror all-warnings>)
+endif()
 set(WARN_FLAGS ${WARN_FLAGS} $<$<COMPILE_LANGUAGE:CXX>:-Werror>)
 
 # CUTLASS slows down compile times when used, so leave it as optional for now

examples/black_scholes.cu

@@ -61,7 +61,7 @@ private:
 
 public:
   BlackScholes(O out, I1 K, I1 V, I1 S, I1 r, I1 T)
-      : out_(out), K_(K), V_(V), S_(S), r_(r), T_(T)  {}
+      : out_(out), V_(V), S_(S), K_(K), r_(r), T_(T)  {}
 
   __device__ inline void operator()(index_t idx)
   {

examples/convolution.cu

@@ -39,7 +39,6 @@ using namespace matx;
 int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv)
 {
   MATX_ENTER_HANDLER();
-  typedef cuda::std::complex<float> complex;
 
   uint32_t iterations = 10;
   constexpr index_t numSamples = 1638400;

examples/mvdr_beamformer.h

@@ -164,9 +164,9 @@ class MVDRBeamformer {
   auto GetCovMatInvView() { return invCovMatView; }
 
 private:
-  index_t num_beams_;
-  index_t num_el_;
-  index_t data_len_;
+  [[maybe_unused]] index_t num_beams_;
+  [[maybe_unused]] index_t num_el_;
+  [[maybe_unused]] index_t data_len_;
   index_t snap_len_;
   cuda::std::complex<float> load_coeff_ = {0.1f, 0.f};
 

examples/recursive_filter.cu

@@ -40,8 +40,6 @@ using namespace matx;
 int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv)
 {
   MATX_ENTER_HANDLER();
-  using complex = cuda::std::complex<float>;
-
   cudaDeviceProp prop;
   cudaGetDeviceProperties(&prop, 0);
 
@@ -70,7 +68,6 @@ int main([[maybe_unused]] int argc, [[maybe_unused]] char **argv)
   cudaEventCreate(&start);
   cudaEventCreate(&stop);
 
-  using OutType = float;
   using InType = float;
   using FilterType = float;
 

include/matx/core/half.h

@@ -417,7 +417,7 @@ __MATX_HOST__ __MATX_DEVICE__ __MATX_INLINE__ bool operator==(const T &lhs,
                                                     const matxHalf<T> &rhs)
 {
   matxHalf<T> tmp{lhs};
-  return lhs == tmp;
+  return rhs == tmp;
 }
 
 /**
@@ -464,7 +464,7 @@ __MATX_HOST__ __MATX_DEVICE__ __MATX_INLINE__ bool operator!=(const T &lhs,
                                                     const matxHalf<T> &rhs)
 {
   matxHalf<T> tmp{lhs};
-  return !(lhs == tmp);
+  return !(rhs == tmp);
 }
 
 /**

include/matx/core/half_complex.h

@@ -515,7 +515,7 @@ __MATX_HOST__ __MATX_DEVICE__ __MATX_INLINE__ bool
 operator==(const T &lhs, const matxHalfComplex<T> &rhs)
 {
   matxHalfComplex<T> tmp{lhs};
-  return lhs == tmp;
+  return rhs == tmp;
 }
 
 /**
@@ -562,7 +562,7 @@ __MATX_HOST__ __MATX_DEVICE__ __MATX_INLINE__ bool
 operator!=(const T &lhs, const matxHalfComplex<T> &rhs)
 {
   matxHalfComplex<T> tmp{lhs};
-  return !(lhs == tmp);
+  return !(rhs == tmp);
 }
 
 
@@ -853,7 +853,7 @@ pow(const T &x, const matxHalfComplex<T> &y)
 {
   cuda::std::complex<float> tmp{static_cast<float>(y.real()),
                                 static_cast<float>(y.imag())};
-  tmp = cuda::std::pow(y, pow);
+  tmp = cuda::std::pow(x, pow);
   return {static_cast<T>(tmp.real()), static_cast<T>(tmp.imag())};
 }
 

include/matx/core/operator_utils.h

@@ -132,8 +132,6 @@ namespace matx {
 
     template <typename Op, typename ValidFunc>
     __MATX_INLINE__ auto GetSupportedTensor(const Op &in, const ValidFunc &fn, matxMemorySpace_t space, cudaStream_t stream = 0) {
-      constexpr int RANK = Op::Rank();
-
       if constexpr (is_matx_transform_op<Op>()) {
         // We can assume that if a transform is passed to the input then PreRun has already completed
         // on the transform and we can use the internal pointer

include/matx/core/print.h

@@ -653,7 +653,7 @@ namespace matx {
    */
   template <typename Op, typename... Args,
             std::enable_if_t<(Op::Rank() > 0 && sizeof...(Args) == 0), bool> = true>
-  void fprint(FILE* fp, const Op &op, Args... dims) {
+  void fprint(FILE* fp, const Op &op, [[maybe_unused]] Args... dims) {
     cuda::std::array<int, Op::Rank()> arr = {0};
     auto tp = cuda::std::tuple_cat(arr);
     cuda::std::apply([&](auto &&...args) { fprint(fp, op, args...); }, tp);

include/matx/core/storage.h

@@ -406,7 +406,7 @@ namespace matx
      */
     void SetData(T *const data) noexcept
     {
-      data_.reset(data_, [](auto){});
+      data_.reset(data, [](auto){});
     }
 
     /**
@@ -423,7 +423,7 @@ namespace matx
      * 
      * @param size Size in bytes to allocate
      */
-    __MATX_INLINE__ T* allocate(size_t size)
+    __MATX_INLINE__ T* allocate([[maybe_unused]] size_t size)
     {
       MATX_THROW(matxInvalidParameter, "Cannot call allocate on a smart pointer storage type");
     }