Setting an array in device memory with a pointer to struct; in cuda

Question

I am trying to initialize an array in memory with pointer to a struct that I create inside a kernel. Here is the code I have so far I don't know what I am doing wrong. I get a segmentation fault if I try to do a cudaMalloc on each item in the array, if I don't then I get an "unspecified launch failure" error.

#include <stdio.h>
#include <stdlib.h>
#include <string.h>

typedef struct {
    int status;
    int location;
    double distance;
} Point;

//Macro for checking cuda errors following a cuda launch or api call
#define cudaCheckError() {\
 cudaError_t e=cudaGetLastError();\
 if(e!=cudaSuccess) {\
   printf("\nCuda failure %s:%d: '%s'\n",__FILE__,__LINE__,cudaGetErrorString(e));\
   exit(0); \
 }\
}

__global__ void kernel1(Point** d_memory, int limit){

  int idx = blockIdx.x * blockDim.x * blockDim.y * blockDim.z 
  + threadIdx.z * blockDim.y * blockDim.x 
  + threadIdx.y * blockDim.x + threadIdx.x;

    if(idx < limit) {

        Point* pt = ( Point *) malloc( sizeof(Point) );
        pt->distance = 10;
        pt->location = -1;
        pt->status = -1;

        d_memory[idx] = pt;
    }
}

__global__ void kernel2(Point** d_memory, int limit){
  int i;
  for (i=0; i<limit;i++){
    printf("%f \n",d_memory[i]->distance);
  }
}

int main(int argc, char *argv[])
{
    int totalGrid = 257*193*129;
    size_t size = sizeof(Point) * totalGrid;
    Point ** d_memory;
    cudaMalloc((void **)&d_memory, size);
    /*
    for(int i=0; i<totalGrid; i++){
        printf("%d\n",i);
        cudaMalloc((void **)&d_memory[i], sizeof(Point));
    }*/
    dim3 bs(16,8,8);
    kernel1<<<6249, bs>>>(d_memory, totalGrid);
    cudaCheckError();

    cudaDeviceSynchronize();

    kernel2<<<1,1>>>(d_memory, totalGrid);
    cudaCheckError();

    cudaFree(d_memory);
    return 0;
}

This is what I used for compiling the code

 nvcc -arch=sm_20 test.cu

Iharob Al Asimi · Accepted Answer · 2015-04-17 22:55:02Z

2

I believe your problem is

Point **d_memory;

it should be

Point *d_memory;

and you should not need the cast to void **, you need it in your code because your pointer as passed is a Point *** and not Point **.

Note that cudaMalloc() will allocate contiguous memory, the Point ** suggests that you want an array of pointers, for which I believe you need something like

Point **d_memory;
cudaMalloc((void **)&d_memory, rows);
for (row = 0 ; row < rows ; ++row)
    cudaMalloc(&d_memory[row], columns * sizeof(Point));

But then, you will need to check that the other objects that take d_memory as a parameter, will treat d_memory accordingly.

Also, cudaMalloc() returns cudaSuccess when the allocation was sucessful, you never check for that.

answered Apr 17, 2015 at 22:55

Iharob Al Asimi

53.1k6 gold badges66 silver badges100 bronze badges

Sign up to request clarification or add additional context in comments.

5 Comments

anupshrestha Over a year ago

I do want d_memory to be an array of pointers. I tried your code, where I set rows = 257*193*129 and columns = 1 but I am still getting a Segmentation fault on the first iteration of the for loop. I am not sure what am I missing.

Iharob Al Asimi Over a year ago

It depends on how you access the array of pointers, columns == 1 doesn't make that much sense.

Robert Crovella Over a year ago

Here's a modification of your code that runs "correctly". However it's just flattening your array as suggested by @iharob. (I've also shortened the length passed to kernel2 just to shorten the runtime. Having a single thread iterate through your entire array will take forever and overflow the device printf buffer anyway.) An array of pointers is more difficult to manage when you are transferring data between host and device. I don't suggest it for a beginner, but if you want to tackle it then search on "cuda 2d array" and start reading.

Robert Crovella Over a year ago

The answer given by @talonmies here covers the details of transferring a 2D array (array of pointers) from host to device. It's not trivial. If that's what you want, then your question is arguably a duplicate of that one.

anupshrestha Over a year ago

I will for the time being flatten my array like @iharob and you suggested and do some reading then try the cuda 2D array. Thanks for your help guys.

Collectives™ on Stack Overflow

Setting an array in device memory with a pointer to struct; in cuda

1 Answer 1

5 Comments

Your Answer

Linked

Hot Network Questions

Collectives™ on Stack Overflow

1 Answer 1

5 Comments

Your Answer

Sign up or log in

Post as a guest

Linked

Related