use pointer in c++ class

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

Do you have any recommendations for pointer (device) member in a class.

Previously, I cudaMalloc() and cudaMemcpy() the pointers in the constructor function and passed the object to a kernel.

Now I have to pass a dynamic array of this objects. I did like this

multipole **pisotopes; 
  pisotopes = (multipole**)malloc(sizeof(multipole)*numIsos);

  for(int i=0;i<numIsos;i++)
    pisotopes[i]=new multipole(data[i]); 
    //copy data in host struct data[i] to device

And in the kernel, I called the device function of the class

isotope[i]->xs_eval_fast(localenergy, sqrt(300.0*KB), sigT, sigA, sigF);

But program exit at this line.

CUDA Exception: Device Illegal Address
The exception was triggered in device 0.

Program received signal CUDA_EXCEPTION_10, Device Illegal Address.
0x0000000000e174a0 in history<<<(512,1,1),(128,1,1)>>> (Info=…, num_src=1000000, devstep=1, numIso=1, isotope=0xe8d760) at src/simulation.cu:49
49 isotope->xs_eval_fast(localenergy, sqrt(300.0*KB), sigT, sigA, sigF);

And I cannot print the variables
(cuda-gdb) p isotope
$2 = (@generic multipole * @generic * @parameter) 0xe8d760
(cuda-gdb) p isotope[0]
Error: Failed to read generic memory at address 0xe8d760 on device 0 sm 10 warp 1 lane 0 (error=7).
(cuda-gdb) p isotope[0]->dev_double[0]
Error: Failed to read generic memory at address 0xe8d760 on device 0 sm 10 warp 1 lane 0 (error=7).

Any ideas?

Thanks!

2

I’ve found why this error occurs but don’t know how to avoid it.

If I pass the

*pisotopes[0]

to the kernel, it works. So what I did in

new multipole(data[i])

worked as expected.
But neither

pisotopes

or

pisotopes[0]

can be successfully passed to the kernel. I should not access the host pointer in kernel.

3

Something approaching a complete, compilable code would probably help. It’s not even clear to me now what your question is.

Nevertheless, this may help:

[url]c++ - cudaMemcpy segmentation fault - Stack Overflow

4

Sorry I couldn’t make it clear.
I have no problem using pointer in class but using class pointer.

The class has pointer, which should be allocated on device.
Then I need an array of this class. Each object in the array must be initialized with different construction inputs.

I created a complete compilable project on Github.
[url]https://github.com/mjlong/classpointer.git[/url]

Thanks!

5

You cannot use pointers to host allocations on the device, and expect anything useful to come of that. Host pointers cannot be dereferenced in device code, and vice versa.

*pbar[0] works, because that is not a pointer, it is an object of the class data, that you are passing by value, as a kernel parameter (therefore the entire object gets copied to the kernel device code).

It seems the difficulty you are having is around passing an array of objects to the device. In order to avoid making this overly complicated, lets reference the array with a single pointer:

data *pbar

rather than a double pointer:

data **pbar

This means, we will need to make sure the underlying object allocation is contiguous.

Here’s a modified version of your main.cu that should get you going:

#include <stdio.h>
#include <stdlib.h>
#include <data.h>
#include <gpuerrchk.h>

__global__ void kernel(data *obj, double* outarray, int numobj){
  int id = threadIdx.x + blockDim.x*blockIdx.x;
  obj[blockIdx.x].foo(threadIdx.x,outarray[id]);
}

int main(int argc, char* argv[]){
//test data for constructor
  double x = 2.0;
  int i = 2;
  int num = 10;
//N = number of data objects in the array
  int N=3;
  //data bar(x,i,num);
  data* pbar;
  pbar = (data*)malloc(sizeof(data)*N);
  for(int j=0;j<N;j++){
    data temp(x, j, num);
    pbar[j]=temp;} //each data object must be initialized with different input
  double *array = (double*)malloc(sizeof(double)*num*N);
  double *d_array;
  gpuErrchk(cudaMalloc((void**)&d_array, sizeof(double)*num*N));
  data *d_pbar;
  gpuErrchk(cudaMalloc(&d_pbar, N*sizeof(data)));
  gpuErrchk(cudaMemcpy(d_pbar, pbar, N*sizeof(data), cudaMemcpyHostToDevice));


  kernel<<<N,num>>>(d_pbar, d_array, N);

  gpuErrchk(cudaMemcpy(array,d_array,sizeof(double)*num*N,cudaMemcpyDeviceToHost));

  for(i=0;i<num*N;i++)
    printf("array[%d]=%6.2f\n",i,array[i]);

  free(array);
  gpuErrchk(cudaFree(d_array));

  return 0;
}

I’ve made a few other changes, but they were mostly just to help me verify correct operation.

6

I’ve moved away from passing class pointer to kernel. I make the array in class accommodate N arrays of size num. That eases me coding and will help coalesce access in my project.

Thanks a lot all the same!
Your method works well. That teaches me the way of passing class pointer. I considered this way but didn’t think out the method

data temp(x,j,num); pbar[j]=temp;

. And failure to use this lead me to the annoying data** mare.