Cuda passing a class to a kernel

[user137079]

Hi all,

I need a templated kernel that use a c++ class as parameter. So i try a short code like this

#include <cstring>
#include <iostream>

typedef struct
{
	double ljWell;
	int n;
	double Rc;
	int first;
	int second;
} ljt;

class A
{
public:
	ljt data;
	__device__ void test()
	{
		printf("just an int please n:%d\n",0);//works
		printf("just an int please n:%d\n",this->data.n);//crash
	}
	__host__ void fillData()
	{
		std::cout<<"A allocate data:"<<cudaMallocManaged((void**)&data,sizeof(ljt))<<"\n";
		cudaDeviceSynchronize();
		data.ljWell = 10.0;
		data.n = 5;
		data.Rc = 3.2;
		data.first = 1;
		data.second = 0;
	}		
};

template <typename T>
__global__ void foo(T* f)
{
	f->test();
}

int main()
{
	auto toto = new A;
	toto->fillData();
	std::cout<<"CPU data read n:"<<toto->data.n<<"\n";
	foo<<<1,1>>>(toto);
	cudaDeviceSynchronize();
	cudaError_t code = cudaGetLastError( );
	std::cout<<"Kernel result :"<<code<<"\n";
	delete(toto);
	return 0;
}

the allocation of data works on CPU side however, if i try to read any value from the kernel foo(), an error 719 occur saying that bad memory access occur.

Wherre i’m wrong?

[Robert_Crovella]

@user137079

Wherre i’m wrong?

new returns a host pointer. So toto points to host memory.

In CUDA, a host pointer cannot be safely dereferenced in device code. This is dereferencing a host pointer:

f->test();
^ ^
| dereference operator
host pointer

One possible approach to fix this would be to copy the object pointed to by toto to device memory (and pass a pointer to that object to your kernel code.)

Another possible approach to fix would be to allocate the toto pointer pointing to managed memory, rather than host memory. You seem to already know how to do that, since you are using cudaMallocManaged already.

A third possible approach would be to pass the object by value, rather than by pointer. Then you would perform f.test() rather than f->test() in kernel code.