[CUDA ERROR] at file \llm.c\train_gpt2.cu:405: too many resources requested for launch (old version does not have this issue - fyi)

[karpathy]

[CUDA ERROR] at file \llm.c\train_gpt2.cu:405: too many resources requested for launch (old version does not have this issue - fyi)

[msharmavikram]

Change the following lines

__global__  void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,
                                    const float*  __restrict__ inp, const float*  __restrict__ weight,
                                    const float* __restrict__ bias, int N, int C) {

to below

__global__ __launch_bounds__(1024,2) 
void layernorm_forward_kernel3(float* __restrict__ out, float* __restrict__ mean, float* __restrict__ rstd,
                                    const float*  __restrict__ inp, const float*  __restrict__ weight,
                                    const float* __restrict__ bias, int N, int C) {

I believe it is running out of registers. So need to forcefully do register spills. This solved the error for me.

[ross-wheeler]

Thank you, Vikram @msharmavikram ! That fixed it! Should this be a generic fix or is it card/gpu specific?

[msharmavikram]

This may not work on all cards. I need to determine the correctness of this fix for different thread block size before submitting a PR. I didn't a chance to work on that yet.

[msharmavikram]

Submitted a PR #126

[msharmavikram]

@ross-wheeler I do not see this issue anymore on the top of the tree. Let me know if you still see the error without the proposed fix.

[rosslwheeler]

@msharmavikram - I still need your above fix. The top of tree did not fix my issue.