[BUG] Tensors are not "modifiable lvalue" in extended lambdas. about matx HOT 4 CLOSED

Hi @sksg, when you're passing a tensor to a device/global function the tensor must first be converted to tensor_impl_t. The reason is that tensor_t maintains some state, like a reference count, that is not usable or wanted in device code. tensor_impl_t strips a tensor to the bare minimum needed to be used on device code. To convert, you use the base_type trait:

So your code may be something like

int main() {
    auto a = matx::make_tensor<float>({10});
    auto a_base = base_type<decltype(a)>(a);
    kernel_launch<<<0, 10>>>([=] __device__() { a_base(threadIdx.x) = threadIdx.x; });  // <-- error line
}

I have not tried this yet to see if it works, but I will test and get back to you today.

from matx.

Hi @sksg, the issue is that there are two definitions of operator(), a const and a non-const version. Since lambdas capture by const, it's choosing the const version that cannot modify an lvalue. Instead you can pass it as a function parameter:

using namespace matx;
template<typename T, typename FUNC>
__global__ void kernel_launch(T t, FUNC kernel) {
    kernel(t);
}

    kernel_launch<<<0, 10>>>(a_base, [=] __device__(auto ab) { 
      ab(threadIdx.x) = threadIdx.x; });

And this works as well. Let me know if that suffices.

from matx.

Hi @cliffburdick

Thanks for the speedy feedback.

I did not know that lambdas were const by default. It then makes perfect sense. I did try it with a mutable lambda instead:

#include <matx.h>

template<typename FUNC>
__global__ void kernel_launch(FUNC kernel) {
    kernel();
}

int main() {
    auto a = matx::make_tensor<float>({10});
    kernel_launch<<<0, 10>>>([=] __device__() mutable { a(threadIdx.x) = threadIdx.x; });  // <-- No errors!
}

Technically, I can verify that the device array a has the correct values. But I see your point wrt. matx::base_type. I do not know if any of the host state will be corrupted by using a directly rather than a_base. The above syntax is very concise, though.

Thanks for helping out. This is perfectly usable for me.

from matx.

Thanks for the report!

from matx.