4

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
xsami
  • 1,312
  • 16
  • 31
anupshrestha
  • 236
  • 5
  • 19

1 Answers1

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.

Iharob Al Asimi
  • 52,653
  • 6
  • 59
  • 97
  • 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. – anupshrestha Apr 17 '15 at 23:55
  • It depends on how you access the array of pointers, `columns == 1` doesn't make that much sense. – Iharob Al Asimi Apr 18 '15 at 00:06
  • 2
    [Here's](http://pastebin.com/uckXpMr1) 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 Apr 18 '15 at 01:05
  • 2
    The answer given by @talonmies [here](http://stackoverflow.com/questions/6137218/how-can-i-add-up-two-2d-pitched-arrays-using-nested-for-loops) 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. – Robert Crovella Apr 18 '15 at 01:09
  • 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. – anupshrestha Apr 19 '15 at 15:52