NVIDIA · cliffburdick · Oct 31, 2024 · Oct 21, 2024 · Oct 21, 2024 · Oct 30, 2024
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -115,7 +115,6 @@ target_include_directories(matx INTERFACE "$<BUILD_INTERFACE:${CMAKE_CURRENT_SOU
 target_include_directories(matx INTERFACE "$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include/matx/kernels>"
 "$<INSTALL_INTERFACE:include/matx/kernels>")
 target_compile_features(matx INTERFACE cxx_std_17 $<BUILD_INTERFACE:cuda_std_17>)
-target_compile_options(matx INTERFACE $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>)
 
 # 11.2 and above required for async allocation
 if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 11.5)

diff --git a/cmake/versions.json b/cmake/versions.json
@@ -1,10 +1,9 @@
 {
   "packages": {
     "CCCL": {
-      "version": "2.4.0",
+      "version": "2.7.0-rc2",
       "git_shallow": true,
       "git_url": "https://github.com/NVIDIA/cccl.git",
-      "git_tag": "1c009d2"
     }
   }
 }
diff --git a/docs_input/build.rst b/docs_input/build.rst
@@ -28,7 +28,7 @@ Required Third-party Dependencies
 ---------------------------------
 
 - `CPM <https://github.com/cpm-cmake/CPM.cmake>`_ (* Included in the project source and does not require a separate download)
-- `CCCL <https://github.com/NVIDIA/cccl>`_ 2.4.0+
+- `CCCL <https://github.com/NVIDIA/cccl>`_ 2.7.0+
 
 
 Optional Third-party Dependencies

diff --git a/include/matx/core/allocator.h b/include/matx/core/allocator.h
@@ -43,8 +43,7 @@
 
 #include "matx/core/error.h"
 #include "matx/core/nvtx.h"
-#include <cuda/std/__algorithm>
-#include <cuda/std/__algorithm>
+#include <cuda/std/functional>
 
 #pragma once
 

diff --git a/include/matx/operators/clone.h b/include/matx/operators/clone.h
@@ -107,7 +107,7 @@ IGNORE_WARNING_POP_GCC
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/collapse.h b/include/matx/operators/collapse.h
@@ -98,7 +98,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }   
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()
@@ -235,7 +235,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }   
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/fftshift.h b/include/matx/operators/fftshift.h
@@ -66,7 +66,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/interleaved.h b/include/matx/operators/interleaved.h
@@ -75,7 +75,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/kronecker.h b/include/matx/operators/kronecker.h
@@ -81,7 +81,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         template <typename ShapeType, typename Executor>

diff --git a/include/matx/operators/permute.h b/include/matx/operators/permute.h
@@ -115,7 +115,7 @@ IGNORE_WARNING_POP_GCC
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         constexpr __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ index_t Size(int32_t dim) const

diff --git a/include/matx/operators/planar.h b/include/matx/operators/planar.h
@@ -72,7 +72,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/r2c.h b/include/matx/operators/r2c.h
@@ -75,7 +75,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }      
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/remap.h b/include/matx/operators/remap.h
@@ -88,7 +88,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/repmat.h b/include/matx/operators/repmat.h
@@ -110,7 +110,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         template <typename ShapeType, typename Executor>

diff --git a/include/matx/operators/reshape.h b/include/matx/operators/reshape.h
@@ -108,7 +108,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         constexpr __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ index_t Size(int32_t dim) const

