[BUG] Const operators don't work with argmax

[deanljohnson]

Describe the Bug
Using argmax with operators of a const data type fails to compile.

To Reproduce

1. Call make_tensor with a const T*
2. Call argmax

This will produce an error such as:

[build] /usr/local/include/matx/transforms/reduce.h(1173): error: expression must be a modifiable lvalue
[build]         smem[threadIdx.x / 32] = in_val;
[build]         ^
[build]           detected during:
[build]             instantiation of "void matx::detail::matxReduceKernel(OutType, InType, ReduceOp, matx::index_t) [with OutType=matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, InType=matx::detail::tensor_impl_t<const float, 1, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, ReduceOp=matx::detail::reduceOpMax<float>]" at line 1395
[build]             instantiation of "void matx::reduce(OutType, TensorIndexType, const InType &, ReduceOp, cudaStream_t, __nv_bool) [with OutType=matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, TensorIndexType=matx::detail::tensor_impl_t<matx::index_t, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, InType=matx::detail::tensor_impl_t<const float, 1, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, ReduceOp=matx::detail::reduceOpMax<float>, <unnamed>=true]" at line 2088
[build]             instantiation of "void matx::argmax_impl(OutType, TensorIndexType &, const InType &, matx::cudaExecutor) [with OutType=matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, TensorIndexType=matx::detail::tensor_impl_t<matx::index_t, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, InType=matx::detail::tensor_impl_t<const float, 1, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>]" at line 69 of /usr/local/include/matx/operators/argmax.h
[build]             instantiation of "void matx::detail::ArgMaxOp<OpA, ORank>::Exec(Out &&, Executor &&) const [with OpA=const matx::tensor_t<const float, 1, matx::basic_storage<matx::raw_pointer_buffer<const float, matx::matx_allocator<const float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>> &, ORank=0, Out=cuda::std::__4::tuple<matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::tensor_impl_t<matx::index_t, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::ArgMaxOp<const matx::tensor_t<const float, 1, matx::basic_storage<matx::raw_pointer_buffer<const float, matx::matx_allocator<const float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>> &, 0>> &, Executor=matx::cudaExecutor &]" at line 100 of /usr/local/include/matx/core/tie.h
[build]             instantiation of "void matx::mtie<Ts...>::Exec(Executor &&) [with Ts=<matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::tensor_impl_t<matx::index_t, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::ArgMaxOp<const matx::tensor_t<const float, 1, matx::basic_storage<matx::raw_pointer_buffer<const float, matx::matx_allocator<const float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>> &, 0>>, Executor=matx::cudaExecutor &]" at line 73 of /usr/local/include/matx/operators/base_operator.h
[build]             instantiation of "void matx::BaseOp<T>::run(Ex &&) [with T=matx::mtie<matx::detail::tensor_impl_t<float, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::tensor_impl_t<matx::index_t, 0, matx::tensor_desc_t<cuda::std::__4::array<long long, 0UL>, cuda::std::__4::array<long long, 0UL>, 0>>, matx::detail::ArgMaxOp<const matx::tensor_t<const float, 1, matx::basic_storage<matx::raw_pointer_buffer<const float, matx::matx_allocator<const float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>> &, 0>>, Ex=matx::cudaExecutor]" at line 112 of /usr/local/include/matx/operators/base_operator.h

Expected Behavior
I would expect argmax to work with both const and non-const inputs

Code Snippets

const float* gpu_data = ...;
auto t_data = matx::make_tensor(gpu_data, {10});
auto data_mag = matx::make_tensor<float>({});
auto data_idx = matx::make_tensor<matx::index_t>({});
(matx::mtie(data_mag , data_idx ) = matx::argmax(t_data)).run(stream);

System Details (please complete the following information):

• OS: Rocky 9
• CUDA version: CUDA 12.6 12.3
• g++ version: 11.4.1

Additional Context
None

[cliffburdick]

Hi @deanljohnson , can you please try this fix: #879. I was not able to get the same error you got, but I got a different error related to mixing const and non-const variables.

[deanljohnson]

Hi @cliffburdick. I tried that fix and I still get the same error. Some minor edits to my code snippet make it directly copy-pasteable to get the error:

const float* gpu_data = nullptr; // Compile time error, so doesn't matter that this is nullptr
auto t_data = matx::make_tensor(gpu_data, {10});
auto data_mag = matx::make_tensor<float>({});
auto data_idx = matx::make_tensor<matx::index_t>({});
(matx::mtie(data_mag , data_idx ) = matx::argmax(t_data)).run();

