# driver.cu(255) : cudaSafeCall() Runtime API error : unspecified launch failure.

**URL:** <https://forums.developer.nvidia.com/t/driver-cu-255-cudasafecall-runtime-api-error-unspecified-launch-failure/16962>\
**Category:** CUDA Programming and Performance\
**Created:** [June 3, 2010, 8:50pm UTC](https://forums.developer.nvidia.com/t/driver-cu-255-cudasafecall-runtime-api-error-unspecified-launch-failure/16962 "2010-06-03T20:50:59Z")\
**Posts on this page:** 1\
**Page:** 1

<div class="post-metadata">

**Author:** ![japnik\_Singh](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@japnik\_Singh](https://forums.developer.nvidia.com/u/japnik_Singh)\
**Post date:** [June 3, 2010, 8:50pm UTC](https://forums.developer.nvidia.com/t/driver-cu-255-cudasafecall-runtime-api-error-unspecified-launch-failure/16962/1 "2010-06-03T20:50:59Z")

</div>

q1) how shall i allocate in device memory a sturct ALLNODES which itself contains pointer to array of struct NODES

//real\_t means FLOAT

typedef struct {

real\_t \*x; // X coordinate of the points

real\_t \*y; // Y coordinate of the points

real\_t \*z; // Z coordinate of the points

real\_t \*den\_pot; // Density/Potential of the points

int num\_pts; // Number of points in the node

} Node;

typedef struct {

int num\_nodes; // Total number of nodes

Node \*N;

real\_t \*xbuffer;

real\_t \*ybuffer;

real\_t \*zbuffer;

real\_t \*den\_potbuffer;

int \*counts;

} AllNodes;

q2) How to debug cuda code in MAC…

my code is

[codebox]#include \<assert.h\>

#include \<stdio.h\>

#include \<stdlib.h\>

//#include \<sys/time.h\>

#include “node.h”

#include “reals.h”

#include “reals\_aligned.h”

#include “evaluate.h”

// includes, kernels

#include \<japnik\_kernel.cu\>

// includes, project

#include \<cutil\_inline.h\>

static void usage (const char\* use) {

fprintf (stderr, “usage: %s \<source\_filename\> \<target\_filename\>”

" \<Ulist\_filename\> \n", use);

}

int main (int argc, char\* argv)

{

AllNodes \*src\_node;

AllNodes \*trg\_node;

AllLists \*ulist;

const char\* input\_src\_filename;

const char\* input\_trg\_filename;

const char\* input\_ulist\_filename;

int src\_node\_count, trg\_node\_count;

int global\_count;

FILE \*fout;

if (argc != 4) {

usage (argv[0]);

return -1;

}

/\*\* Command line input \*/

input\_src\_filename = argv[1];

input\_trg\_filename = argv[2];

input\_ulist\_filename = argv[3];

/\* /CLOCKING; //removed

struct stopwatch\_t\* timer = NULL;

long double t\_u;

stopwatch\_init ();

timer = stopwatch\_create ();

\*/

/\*\* Setup to load points and create data structure for nodes/lists \*/

src\_node = load\_nodes (input\_src\_filename);

trg\_node = load\_nodes (input\_trg\_filename);

trg\_node\_count = trg\_node-\>num\_nodes;

src\_node\_count = src\_node-\>num\_nodes;

ulist = load\_lists (input\_ulist\_filename, trg\_node\_count);

fprintf(stderr, “Finished setup\n”);

/\* Assign random values to all the target points \*/

global\_count = 0;

for (int i = 0; i \< trg\_node\_count; i++) {

get\_value(trg\_node-\>counts[i], trg\_node-\>xbuffer + global\_count);

get\_value(trg\_node-\>counts[i], trg\_node-\>ybuffer + global\_count);

get\_value(trg\_node-\>counts[i], trg\_node-\>zbuffer + global\_count);

set\_zero(trg\_node-\>counts[i], trg\_node-\>den\_potbuffer + global\_count);

global\_count += trg\_node-\>counts[i];

}

global\_count = 0;

for (int i = 0; i \< trg\_node\_count; i++) {

trg\_node-\>N[i].x = trg\_node-\>xbuffer + global\_count;

trg\_node-\>N[i].y = trg\_node-\>ybuffer + global\_count;

trg\_node-\>N[i].z = trg\_node-\>zbuffer + global\_count;

trg\_node-\>N[i].den\_pot = trg\_node-\>den\_potbuffer + global\_count;

trg\_node-\>N[i].num\_pts = trg\_node-\>counts[i];

global\_count += trg\_node-\>counts[i];

}

/\* Assign random values to all the source points \*/

global\_count = 0;

for (int i = 0; i \< src\_node\_count; i++) {

get\_value(src\_node-\>counts[i], src\_node-\>xbuffer + global\_count);

get\_value(src\_node-\>counts[i], src\_node-\>ybuffer + global\_count);

get\_value(src\_node-\>counts[i], src\_node-\>zbuffer + global\_count);

get\_value(src\_node-\>counts[i], src\_node-\>den\_potbuffer + global\_count);

global\_count += src\_node-\>counts[i];

}

global\_count = 0;

for (int i = 0; i \< src\_node\_count; i++) {

src\_node-\>N[i].x = src\_node-\>xbuffer + global\_count;

src\_node-\>N[i].y = src\_node-\>ybuffer + global\_count;

src\_node-\>N[i].z = src\_node-\>zbuffer + global\_count;

src\_node-\>N[i].den\_pot = src\_node-\>den\_potbuffer + global\_count;

src\_node-\>N[i].num\_pts = src\_node-\>counts[i];

global\_count += src\_node-\>counts[i];

}

for (int i = 0; i \< trg\_node\_count; i++) {

ulist-\>L[i].node\_list = ulist-\>nodelist\_buffer + ulist-\>offsets[i];

ulist-\>L[i].num\_listnodes = ulist-\>counts[i];

}

printf(“random number assigned inhost memory”);

//…for allocating and copying memory to device

int BLOCK\_SIZE=8;

int threadsPerBlock= BLOCK\_SIZE \* BLOCK\_SIZE;

//for target

size\_t size=0; //just chk size\_t is predefined

Node N[trg\_node\_count]; //to be passed as argument to multiplicat kernel

for(int i=0; i\<trg\_node\_count;i++)

{

size=trg\_node-\>N[i].num\_pts\*sizeof(real\_t); //assigning sapce to node to the nearest multiple of thredsperblock

N[i].num\_pts=trg\_node-\>N[i].num\_pts;

cutilSafeCall(cudaMalloc((void\*\*)&N[i].x,size));

cutilSafeCall(cudaMalloc((void\*\*)&N[i].y,size));

cutilSafeCall(cudaMalloc((void\*\*)&N[i].z,size));

cutilSafeCall(cudaMalloc((void\*\*)&N[i].den\_pot,size));

cutilSafeCall(cudaMemcpy(N[i].x,trg\_node-\>N[i].x,size,cudaMemcpyHostToDevice));

cutilSafeCall(cudaMemcpy(N[i].y,trg\_node-\>N[i].y,size,cudaMemcpyHostToDevice));

cutilSafeCall(cudaMemcpy(N[i].z,trg\_node-\>N[i].z,size,cudaMemcpyHostToDevice));

//cudaMemcpy(dev\_ptr\_trg\_den[i],trg\_node-\>N[i].den\_pot,size,cudaMemcpyHostToDevice); //to be calculted

}

//…

//for source

AllNodes src; //to be passed to kernel

//src.counts=src\_node-\>counts;

src.num\_nodes=src\_node-\>num\_nodes;

size=src.num\_nodes\*sizeof(Node);

//printf (“%s \n”, “A string”); //debuggin tym

cutilSafeCall(cudaMalloc((void\*\*)&src.N,size));

//cutilSafeCall(cudaMemcpy((void\*\*)&src.N,src\_node-\>N,size,cudaMemcpyHostToDevice)); //ques if num\_pts in each node gets copied

for(int i=0; i\<src\_node\_count;i++)

{

size=src\_node-\>N[i].num\_pts\*sizeof(real\_t);

cutilSafeCall(cudaMalloc((void\*\*)&src.N[i].x,size));

cutilSafeCall(cudaMalloc((void\*\*)&src.N[i].y,size));

cutilSafeCall(cudaMalloc((void\*\*)&src.N[i].z,size));

cutilSafeCall(cudaMalloc((void\*\*)&src.N[i].den\_pot,size));

cutilSafeCall(cudaMemcpy(src.N[i].x,src\_node-\>N[i].x,size,cudaMemcpyHostToDevice));

cutilSafeCall(cudaMemcpy(src.N[i].y,src\_node-\>N[i].y,size,cudaMemcpyHostToDevice));

cutilSafeCall(cudaMemcpy(src.N[i].z,src\_node-\>N[i].z,size,cudaMemcpyHostToDevice));

cutilSafeCall(cudaMemcpy(src.N[i].den\_pot,src\_node-\>N[i].den\_pot,size,cudaMemcpyHostToDevice));

}

//…

//for ulist

List L[trg\_node\_count];

for(int i=0;i\<trg\_node\_count;i++)

{

L[i].num\_listnodes=ulist-\>L[i].num\_listnodes;

size=ulist-\>L[i].num\_listnodes\*sizeof(int);

cutilSafeCall(cudaMalloc((void\*\*)&L[i].node\_list,size));

cutilSafeCall(cudaMemcpy(L[i].node\_list,ulist-\>L[i].node\_list,size,cudaMemcpyHostToDevice));

}

printf(“\n memory allocation done for device \n”);

//memory allocation done…

dim3 dimBlock;

dimBlock.x=BLOCK\_SIZE;

dimBlock.y=BLOCK\_SIZE;

for(int i=0;i\<trg\_node\_count;i++)

{

printf(“enterin for loop %i”,i);

size=src\_node-\>N[i].num\_pts \* sizeof(real\_t);

dim3 dimgrid;

dimgrid.x= (trg\_node-\>N[i].num\_pts + threadsPerBlock - 1) / threadsPerBlock ; //need to ckeck again…if num blocks

dir\_eval\<\<\<dimgrid,dimBlock\>\>\>(N[i],L[i],src,threadsPerBlock);

printf(" copying memeory from device to host \n");

cutilSafeCall(cudaMemcpy(trg\_node-\>N[i].den\_pot,N[i].den\_pot,size,cudaMemcpyDeviceToHost));

}

/\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*

IN LAST LINE OF ABOVE FOR LOOP I M GETTING THE ERROR…driver.cu(255) : cudaSafeCall() Runtime API error : unspecified launch failure.

\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*/

//freeing memory

```
//targets and ulists

for(int i=0; i<trg_node_count;i++)

{

	cudaFree(N[i].x);

	cudaFree(N[i].y);

	cudaFree(N[i].z);

	cudaFree(N[i].den_pot);

	cudaFree(L[i].node_list);

}

//source

cudaFree(src.N);

for (int i=0; i<src_node_count; i++) 

{

	cudaFree(src.N[i].x);		

	cudaFree(src.N[i].y);		

	cudaFree(src.N[i].z);		

	cudaFree(src.N[i].den_pot);		

}

```

//\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*\*

//kernel its another file but i am pasting it here

#ifndef M\_PI

#define M\_PI 3.1415926535897932385

#endif

#define OOFP 1.0/(4.0 \* M\_PI)

**host**  **device** int nearthreadsPerBlock(int num,int threadsPerBlock)

{

```
while (num%threadsPerBlock!=0) {

	num++;

}

return num;

```

}

**global** void dir\_eval(Node A,List B,AllNodes C,int threadsPerBlock)

{

int i=blockIdx.x \* threadsPerBlock + (threadIdx.y\*blockDim.x)+threadIdx.x;

//printf(“%i”,i);

if(i\<A.num\_pts)

{

A.den\_pot[i]=7; //initializing it to 0… actually this is to be found

}

for(int j=0;j\<B.num\_listnodes;j++)

{

A.den\_pot[i]=A.den\_pot[i]+1;

for(int k=0;k\<nearthreadsPerBlock(C.N[B.node\_list[j]].num\_pts,threadsPe

rBlock)/threadsPerBlock;k++)

{

**shared** float4 point\_src[64] ; //number of threads in a block

int l=i%threadsPerBlock;

point\_src[l].x=C.N[B.node\_list[j]].x[(threadsPerBlock\*k)+l]; // each thread puts one paoints stuff in shared memory

point\_src[l].y=C.N[B.node\_list[j]].y[k\*threadsPerBlock+l]; // each thread puts one paoints stuff in shared memory

point\_src[l].z=C.N[B.node\_list[j]].z[k\*threadsPerBlock+l]; // each thread puts one paoints stuff in shared memory

point\_src[l].w=C.N[B.node\_list[j]].den\_pot[k\*threadsPerBlock

+l]; // each thread puts one paoints stuff in shared memory

\_\_syncthreads(); //to enure that all threads 've loaded shard memory

for(int m=0;m\< threadsPerBlock;m++)

{

```
if(i<A.num_pts){

```

real\_t x,y,z,r,r2;

//oofp = (real\_t) (1.0/(4.0 \* M\_PI));

x=A.x[i]-point\_src[m].x;

y=A.y[i]-point\_src[m].y;

z=A.z[i]-point\_src[m].z;

r2=x_x + y_y + z\*z;

r=sqrt(r2); //assuming sqrt predefined

A.den\_pot[i]+=(OOFP/r)\*point\_src[m].w; //finall answer

```
}

```

}

\_\_syncthreads(); //to ensure all threads are done with this calc

}

}

}

printf(“done”);

fout = fopen (“direct\_japnik”, “w”);

for (int i = 0; i \< trg\_node\_count; i++) {

for (int j = 0; j \< trg\_node-\>N[i].num\_pts; j++) {

fprintf (fout, "%lf ", trg\_node-\>N[i].den\_pot[j]);

}

fprintf (fout, “\n”);

}

fclose (fout);

//stopwatch\_destroy (timer);

free\_lists(ulist);

free\_nodes(src\_node); free\_nodes(trg\_node);

free(ulist);

free(src\_node); free(trg\_node);

return 0;

}

/\* ----------------------------------------------------------------------------------------------------------

- eof

\*/

/\*

- japnik.h

- 
- 
- Created by Japnik Singh on 5/26/10.

- Copyright 2010 **MyCompanyName**. All rights reserved.

- 

\*/

[/codebox]