Databases
StackOverflow2013
Postsdbo
POCUDA DE kernel not launching
 

Note that there are some explanatory texts on larger screens.

plurals
  1. POCUDA DE kernel not launching
    primarykey
    data
    text
    <p>I'm trying to do differential evolution on CUDA, but the problem is that kernel which is responsible for "Mutation, Crossover, Evaluation, Selection" never gets launched.</p> <p>Any help?</p> <p>Here's the entire code:</p> <pre><code>#include &lt;iostream&gt; #include &lt;curand_kernel.h&gt; using namespace std; /**** ERROR HANDLING ****/ static void HandleError(cudaError_t err,const char *file, int line ) { if (err != cudaSuccess) { printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line ); system("pause"); exit( EXIT_FAILURE ); } } #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ )) /**** HOST AND DEVICE CONSTANTS****/ const int hNP=100, hD=31, hN=10; __constant__ int NP, D, N; __constant__ float Cr, F; /*** EVAL FUNCTION******/ __device__ float lennardJones(float a[3], float b[3]) { float distance = sqrt((a[0] - b[0]) * (a[0] - b[0]) + (a[1] - b[1]) * (a[1] - b[1]) + (a[2] - b[2]) * (a[2] - b[2])); float distance6 = distance * distance * distance * distance * distance * distance; float distance12 = distance6 * distance6; return 1/distance12 - 2/distance6; } /**** RANDOM GENERATORS***/ __device__ float rndFloat(curandState* globalState, int id) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM; } __device__ int rndInt(curandState* globalState, int id, int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM*max; } __device__ float rndFloat(curandState* globalState, int id, int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM*max; } __device__ float rndFloat(curandState* globalState, int id, int min,int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return min+RANDOM*(max-min); } /*** SEEDS ****/ __global__ void setup_kernel (curandState * state, unsigned long seed) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id &lt; NP) curand_init(seed, id, 0,&amp;state[id]); } /**** DIFFERENTIAL EVOLUTION: INITIALIZATION ***/ __global__ void kernelE(curandState* globalState, float *population) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id &lt; NP) { //init, just populating array with some specific numbers population[D*id]=0; population[D*id+N]=0; population[D*id +2*N]=0; population[D*id+1]=rndFloat(globalState,threadIdx.x,4); population[D*id+N+1]=0; population[D*id +2*N+1]=0; for(int i=2; i&lt;N; i++){ float min= -4 - 1/4*abs((int)((i-4)/3)); float max= 4 + 1/4*abs((int)((i-4)/3)); if(i==2) { population[D*id+2]=rndFloat(globalState,threadIdx.x,3.14159265359); population[D*id+N+2]=rndFloat(globalState,threadIdx.x,min,max); population[D*id +2*N+2]=0; } else { population[D*id +i]=rndFloat(globalState,threadIdx.x,min,max); population[D*id+N+i]=rndFloat(globalState,threadIdx.x,min,max); population[D*id +2*N+i]=rndFloat(globalState,threadIdx.x,min,max); } } //eval float e=0; for(int i=0; i&lt;N; i++) { for(int j=0; j&lt;i; j++) { float a[]={population[D*id +i], population[D*id+N+i], population[D*id +2*N+i]}, b[]={population[D*id +j],population[D*id +j+N], population[D*id +2*N+j]}; e += lennardJones(a,b); } } population[D*id + D-1]=e; } } /**** DIFFERENTIAL EVOLUTION: MUTATION INDICES ****/ __global__ void kernelP(curandState* globalState, int *mutation) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id&lt;NP) { int a = rndInt(globalState, id, NP),b = rndInt(globalState, id, NP),c= rndInt(globalState, id, NP); while(a == id){a = rndInt(globalState, id, NP);} while(b == a &amp;&amp; b==id){b=rndInt(globalState, id, NP);} while(c == a &amp;&amp; c== b &amp;&amp; c ==id){c=rndInt(globalState, id, NP);} mutation[D*id+0]=a; mutation[D*id+1]=b; mutation[D*id+2]=c; } } /**** DIFFERENTIAL EVOLUTION: MUTATION, CROSSOVER, EVALUATION AND SELECTION ***/ __global__ void kernelMCER(curandState* globalState, float *population, int *mutation, float *pop) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id&lt;NP) { int a=mutation[D*id+0], b=mutation[D*id+1], c=mutation[D*id+2]; //DE mutation and crossover int j=rndInt(globalState, id, NP); for(int i=0; i&lt;D-1; i++) { //DE mutation pop[D*id+i]= population[D*a +i] + F*(population[D*b +i]-population[D*c +i]); //DE crossover if(Cr &gt; rndFloat(globalState, id) &amp;&amp; i!= j) pop[D*id+i]=population[D*id +i]; } // Eval pop[D*id+D-1]=0; for(int i=0; i&lt;N; i++) { for(int j=0; j&lt;i; j++) { float a[]={pop[D*id+i], pop[D*id+N+i], pop[D*id+2*N+i]}, b[]={pop[D*id+j],pop[D*id+N+j], pop[D*id+2*N+j]}; pop[D*id+D-1] += lennardJones(a,b); } } __syncthreads(); //DE selection if(pop[D*id+D-1] &lt; population[D*id +D-1]) { for(int i=0; i&lt;D; i++) population[D*id +i]=pop[D*id+i]; } } } void getBestScore(float *hpopulation) { int max=0; for(int i=1; i&lt;hNP; i++) { if(hpopulation[hD*max+hD-1] &gt; hpopulation[hD*i+hD-1]) max=i; } for(int j=0; j&lt;hN; j++) cout&lt;&lt;"Atom "&lt;&lt;(j+1)&lt;&lt;": ("&lt;&lt;hpopulation[hD*max+j]&lt;&lt;", "&lt;&lt;hpopulation[hD*max+hN+j]&lt;&lt;", "&lt;&lt;hpopulation[hD*max+hN*2+j]&lt;&lt;") "&lt;&lt;endl; cout&lt;&lt;"Result: "&lt;&lt;hpopulation[hD*max+hD-1]&lt;&lt;endl; } int main() { cudaEvent_t start,stop; HANDLE_ERROR(cudaEventCreate(&amp;start)); HANDLE_ERROR(cudaEventCreate(&amp;stop)); HANDLE_ERROR(cudaEventRecord(start,0)); int device, st=100; float hCr=0.6f, hF=0.8f; cudaDeviceProp prop; HANDLE_ERROR(cudaGetDevice(&amp;device)); HANDLE_ERROR(cudaGetDeviceProperties(&amp;prop, device)); // int SN = prop.maxThreadsPerBlock; //512 threads per block //int SB = (hNP+(SN-1))/SN; //constants NP, D, N, Cr, F HANDLE_ERROR(cudaMemcpyToSymbol(N, &amp;hN, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(NP, &amp;hNP, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(D, &amp;hD, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(F, &amp;hF, sizeof(float))); HANDLE_ERROR(cudaMemcpyToSymbol(Cr, &amp;hCr, sizeof(float))); //seeds curandState* devStates; HANDLE_ERROR(cudaMalloc (&amp;devStates, hNP*sizeof(curandState))); setup_kernel &lt;&lt;&lt; 1, hNP&gt;&gt;&gt; (devStates, 50); //population float *population, *pop; float hpopulation[hNP*hD]; HANDLE_ERROR(cudaMalloc((void**)&amp;population, hNP*hD*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&amp;pop, hNP*hD*sizeof(float))); //mutation int *mutation, *mutation1; int *hmutation; HANDLE_ERROR(cudaHostAlloc((void**)&amp;hmutation, hNP*3*sizeof(int), cudaHostAllocDefault)); HANDLE_ERROR(cudaMalloc((void**)&amp;mutation, hNP*3*sizeof(int))); HANDLE_ERROR(cudaMalloc((void**)&amp;mutation1, hNP*3*sizeof(int))); //stream cudaStream_t stream_i, stream_j; HANDLE_ERROR(cudaStreamCreate(&amp;stream_i)); HANDLE_ERROR(cudaStreamCreate(&amp;stream_j)); kernelE&lt;&lt;&lt;1,hNP, 0,stream_i&gt;&gt;&gt;(devStates,population); kernelP&lt;&lt;&lt;1,hNP, 0,stream_j&gt;&gt;&gt;(devStates,mutation); while(st != 0) { /*** COPYING MUTATION INDICES***/ HANDLE_ERROR(cudaMemcpyAsync(hmutation, mutation,hNP*3*sizeof(int), cudaMemcpyDeviceToHost, stream_j)); HANDLE_ERROR(cudaMemcpyAsync(mutation1, hmutation,hNP*3*sizeof(int), cudaMemcpyHostToDevice, stream_i)); /**** CALLING KERNELS****/ kernelP&lt;&lt;&lt;1,hNP,0,stream_j&gt;&gt;&gt;(devStates,mutation); kernelMCER&lt;&lt;&lt;1,hNP,0,stream_i&gt;&gt;&gt;(devStates,population,mutation1,pop); st--; //HANDLE_ERROR(cudaStreamSynchronize(stream_i)); //HANDLE_ERROR(cudaMemcpy(hpopulation, population, hNP*hD*sizeof(float), cudaMemcpyDeviceToHost)); //getBestScore(hpopulation); //cin.get(); } HANDLE_ERROR(cudaStreamSynchronize(stream_i)); HANDLE_ERROR(cudaMemcpy(hpopulation, population, hNP*hD*sizeof(float), cudaMemcpyDeviceToHost)); getBestScore(hpopulation); cudaEventRecord(stop,0); cudaEventSynchronize(stop); float time; HANDLE_ERROR(cudaEventElapsedTime(&amp;time, start, stop)); cout&lt;&lt;endl&lt;&lt;"Tme: "&lt;&lt;time/1000&lt;&lt;"s"&lt;&lt;endl; HANDLE_ERROR(cudaEventDestroy(start)); HANDLE_ERROR(cudaEventDestroy(stop)); HANDLE_ERROR(cudaStreamDestroy(stream_i)); HANDLE_ERROR(cudaStreamDestroy(stream_j)); HANDLE_ERROR(cudaFree(population)); HANDLE_ERROR(cudaFree(pop)); HANDLE_ERROR(cudaFreeHost(hmutation)); HANDLE_ERROR(cudaFree(mutation1)); HANDLE_ERROR(cudaFree(devStates)); system("pause"); return 0; } </code></pre> <p>UPDATE - Solution:</p> <pre><code>#include &lt;iostream&gt; #include &lt;curand_kernel.h&gt; using namespace std; /**** ERROR HANDLING ****/ static void HandleError(cudaError_t err,const char *file, int line ) { if (err != cudaSuccess) { printf( "%s in %s at line %d\n", cudaGetErrorString( err ), file, line ); system("pause"); exit( EXIT_FAILURE ); } } #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ )) /**** HOST AND DEVICE CONSTANTS****/ const int hNP=100, hD=31, hN=10; __constant__ int NP, D, N; __constant__ float Cr, F; /*** EVAL FUNCTION******/ __device__ float lennardJones(float a[3], float b[3]) { float distance = sqrt((a[0] - b[0]) * (a[0] - b[0]) + (a[1] - b[1]) * (a[1] - b[1]) + (a[2] - b[2]) * (a[2] - b[2])); float distance6 = distance * distance * distance * distance * distance * distance; float distance12 = distance6 * distance6; return 1/distance12 - 2/distance6; } /**** RANDOM GENERATORS***/ __device__ float rndFloat(curandState* globalState, int id) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM; } __device__ int rndInt(curandState* globalState, int id, int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM*max; } __device__ float rndFloat(curandState* globalState, int id, int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return RANDOM*max; } __device__ float rndFloat(curandState* globalState, int id, int min,int max) { curandState localState = globalState[id]; float RANDOM = curand_uniform(&amp;localState); globalState[id] = localState; return min+RANDOM*(max-min); } /*** SEEDS ****/ __global__ void setup_kernel (curandState * state, unsigned long seed) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id &lt; NP) curand_init(seed, id, 0,&amp;state[id]); } /**** DIFFERENTIAL EVOLUTION: INITIALIZATION ***/ __global__ void kernelE(curandState* globalState, float *population) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id &lt; NP) { //init, just populating array with some specific numbers population[D*id]=0; population[D*id+N]=0; population[D*id +2*N]=0; population[D*id+1]=rndFloat(globalState,threadIdx.x,4); population[D*id+N+1]=0; population[D*id +2*N+1]=0; for(int i=2; i&lt;N; i++){ float min= -4 - 1/4*abs((int)((i-4)/3)); float max= 4 + 1/4*abs((int)((i-4)/3)); if(i==2) { population[D*id+2]=rndFloat(globalState,threadIdx.x,3.14159265359); population[D*id+N+2]=rndFloat(globalState,threadIdx.x,min,max); population[D*id +2*N+2]=0; } else { population[D*id +i]=rndFloat(globalState,threadIdx.x,min,max); population[D*id+N+i]=rndFloat(globalState,threadIdx.x,min,max); population[D*id +2*N+i]=rndFloat(globalState,threadIdx.x,min,max); } } //eval float e=0; for(int i=0; i&lt;N; i++) { for(int j=0; j&lt;i; j++) { float a[]={population[D*id +i], population[D*id+N+i], population[D*id +2*N+i]}, b[]={population[D*id +j],population[D*id +j+N], population[D*id +2*N+j]}; e += lennardJones(a,b); } } population[D*id + D-1]=e; } } /**** DIFFERENTIAL EVOLUTION: MUTATION INDICES ****/ __global__ void kernelP(curandState* globalState, int *mutation) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id&lt;NP) { int a = rndInt(globalState, id, NP),b = rndInt(globalState, id, NP),c= rndInt(globalState, id, NP); while(a == id){a = rndInt(globalState, id, NP);} while(b == a &amp;&amp; b==id){b=rndInt(globalState, id, NP);} while(c == a &amp;&amp; c== b &amp;&amp; c ==id){c=rndInt(globalState, id, NP);} mutation[3*id+0]=a; mutation[3*id+1]=b; mutation[3*id+2]=c; } } /**** DIFFERENTIAL EVOLUTION: MUTATION, CROSSOVER, EVALUATION AND SELECTION ***/ __global__ void kernelMCER(curandState* globalState, float *population, int *mutation, float *pop) { int id= threadIdx.x+blockIdx.x*blockDim.x; if(id&lt;NP) { int a=mutation[3*id+0], b=mutation[3*id+1], c=mutation[3*id+2]; //DE mutation and crossover int j=rndInt(globalState, id, NP); for(int i=0; i&lt;D-1; i++) { //DE mutation pop[D*id+i]= population[D*a +i] + F*(population[D*b +i]-population[D*c +i]); //DE crossover if(Cr &gt; rndFloat(globalState, id) &amp;&amp; i!= j) pop[D*id+i]=population[D*id +i]; } // Eval pop[D*id+D-1]=0; for(int i=0; i&lt;N; i++) { for(int j=0; j&lt;i; j++) { float a[]={pop[D*id+i], pop[D*id+N+i], pop[D*id+2*N+i]}, b[]={pop[D*id+j],pop[D*id+N+j], pop[D*id+2*N+j]}; pop[D*id+D-1] += lennardJones(a,b); } } __syncthreads(); //DE selection if(pop[D*id+D-1] &lt; population[D*id +D-1]) { for(int i=0; i&lt;D; i++) population[D*id +i]=pop[D*id+i]; } } } void getBestScore(float *hpopulation) { int max=0; for(int i=1; i&lt;hNP; i++) { if(hpopulation[hD*max+hD-1] &gt; hpopulation[hD*i+hD-1]) max=i; } for(int j=0; j&lt;hN; j++) cout&lt;&lt;"Atom "&lt;&lt;(j+1)&lt;&lt;": ("&lt;&lt;hpopulation[hD*max+j]&lt;&lt;", "&lt;&lt;hpopulation[hD*max+hN+j]&lt;&lt;", "&lt;&lt;hpopulation[hD*max+hN*2+j]&lt;&lt;") "&lt;&lt;endl; cout&lt;&lt;"Result: "&lt;&lt;hpopulation[hD*max+hD-1]&lt;&lt;endl; } int main() { cudaEvent_t start,stop; HANDLE_ERROR(cudaEventCreate(&amp;start)); HANDLE_ERROR(cudaEventCreate(&amp;stop)); HANDLE_ERROR(cudaEventRecord(start,0)); int device, st=100; float hCr=0.6f, hF=0.8f; cudaDeviceProp prop; HANDLE_ERROR(cudaGetDevice(&amp;device)); HANDLE_ERROR(cudaGetDeviceProperties(&amp;prop, device)); // int SN = prop.maxThreadsPerBlock; //512 threads per block //int SB = (hNP+(SN-1))/SN; //constants NP, D, N, Cr, F HANDLE_ERROR(cudaMemcpyToSymbol(N, &amp;hN, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(NP, &amp;hNP, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(D, &amp;hD, sizeof(int))); HANDLE_ERROR(cudaMemcpyToSymbol(F, &amp;hF, sizeof(float))); HANDLE_ERROR(cudaMemcpyToSymbol(Cr, &amp;hCr, sizeof(float))); //seeds curandState* devStates; HANDLE_ERROR(cudaMalloc (&amp;devStates, hNP*sizeof(curandState))); setup_kernel &lt;&lt;&lt; 1, hNP&gt;&gt;&gt; (devStates, 50); //population float *population, *pop; float hpopulation[hNP*hD]; HANDLE_ERROR(cudaMalloc((void**)&amp;population, hNP*hD*sizeof(float))); HANDLE_ERROR(cudaMalloc((void**)&amp;pop, hNP*hD*sizeof(float))); //mutation int *mutation, *mutation1; int *hmutation; HANDLE_ERROR(cudaHostAlloc((void**)&amp;hmutation, hNP*3*sizeof(int), cudaHostAllocDefault)); HANDLE_ERROR(cudaMalloc((void**)&amp;mutation, hNP*3*sizeof(int))); HANDLE_ERROR(cudaMalloc((void**)&amp;mutation1, hNP*3*sizeof(int))); //stream cudaStream_t stream_i, stream_j; HANDLE_ERROR(cudaStreamCreate(&amp;stream_i)); HANDLE_ERROR(cudaStreamCreate(&amp;stream_j)); kernelE&lt;&lt;&lt;1,hNP, 0,stream_i&gt;&gt;&gt;(devStates,population); kernelP&lt;&lt;&lt;1,hNP, 0,stream_j&gt;&gt;&gt;(devStates,mutation); while(st != 0) { /*** COPYING MUTATION INDICES***/ HANDLE_ERROR(cudaMemcpyAsync(hmutation, mutation,hNP*3*sizeof(int), cudaMemcpyDeviceToHost, stream_j)); HANDLE_ERROR(cudaMemcpyAsync(mutation1, hmutation,hNP*3*sizeof(int), cudaMemcpyHostToDevice, stream_i)); /**** CALLING KERNELS****/ kernelP&lt;&lt;&lt;1,hNP,0,stream_j&gt;&gt;&gt;(devStates,mutation); kernelMCER&lt;&lt;&lt;1,hNP,0,stream_i&gt;&gt;&gt;(devStates,population,mutation1,pop); st--; //HANDLE_ERROR(cudaStreamSynchronize(stream_i)); //HANDLE_ERROR(cudaMemcpy(hpopulation, population, hNP*hD*sizeof(float), cudaMemcpyDeviceToHost)); //getBestScore(hpopulation); //cin.get(); } HANDLE_ERROR(cudaStreamSynchronize(stream_i)); HANDLE_ERROR(cudaMemcpy(hpopulation, population, hNP*hD*sizeof(float), cudaMemcpyDeviceToHost)); getBestScore(hpopulation); cudaEventRecord(stop,0); cudaEventSynchronize(stop); float time; HANDLE_ERROR(cudaEventElapsedTime(&amp;time, start, stop)); cout&lt;&lt;endl&lt;&lt;"Tme: "&lt;&lt;time/1000&lt;&lt;"s"&lt;&lt;endl; HANDLE_ERROR(cudaEventDestroy(start)); HANDLE_ERROR(cudaEventDestroy(stop)); HANDLE_ERROR(cudaStreamDestroy(stream_i)); HANDLE_ERROR(cudaStreamDestroy(stream_j)); HANDLE_ERROR(cudaFree(population)); HANDLE_ERROR(cudaFree(pop)); HANDLE_ERROR(cudaFreeHost(hmutation)); HANDLE_ERROR(cudaFree(mutation1)); HANDLE_ERROR(cudaFree(devStates)); system("pause"); return 0; } </code></pre>
    singulars
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PTQuestion
    1. USuser755974
    1. USuser755974
    plurals
    1. PL
      singulars
      1. LTLinked
    1. This table or related slice is empty.
    1. This table or related slice is empty.
    1. PO
      singulars
      1. PTAnswer
    1. VO
      singulars
      1. POCUDA DE kernel not launching
      1. This table or related slice is empty.
      1. VTClose
    2. VO
      singulars
      1. POCUDA DE kernel not launching
      1. This table or related slice is empty.
      1. VTClose
    1. COYou can do [error checking on kernel launches](http://stackoverflow.com/questions/14038589/what-is-the-canonical-way-to-check-for-errors-using-the-cuda-runtime-api) too. Why don't you do that and report back what kind of error you get? Also, what happens when you run your code with `cuda-memcheck`? My guess is your kernel is launching but you have some sort of error that occurs in device code, causing the kernel to terminate early.
      singulars
      1. POCUDA DE kernel not launching
      1. USRobert Crovella
    2. COerror checking on kernel shows no errors. cuda-memcheck also show no error. I also installed latest nsight with appropriate driver on VS 2010 but when I press "Start CUDA Debugging" nothing happens... breakpoints that I have placed inside kernel are ignored. I have no idea what to do next.
      singulars
      1. POCUDA DE kernel not launching
      1. USuser755974
    3. COUpdate the code you've posted to show how you are doing kernel error checking. Even better would be to post a complete, compilable code. Unless you have a logic error in your host code outside of what you have posted here, or are completely misinterpreting the data produced by your code, it's not possible for a kernel to not launch, throw no errors, and show no errors in cuda-memcheck.
      singulars
      1. POCUDA DE kernel not launching
      1. USRobert Crovella
 

Querying!

 
Guidance
A row detail

Detail views are divided into sections. All the information in the data section comes from columns in the selected row. The other sections display data from other, related rows.

Related data can be related in a to-one or a to-many fashion. Captions of data related in a to-many fashion link to a list view showing a filtered view of the table.

Try moving around until you find a non-empty to-many entry and click on the label to get to one. You can move back to the root by clicking on the database name in the header.

SQuiL has stopped working due to an internal error.

If you are curious you may find further information in the browser console, which is accessible through the devtools (F12).

Reload