diff --git a/include/matx/operators/reverse.h b/include/matx/operators/reverse.h
@@ -80,7 +80,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/scalar_ops.h b/include/matx/operators/scalar_ops.h
@@ -235,7 +235,20 @@ template <typename T> static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto
     return matx::sin(v1);
   }
   else {
+    // This should go away once CCCL adds <cmath> support
+#ifdef __CUDACC__
+    if constexpr (std::is_same_v<T, float>) {
+      return ::sinf(v1);
+    }
+    else if constexpr (std::is_same_v<T, double>) {
+      return ::sin(v1);
+    }
+    else {
+      matx::sin(v1);
+    }
+#else
     return cuda::std::sin(v1);
+#endif 
   }
 }
 template <typename T> struct SinF {
@@ -250,7 +263,20 @@ template <typename T> static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto
     return matx::cos(v1);
   }
   else {
+    // This should go away once CCCL adds <cmath> support
+#ifdef __CUDACC__
+    if constexpr (std::is_same_v<T, float>) {
+      return ::cosf(v1);
+    }
+    else if constexpr (std::is_same_v<T, double>) {
+      return ::cos(v1);
+    }
+    else {
+      matx::cos(v1);
+    }
+#else
     return cuda::std::cos(v1);
+#endif  
   }
 }
 template <typename T> struct CosF {
@@ -691,7 +717,8 @@ static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto _internal_isinf(T v1)
     return cuda::std::isinf(static_cast<typename castType::value_type>(v1.real())) || cuda::std::isinf(static_cast<typename castType::value_type>(v1.imag()));
   } else {
     return cuda::std::isinf(static_cast<castType>(v1));
-  } 
+  }  
+
 }
 template <typename T>
 struct IsInf {

diff --git a/include/matx/operators/select.h b/include/matx/operators/select.h
@@ -69,7 +69,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(index_t i)
         {
-          return std::as_const(*this).template operator()(i);
+          return cuda::std::as_const(*this).template operator()(i);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/self.h b/include/matx/operators/self.h
@@ -68,7 +68,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/shift.h b/include/matx/operators/shift.h
@@ -91,7 +91,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         template <typename ShapeType, typename Executor>

diff --git a/include/matx/operators/slice.h b/include/matx/operators/slice.h
@@ -140,7 +140,7 @@ namespace matx
         template <typename... Is>
         __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
         {
-          return std::as_const(*this).template operator()(indices...);
+          return cuda::std::as_const(*this).template operator()(indices...);
         }
 
         static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank()

diff --git a/include/matx/operators/stack.h b/include/matx/operators/stack.h
@@ -147,7 +147,7 @@ namespace matx
       template <typename... Is>
       __MATX_INLINE__ __MATX_DEVICE__ __MATX_HOST__ decltype(auto) operator()(Is... indices)
       {
-        return std::as_const(*this).template operator()(indices...);
+        return cuda::std::as_const(*this).template operator()(indices...);
       }
 
       static __MATX_INLINE__ constexpr __MATX_HOST__ __MATX_DEVICE__ int32_t Rank() noexcept

diff --git a/include/matx/operators/toeplitz.h b/include/matx/operators/toeplitz.h
@@ -35,6 +35,7 @@
 
 #include "matx/core/type_utils.h"
 #include "matx/operators/base_operator.h"
+#include <cuda/std/__algorithm/transform.h>
 
 namespace matx
 {

diff --git a/include/matx/transforms/reduce.h b/include/matx/transforms/reduce.h
@@ -2069,7 +2069,7 @@ void __MATX_INLINE__ argmax_impl(OutType dest, TensorIndexType &idest, const InT
 #ifdef __CUDACC__
   MATX_NVTX_START("argmax_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API)
 
-  const auto initial_value = cuda::std::make_tuple(static_cast<matx::index_t>(-1), std::numeric_limits<typename InType::value_type>::lowest());
+  const auto initial_value = thrust::make_tuple(static_cast<matx::index_t>(-1), std::numeric_limits<typename InType::value_type>::lowest());
   using reduce_param_type = typename detail::ReduceParams_t<typename detail::CustomArgMaxCmp, decltype(initial_value)>;
   auto reduce_params = reduce_param_type{detail::CustomArgMaxCmp{}, initial_value};
 
@@ -2219,7 +2219,7 @@ void __MATX_INLINE__ argmin_impl(OutType dest, TensorIndexType &idest, const InT
 #ifdef __CUDACC__
   MATX_NVTX_START("argmin_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API)
 
-  const auto initial_value = cuda::std::make_tuple(static_cast<matx::index_t>(-1), std::numeric_limits<typename InType::value_type>::max());
+  const auto initial_value = thrust::make_tuple(static_cast<matx::index_t>(-1), std::numeric_limits<typename InType::value_type>::max());
   using reduce_param_type = typename detail::ReduceParams_t<typename detail::CustomArgMinCmp, decltype(initial_value)>;
   auto reduce_params = reduce_param_type{detail::CustomArgMinCmp{}, initial_value};
 
@@ -2307,7 +2307,7 @@ void __MATX_INLINE__ argminmax_impl(OutType destmin, TensorIndexType &idestmin,
 #ifdef __CUDACC__
   MATX_NVTX_START("argminmax_impl(" + get_type_str(in) + ")", matx::MATX_NVTX_LOG_API)
 
-  const auto initial_value = cuda::std::make_tuple(
+  const auto initial_value = thrust::make_tuple(
     static_cast<matx::index_t>(-1),
     std::numeric_limits<typename InType::value_type>::max(),
     static_cast<matx::index_t>(-1),