I first define a C++ class that allocates data I will use in the following computation. After that, I want to let the device function directly access the private variable. I get the error of illegal memory access. My target is to let the device function call another class member function which is also a device function. I hope that I am clear. Thank you !
#include <cstddef>
#include <memory>
#include <utility>
#include <iostream>
#include <cuda_runtime.h>
#include <cuda/std/atomic>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
class static_map {
public:
void initialize() {
capacity = 32;
gpuErrchk(cudaMalloc((void**)&vals, capacity * sizeof(int32_t)));
gpuErrchk(cudaMalloc((void**)&keys, capacity * sizeof(cuda::atomic<int32_t>)));
}
__device__ void init() {
int tid = threadIdx.x + blockDim.x * blockIdx.x;
keys[tid].store(1);
vals[tid] = 1;
}
__device__ void access() {
int tid = threadIdx.x + blockDim.x * blockIdx.x;
printf("Thread id %d, local key %d\n", tid, keys[tid].load());
}
private:
cuda::atomic<int32_t>* keys;
int32_t* vals;
int capacity;
};
__global__ void call_device_kernel(static_map* map) {
int tid = threadIdx.x + blockDim.x * blockIdx.x;
map->init();
}
int main() {
static_map* map = new static_map();
map->initialize();
int block_size = 32;
int grid_size = 1;
call_device_kernel<<<grid_size, block_size>>>(map);
gpuErrchk(cudaDeviceSynchronize());
}