It seems like to me that using value_type = typename InType::value_type; in matxReduceKernel is resolving to const float in my example which is what is causing the problem since then the smem values are const.

[cliffburdick]

Hi @cliffburdick. I tried that fix and I still get the same error. Some minor edits to my code snippet make it directly copy-pasteable to get the error:

const float* gpu_data = nullptr; // Compile time error, so doesn't matter that this is nullptr
auto t_data = matx::make_tensor(gpu_data, {10});
auto data_mag = matx::make_tensor<float>({});
auto data_idx = matx::make_tensor<matx::index_t>({});
(matx::mtie(data_mag , data_idx ) = matx::argmax(t_data)).run();

It seems like to me that using value_type = typename InType::value_type; in matxReduceKernel is resolving to const float in my example which is what is causing the problem since then the smem values are const.

Hi @deanljohnson , I cannot reproduce it with my compilers:

I have no name!@62e8db451b34:/repro/tmp/MatX/build$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Jan_15_19:20:09_PST_2025
Cuda compilation tools, release 12.8, V12.8.61
Build cuda_12.8.r12.8/compiler.35404655_0

gcc (Ubuntu 11.4.0-1ubuntu1~22.04) 11.4.0

I can try CTK 12.6 and see if it compiles, but CTK 11.8 seems to work. Our gcc is almost identical.

[deanljohnson]

I am using the latest MatX release v0.9 - perhaps something related has changed on MatX master that fixes the issue. I will try that at some point.

[cliffburdick]

Yes, please try that. I just tried CUDA 12.6 and it also compiles fine.

[deanljohnson]

I tried with matx master and got the same error. However, I realized I actually had the CUDA version I was using wrong. I am using CUDA 12.3.

[deanljohnson]

Note I also get a const related compile error when calling pwelch with a const time domain input:

[build] /usr/local/include/matx/operators/pwelch.h(135): error: no instance of overloaded function "matx::matxFree" matches the argument list
[build]             argument types are: (const cuda::std::__4::complex<float> *)
[build]             matxFree(ptr);
[build]             ^
[build]           detected during:
[build]             instantiation of "void matx::detail::PWelchOp<OpX, OpW>::PostRun(ShapeType &&, Executor &&) const noexcept [with OpX=matx::tensor_t<const cuda::std::__4::complex<float>, 1, matx::basic_storage<matx::raw_pointer_buffer<const cuda::std::__4::complex<float>, matx::matx_allocator<const cuda::std::__4::complex<float>>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, OpW=matx::detail::matxGenerator1D_t<matx::detail::Hamming<float>, 0, cuda::std::__4::array<long long, 1UL>>, ShapeType=cuda::std::__4::array<long long, 1UL>, Executor=matx::cudaExecutor &]" at line 180 of /usr/local/include/matx/operators/set.h
[build]             instantiation of "void matx::detail::set<T, Op>::TransformExec(ShapeType &&, Executor &&) const noexcept [with T=matx::tensor_t<float, 1, matx::basic_storage<matx::raw_pointer_buffer<float, matx::matx_allocator<float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, Op=matx::detail::PWelchOp<matx::tensor_t<const cuda::std::__4::complex<float>, 1, matx::basic_storage<matx::raw_pointer_buffer<const cuda::std::__4::complex<float>, matx::matx_allocator<const cuda::std::__4::complex<float>>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, matx::detail::matxGenerator1D_t<matx::detail::Hamming<float>, 0, cuda::std::__4::array<long long, 1UL>>>, ShapeType=cuda::std::__4::array<long long, 1UL>, Executor=matx::cudaExecutor &]" at line 77 of /usr/local/include/matx/operators/base_operator.h
[build]             instantiation of "void matx::BaseOp<T>::run(Ex &&) [with T=matx::detail::set<matx::tensor_t<float, 1, matx::basic_storage<matx::raw_pointer_buffer<float, matx::matx_allocator<float>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, matx::detail::PWelchOp<matx::tensor_t<const cuda::std::__4::complex<float>, 1, matx::basic_storage<matx::raw_pointer_buffer<const cuda::std::__4::complex<float>, matx::matx_allocator<const cuda::std::__4::complex<float>>>>, matx::tensor_desc_t<cuda::std::__4::array<long long, 1UL>, cuda::std::__4::array<long long, 1UL>, 1>>, matx::detail::matxGenerator1D_t<matx::detail::Hamming<float>, 0, cuda::std::__4::array<long long, 1UL>>>>, Ex=matx::cudaExecutor]" at line 112 of /usr/local/include/matx/operators/base_operator.h

[deanljohnson]

Retried this with matx master and I no longer get the compile error with argmax