diff --git a/CMakeLists.txt b/CMakeLists.txt index bc1bb2ee..602eca46 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -115,7 +115,6 @@ target_include_directories(matx INTERFACE "$" "$") target_compile_features(matx INTERFACE cxx_std_17 $) -target_compile_options(matx INTERFACE $<$:--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 index f3b58f5f..8bd67199 100644 --- 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" } } } \ No newline at end of file diff --git a/docs_input/build.rst b/docs_input/build.rst index 2e6fbbe2..aad03193 100644 --- a/docs_input/build.rst +++ b/docs_input/build.rst @@ -28,7 +28,7 @@ Required Third-party Dependencies --------------------------------- - `CPM `_ (* Included in the project source and does not require a separate download) -- `CCCL `_ 2.4.0+ +- `CCCL `_ 2.7.0+ Optional Third-party Dependencies diff --git a/include/matx/core/allocator.h b/include/matx/core/allocator.h index 7a2eb484..b30e1c85 100644 --- 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 -#include +#include #pragma once diff --git a/include/matx/operators/clone.h b/include/matx/operators/clone.h index dfe11a0c..47b77cdf 100644 --- a/include/matx/operators/clone.h +++ b/include/matx/operators/clone.h @@ -107,7 +107,7 @@ IGNORE_WARNING_POP_GCC template __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 index 88e4c03f..537424aa 100644 --- a/include/matx/operators/collapse.h +++ b/include/matx/operators/collapse.h @@ -98,7 +98,7 @@ namespace matx template __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 __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 index d2bbe6f4..e089f740 100644 --- a/include/matx/operators/fftshift.h +++ b/include/matx/operators/fftshift.h @@ -66,7 +66,7 @@ namespace matx template __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 index a2c1f13d..921a3925 100644 --- a/include/matx/operators/interleaved.h +++ b/include/matx/operators/interleaved.h @@ -75,7 +75,7 @@ namespace matx template __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 index 6b734d60..bb03c44d 100644 --- a/include/matx/operators/kronecker.h +++ b/include/matx/operators/kronecker.h @@ -81,7 +81,7 @@ namespace matx template __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 diff --git a/include/matx/operators/permute.h b/include/matx/operators/permute.h index a8acaccc..bf366e24 100644 --- a/include/matx/operators/permute.h +++ b/include/matx/operators/permute.h @@ -115,7 +115,7 @@ IGNORE_WARNING_POP_GCC template __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 index e80fed24..4bb16cd4 100644 --- a/include/matx/operators/planar.h +++ b/include/matx/operators/planar.h @@ -72,7 +72,7 @@ namespace matx template __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 index d013748d..59d268b6 100644 --- a/include/matx/operators/r2c.h +++ b/include/matx/operators/r2c.h @@ -75,7 +75,7 @@ namespace matx template __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 index bac5b152..1419e9e4 100644 --- a/include/matx/operators/remap.h +++ b/include/matx/operators/remap.h @@ -88,7 +88,7 @@ namespace matx template __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 index a45053e8..87366a93 100644 --- a/include/matx/operators/repmat.h +++ b/include/matx/operators/repmat.h @@ -110,7 +110,7 @@ namespace matx template __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 diff --git a/include/matx/operators/reshape.h b/include/matx/operators/reshape.h index 1b8f2628..8df543f2 100644 --- a/include/matx/operators/reshape.h +++ b/include/matx/operators/reshape.h @@ -108,7 +108,7 @@ namespace matx template __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 index 83b7c70d..0964de47 100644 --- a/include/matx/operators/reverse.h +++ b/include/matx/operators/reverse.h @@ -80,7 +80,7 @@ namespace matx template __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 index ac743094..44113853 100644 --- a/include/matx/operators/scalar_ops.h +++ b/include/matx/operators/scalar_ops.h @@ -235,7 +235,20 @@ template static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto return matx::sin(v1); } else { + // This should go away once CCCL adds support +#ifdef __CUDACC__ + if constexpr (std::is_same_v) { + return ::sinf(v1); + } + else if constexpr (std::is_same_v) { + return ::sin(v1); + } + else { + matx::sin(v1); + } +#else return cuda::std::sin(v1); +#endif } } template struct SinF { @@ -250,7 +263,20 @@ template static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto return matx::cos(v1); } else { + // This should go away once CCCL adds support +#ifdef __CUDACC__ + if constexpr (std::is_same_v) { + return ::cosf(v1); + } + else if constexpr (std::is_same_v) { + return ::cos(v1); + } + else { + matx::cos(v1); + } +#else return cuda::std::cos(v1); +#endif } } template struct CosF { @@ -691,7 +717,8 @@ static __MATX_INLINE__ __MATX_HOST__ __MATX_DEVICE__ auto _internal_isinf(T v1) return cuda::std::isinf(static_cast(v1.real())) || cuda::std::isinf(static_cast(v1.imag())); } else { return cuda::std::isinf(static_cast(v1)); - } + } + } template struct IsInf { diff --git a/include/matx/operators/select.h b/include/matx/operators/select.h index fa37d006..2ce10a1d 100644 --- a/include/matx/operators/select.h +++ b/include/matx/operators/select.h @@ -69,7 +69,7 @@ namespace matx template __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 index fd7f5157..94d50a7e 100644 --- a/include/matx/operators/self.h +++ b/include/matx/operators/self.h @@ -68,7 +68,7 @@ namespace matx template __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 index 21ab429a..a8ccaac5 100644 --- a/include/matx/operators/shift.h +++ b/include/matx/operators/shift.h @@ -91,7 +91,7 @@ namespace matx template __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 diff --git a/include/matx/operators/slice.h b/include/matx/operators/slice.h index f1baf8dc..4453608c 100644 --- a/include/matx/operators/slice.h +++ b/include/matx/operators/slice.h @@ -140,7 +140,7 @@ namespace matx template __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 index af8ac5f4..2d723975 100644 --- a/include/matx/operators/stack.h +++ b/include/matx/operators/stack.h @@ -147,7 +147,7 @@ namespace matx template __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 index a53e33c3..5527be87 100644 --- 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 namespace matx { diff --git a/include/matx/transforms/reduce.h b/include/matx/transforms/reduce.h index fc231833..19d3e4f2 100644 --- 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(-1), std::numeric_limits::lowest()); + const auto initial_value = thrust::make_tuple(static_cast(-1), std::numeric_limits::lowest()); using reduce_param_type = typename detail::ReduceParams_t; 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(-1), std::numeric_limits::max()); + const auto initial_value = thrust::make_tuple(static_cast(-1), std::numeric_limits::max()); using reduce_param_type = typename detail::ReduceParams_t; 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(-1), std::numeric_limits::max(), static_cast(-1), diff --git a/test/00_operators/OperatorTests.cu b/test/00_operators/OperatorTests.cu index 847bdad1..4d13e29c 100644 --- a/test/00_operators/OperatorTests.cu +++ b/test/00_operators/OperatorTests.cu @@ -180,14 +180,12 @@ TYPED_TEST(OperatorTestsAllExecs, GetString) auto op1 = C = A; auto op2 = C = A + B + (TestType)5; auto op3 = C = A / B; - auto op4 = C = cos(A * B) + sin(C); - auto op6 = (op1,op2,op3,op4); + auto op4 = (op1,op2,op3); std::cout << "op1: " << op1.str() << std::endl; std::cout << "op2: " << op2.str() << std::endl; std::cout << "op3: " << op3.str() << std::endl; std::cout << "op4: " << op4.str() << std::endl; - std::cout << "op6: " << op6.str() << std::endl; MATX_EXIT_HANDLER(); } @@ -518,7 +516,7 @@ TYPED_TEST(OperatorTestsFloatNonComplexAllExecs, FMod) MATX_EXIT_HANDLER(); } -TYPED_TEST(OperatorTestsFloatAllExecs, TrigFuncs) +TYPED_TEST(OperatorTestsFloatNonComplexAllExecs, TrigFuncs) { MATX_ENTER_HANDLER(); using TestType = cuda::std::tuple_element_t<0, TypeParam>; @@ -2105,6 +2103,58 @@ TYPED_TEST(OperatorTestsComplexTypesAllExecs, OperatorFuncDivComplex) MATX_EXIT_HANDLER(); } +TYPED_TEST(OperatorTestsFloatNonComplexAllExecs, IsNanInf) +{ + MATX_ENTER_HANDLER(); + using TestType = cuda::std::tuple_element_t<0, TypeParam>; + using ExecType = cuda::std::tuple_element_t<1, TypeParam>; + using inner_type = typename inner_op_type_t::type; + + ExecType exec{}; + + auto nan = make_tensor({}); + using conversionType = typename matx::detail::value_promote_t; + if constexpr(matx::is_complex_v) { + nan() = TestType(std::numeric_limits::quiet_NaN()); + } else { + nan() = std::numeric_limits::quiet_NaN(); + } + auto tob = make_tensor({}); + // example-begin nan-test-1 + (tob = matx::isnan(nan)).run(exec); + // example-end nan-test-1 + exec.sync(); + EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), std::is_floating_point_v ? true : false)); + + auto notnanorinf = make_tensor({}); + if constexpr(matx::is_complex_v) { + notnanorinf() = TestType(0); + } else { + notnanorinf() = 0; + } + (tob = matx::isnan(notnanorinf)).run(exec); + exec.sync(); + EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), false)); + + auto inf = make_tensor({}); + if constexpr(matx::is_complex_v) { + inf() = TestType(std::numeric_limits::infinity()); + } else { + inf() = std::numeric_limits::infinity(); + } + // example-begin inf-test-1 + (tob = matx::isinf(inf)).run(exec); + // example-end inf-test-1 + exec.sync(); + EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), std::is_floating_point_v ? true : false)); + + (tob = matx::isinf(notnanorinf)).run(exec); + exec.sync(); + EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), false)); + + MATX_EXIT_HANDLER(); +} + TYPED_TEST(OperatorTestsNumericAllExecs, OperatorFuncs) { MATX_ENTER_HANDLER(); @@ -2171,46 +2221,6 @@ TYPED_TEST(OperatorTestsNumericAllExecs, OperatorFuncs) res = c * c * (c + c) / c + three; EXPECT_TRUE(MatXUtils::MatXTypeCompare(tov0(), res, 0.07)); - auto nan = make_tensor({}); - using conversionType = typename matx::detail::value_promote_t; - if constexpr(matx::is_complex_v) { - nan() = TestType(std::numeric_limits::quiet_NaN()); - } else { - nan() = std::numeric_limits::quiet_NaN(); - } - auto tob = make_tensor({}); - // example-begin nan-test-1 - (tob = matx::isnan(nan)).run(exec); - // example-end nan-test-1 - exec.sync(); - EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), std::is_floating_point_v ? true : false)); - - auto notnanorinf = make_tensor({}); - if constexpr(matx::is_complex_v) { - notnanorinf() = TestType(0); - } else { - notnanorinf() = 0; - } - (tob = matx::isnan(notnanorinf)).run(exec); - exec.sync(); - EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), false)); - - auto inf = make_tensor({}); - using conversionType = typename matx::detail::value_promote_t; - if constexpr(matx::is_complex_v) { - inf() = TestType(std::numeric_limits::infinity()); - } else { - inf() = std::numeric_limits::infinity(); - } - // example-begin inf-test-1 - (tob = matx::isinf(inf)).run(exec); - // example-end inf-test-1 - exec.sync(); - EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), std::is_floating_point_v ? true : false)); - - (tob = matx::isinf(notnanorinf)).run(exec); - exec.sync(); - EXPECT_TRUE(MatXUtils::MatXTypeCompare(tob(), false)); MATX_EXIT_HANDLER(); }