CUDA unspecified launch failure error

Question

I have the following code http://pastebin.com/vLeD1GJm wich works just fine, but if I increase:

#define GPU_MAX_PW 100000000

to:

#define GPU_MAX_PW 1000000000

Then I receive:

frederico@zeus:~/Dropbox/coisas/projetos/delta_cuda$ optirun ./a
block size = 97657 grid 48828 grid 13951

unspecified launch failure in a.cu at line 447.. err number 4

I'm running this on a GTX 675M which has 2GB of memory. And the second definition of GPU_MAX_PW will have around 1000000000×2÷1024÷1024 = 1907 MB, so I'm not out of memory. What can be the problem since I'm only allocating more memory? Maybe the grid and block configuration become impossible?

Note that the error is pointing to this line:

HANDLE_ERROR(cudaMemcpy(gwords, gpuHashes, sizeof(unsigned short) * GPU_MAX_PW, cudaMemcpyDeviceToHost));

I added a code verification right after cudaThreadSynchronize(); and the error is there, not at copping the results — Frederico Schardong
If you run nvidia-smi -a (with nothing running on the GPU) how much free memory is reported? — Robert Crovella
Is your GTX 675M also displaying the user interface? How long is your kernel running with #define GPU_MAX_PW 1000000000 before the error comes up? You are probably running into a timeout as the OS reclaims the GPU for UI display. — tera
@tera No, I'm using my intel graphic card for display, the GTX 675M is only being used when I call it. The time was around 3 seconds — Frederico Schardong

Robert Crovella Robert Crovella · Accepted Answer · 2012-11-01T03:02:59

First of all you have your sizes listed incorrectly. The program works for 10,000,000 and not 100,000,000 (whereas you said it works for 100,000,000 and not 1,000,000,000). So memory size is not the issue, and your calculations there are based on the wrong numbers.

calculate_grid_parameters is messed up. The objective of this function is to figure out how many blocks are needed and therefore grid size, based on the GPU_MAX_PW specifying the total number of threads needed and 1024 threads per block (hard coded). The line that prints out block size = grid ... grid ... actually has the clue to the problem. For GPU_MAX_PW of 100,000,000, this function correctly computes that 100,000,000/1024 = 97657 blocks are needed. However, the grid dimensions are computed incorrectly. The grid dimensions grid.x * grid.y should equal the total number of blocks desired (approximately). But this function has decided that it wants grid.x of 48828 and grid.y of 13951. If I multiply those two, I get 681,199,428, which is much larger than the desired total block count of 97657. Now if I then launch a kernel with requested grid dimensions of 48828 (x) and 13951 (y), and also request 1024 threads per block, I have requested 697,548,214,272 total threads in that kernel launch. First of all this is not your intent, and secondly, while at the moment I can't say exactly why, this is apparently too many threads. Suffice it to say that this overall grid request exceeds some resource limitation of the machine.

Note that if you drop from 100,000,000 to 10,000,000 for GPU_MAX_PW, the grid calculation becomes "sensible", I get:

block size = 9766 grid 9766 grid 1

and no launch failure.

CUDA unspecified launch failure error

1 Answers