1

I am trying to implement an algorithm in cuda and I need to allocate an Array of Pointers that point to an Array of Structs. My struct is, lets say:

    typedef struct {
       float x, y; 
    } point;

I know that If I want to preserve the arrays for multiple kernel calls I have to control them from the host, is that right? The initialization of the pointers must be done from within the kernel. To be more specific, the Array of Struct P will contain random order of cartesian points while the dev_S_x will be a sorted version as to x coordinate of the points in P.

I have tried with:

__global__ void test( point *dev_P, point **dev_S_x) {
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

    dev_P[tid].x = 3.141516;
    dev_P[tid].y = 3.141516;
    dev_S_x[tid] = &dev_P[tid];
   ...
}

and:

 int main( void ) {
     point *P, *dev_P, **S_x, *dev_S_x;
     P   = (point*)  malloc (N * sizeof (point) );
     S_x = (point**) malloc (N * sizeof (point*));

     // allocate the memory on the GPU
     cudaMalloc( (void**)  &dev_P,   N * sizeof(point) );
     cudaMalloc( (void***)  &dev_S_x, N * sizeof(point*));

     // copy the array P to the GPU
     cudaMemcpy( dev_P, P,  N * sizeof(point),  cudaMemcpyHostToDevice);
     cudaMemcpy( dev_S_x,S_x,N * sizeof(point*), cudaMemcpyHostToDevice);

     test <<<1, 1 >>>( dev_P, &dev_S_x);
        ...
     return 0;
}

which leads to many

First-chance exception at 0x000007fefcc89e5d (KernelBase.dll) in Test_project_cuda.exe: Microsoft C++ exception: cudaError_enum at memory location 0x0020f920.. Critical error detected c0000374

Am I doing something wrong in the cudamalloc of the array of pointers or is it something else? Is the usage of (void***) correct? I would like to use for example dev_S_x[tid]->x or dev_S_x[tid]->y from within the kernels pointing to device memory addresses. Is that feasible? Thanks in advance

6
  • Leaving aside the other syntax errors, why are you passing a pointer to dev_S_x (which isn't defined anywhere) to the kernel? That is a host address.
    – talonmies
    Jul 31, 2013 at 9:49
  • @talonmies Sorry I forgot to write *dev_S_x in main. Is that what you meant?
    – BugShotGG
    Jul 31, 2013 at 9:54
  • 1
    Why are you defining your own struct of couples of floats and you are not using float2?
    – Vitality
    Jul 31, 2013 at 9:57
  • @JackOLantern You are very right about this but its not the main issue here. Lets assume that this is not a problem and in real usage the struct has more than 2 x, y floats.
    – BugShotGG
    Jul 31, 2013 at 10:00
  • @GeoPapas: For future reference, if you want to ask a question about why code doesn't work, please post the shortest, simplest, complete and compilable code which will reproduce the problem. It is often very difficult to answer these sorts of questions otherwise.
    – talonmies
    Jul 31, 2013 at 10:52

1 Answer 1

1

dev_S_x should be declared as point ** and should be passed to the kernel as a value (i.e. test <<<1, 1 >>>(dev_P, dev_S_x);).

Putting that to one side, what you describe sounds like a natural fit for Thrust, which will give you a simpler memory management strategy and access to fast sort routines.

2
  • @Tom Thanks I was unsure of how & works in kernel call. Now it works. Regarding Thrust, I may use it if I am stuck on my implementation.
    – BugShotGG
    Jul 31, 2013 at 10:18
  • Glad it helped. You could also consider comparing with Thrust anyway, since it allows you to focus on other parts of your algorithm by providing fast, optimised sort routines. Good luck!
    – Tom
    Jul 31, 2013 at 10:35

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.