ULF on a simple example CUDA program

Thares32 · February 17, 2009, 1:35pm

Hallo everyone,

while reading the programming manual I’ve tried to implement a simple parallel vector adding programm: it takes 3 arrays A,B,C with 5 elements each, copies the stuff into global device memory, executes a global kernel on 1 block with 5 threads, this copies the stuff into shared memory and uses 5 threads to add the 5 elements (A[tx]+B[tx]=C[tx]), then the resulting C array is copied back into global memory, the kernel finishes and C is copied into the final host array.

Or so it should happen, but it doesn’t - an “unspecified launch failure” at line 82 (the line where the result matrix is copied from device to host memory).

I’ve tried deviceemulation to check whether the threads go out of array bounds but they don’t.

Valgrind shows a memory misread by the GNU linker.

Any idea what might be going wrong?

Here is the code (I’ve included one call from the cutil at the beginning)

[codebox]#include <stdio.h>

#define ARRSIZE 5

#define CUDA_SAFE_CALL_NO_SYNC( call) do { \

cudaError err = call; \

if( cudaSuccess != err) { \

  fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \

          __FILE__, __LINE__, cudaGetErrorString( err) );              \

  exit(EXIT_FAILURE);                                                  \

} } while (0)

global void vecAdd(float *A,float *B,float *C)

{

int tx=threadIdx.x;

// Declaration of the shared memory array A,B,C

__shared__ float As[ARRSIZE];

__shared__ float Bs[ARRSIZE];

__shared__ float Cs[ARRSIZE];

// Load the arrays from device memory to shared mem,

// each thread loads one element

As[tx]=A[tx];

Bs[tx]=B[tx];

Cs[tx]=C[tx];

// Synchronize to make sure the matrices are loaded

__syncthreads();

//compute teh stuff threadwise

Cs[tx]=As[tx]+Bs[tx];

// And sync…

__syncthreads();

//load C into device mem (each thread one elem)

C[tx]=Cs[tx];

//printf("I'm thread nr. %d and I'm setting C[%d]=%d\n",threadIdx.x,tx,C[tx]);

}

void vecAddser(float *A, float *B, float *C, float n)

{

for (int i=0;i<n;i++) C[i]=A[i]+B[i];

}

int main()

{

float A={1,2,3,4,5};

float B={6,7,8,9,0};

float C={0,0,0,0,0};

float reference={0,0,0,0,0};

float size = ARRSIZE * sizeof(int);

int* Aondev;

int* Bondev;

int* Condev;

//copy the arrays to global mem

CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&Aondev, size));

CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(Aondev, A, size, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&Bondev, size));

CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(Bondev, B, size, cudaMemcpyHostToDevice));

CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void**)&Condev, size));

CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(Condev, C, size, cudaMemcpyHostToDevice));

//Kernel invocation which says: exec vecAdd on a Grid containing one Block which contains 5 Threads

vecAdd<<<1, ARRSIZE>>>(A,B,C);

vecAddser(A,B,reference,ARRSIZE); //serial stuff for reference

// Read C from the device

CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(C, Condev, size, cudaMemcpyDeviceToHost)); //line 82

// Free device memory

cudaFree(Aondev);

cudaFree(Bondev);

cudaFree(Condev);

for (int i=0;i<5;i++) printf(“C[%d]=%d\t”,i,C[i]);

printf(“\n”);

for (int i=0;i<5;i++) printf(“R[%d]=%d\t”,i,reference[i]);

printf(“\nDone.\n”);

}[/codebox]

The output:

[codebox]Cuda error in file ‘vectorf.cu’ in line 82 : unspecified launch failure.[/codebox]

Here is the valgrind output:

[codebox]==19956== Memcheck, a memory error detector.

==19956== Using LibVEX rev 1804, a library for dynamic binary translation.

==19956== Using valgrind-3.3.0-Debian, a dynamic binary instrumentation framework.

==19956== For more details, rerun with: -v

==19956==

==19956== Invalid read of size 4

==19956== at 0x40151E3: (within /lib/ld-2.7.so)

==19956== by 0x4005C59: (within /lib/ld-2.7.so)

==19956== by 0x4007A87: (within /lib/ld-2.7.so)

==19956== by 0x400BC06: (within /lib/ld-2.7.so)