I'm trying to create a struct of arrays with arrays inside and load them onto the GPU. I think I followed the steps to do this correctly.
- Create a struct on the CPU using malloc.
cudamallocthe arrays to the struct.- Create a struct on the GPU using cudamalloc
- Copy the CPU struct onto the GPU struct.
When I run this code, it will correctly work as long as I don't change the value p[i].c[0] in the kernel function. If I delete the line p[i].c[0] = 3.3; then it outputs the expected results. When I leave it as is, it outputs random numbers for all of the values. I would like to be able to update the values in the array using the kernel function.
What could be wrong?
Here is my code:
#include <stdio.h>
#include <cuda_runtime.h>
#include <iostream>
#include <fstream>
#include <sstream>
#include <cstdio>
#include <fcntl.h>
#include <unistd.h>
#include <assert.h>
#include <omp.h>
#include <vector>
#include <sys/time.h>
float cData[]
{
1,
2,
3,
4,
5,
6,
7,
8,
9,
10,
11,
12,
13,
14,
15,
16
};
float dData[]
{
1,
2,
3,
4,
5,
6,
7,
8,
9,
10,
11,
12,
13,
14,
15,
16
};
typedef struct
{
float a, b;
float* c;
float* d;
} point;
__global__ void testKernel(point *p){
int i = blockIdx.x * blockDim.x + threadIdx.x;
p[i].a = 1.1;
p[i].b = 2.2;
p[i].c[0] = 3.3;
}
void checkerror(cudaError_t error, char* descrp){
if (error != 0){
printf("%s error code: %d \n", descrp, error);
}
}
extern "C" int main()
{
printf("starting gpuCode\n");
int *dev_a;
// set number of points
int numPoints = 16,
gpuBlockSize = 4,
pointSize = sizeof(point),
numBytes = numPoints * pointSize,
gpuGridSize = numPoints / gpuBlockSize;
cudaError_t err = cudaSuccess;
printf("initialized variables\n");
// allocate memory
point *cpuPointArray,
*gpuPointArray,
*outPointArray;
cpuPointArray = (point*)malloc(numBytes); //create the cpuPointArray struct on the cpu
outPointArray = (point*)malloc(numBytes); //create the outPointArray struct on the cpu
printf("load cpuPointArray struct with default values\n");
for (int k=0; k<16; k++){
err = cudaMalloc( (void**)&cpuPointArray[k].c, 16*sizeof(float) );
checkerror(err, "assigning cuda pointer c");
err = cudaMalloc( (void**)&cpuPointArray[k].d, 16*sizeof(float) );
checkerror(err, "assigning cuda pointer d");
cpuPointArray[k].a = 16;
cpuPointArray[k].b = 16;
}
for (int k=0; k<16; k++){
printf("top loop %d\n", k);
err = cudaMemcpy(cpuPointArray[k].c, cData, 16*sizeof(float), cudaMemcpyHostToDevice);
printf("after cdata\n");
checkerror(err, "copying cdata to gpu array c" );
err = cudaMemcpy(cpuPointArray[k].d, dData, 16*sizeof(float), cudaMemcpyHostToDevice);
printf("after ddata\n");
checkerror(err, "copying ddata to gpu array d");
printf("bottom of loop %d\n", k);
}
err = cudaMalloc((void**)&gpuPointArray, numBytes); //allocate memory on the gpu for the cpu point array
checkerror(err, "allocating memory for gpuPointArray");
err = cudaMemcpy(gpuPointArray,cpuPointArray,sizeof(cpuPointArray), cudaMemcpyHostToDevice); //copy the cpu point array onto the gpu
checkerror(err, "copying cpuPointArray to gpuPointArray");
printf("loaded the struct into the kernel\n");
for(int i = 0; i < numPoints; ++i)
{
printf("point.a: %f, point.b: %f ************************\n",cpuPointArray[i].a,cpuPointArray[i].b);
printf("cuda mem location point.c: %d point.d: %d\n",&cpuPointArray[i].c, &cpuPointArray[i].d);
}
// launch kernel
testKernel<<<gpuGridSize,gpuBlockSize>>>(gpuPointArray);
printf("returned the struct from the kernel\n");
err = cudaMemcpy(outPointArray,gpuPointArray,numBytes, cudaMemcpyDeviceToHost);
checkerror(err, "copying gpuPointArray to cpuPointArray");
printf("after gpu copy to cpu\n");
for (int k=0; k<16; k++){
printf("creating memory on cpu for array c\n");
outPointArray[k].c = (float*)malloc(16*sizeof(float));
printf("creating memory on cpu for array d\n");
outPointArray[k].d = (float*)malloc(16*sizeof(float));
printf("copying memory values onto cpu array c\n");
err = cudaMemcpy(outPointArray[k].c, cpuPointArray[k].c, 16*sizeof(float), cudaMemcpyDeviceToHost);
checkerror(err, "copy array c from gpu to cpu");
printf("copying memory values onto cpu array c\n");
err = cudaMemcpy(outPointArray[k].d, cpuPointArray[k].d, 16*sizeof(float), cudaMemcpyDeviceToHost);
checkerror(err, "copy array d from gpu to cpu");
printf("bottom of loop %d\n", k);
}
// retrieve the results
printf("testKernel results:\n");
for(int i = 0; i < numPoints; ++i)
{
printf("point.a: %f, point.b: %f ************************\n",outPointArray[i].a,outPointArray[i].b);
for (int j=0; j<16; j++){
printf("point.c: %f point.d: %f\n",outPointArray[i].c[j], outPointArray[i].d[j]);
}
}
// deallocate memory
free(cpuPointArray);
cudaFree(gpuPointArray);
return 0;
}
