NaN

[jon-hotaisle]

Just doing a bit of debugging.

"val loss" output nan, so I figured start there...

val loss 1 nan
val loss 2 nan

But digging higher up, val_num_batches is set to 20, so I'm not sure how this is turning into nan so easily. Feels like something else is up...

[jon-hotaisle]

@anthonix bump.

[anthonix]

Will try and reproduce -- on the list of things to do when I have some spare cycles

[jon-hotaisle]

Oh, I'm blind (and probably dumb). val_loss must be 0, hence the nan. so it must be something in that gpt2_validate() returning all zeros.

[anthonix]

In the mean time, can you verify some other training works, like AMD's tinyllama code they recently released? Or their JAX GPT2 training?

[phani544]

I think it is happening here
779 cudaCheck(hipMemcpy(model->cpu_losses, acts.losses, B * T * sizeof(float), hipMemcpyDeviceToHost));

gdb trace

Thread 1 "train_gpt2amd" hit Breakpoint 1, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:779
779 cudaCheck(hipMemcpy(model->cpu_losses, acts.losses, B * T * sizeof(float), hipMemcpyDeviceToHost));
(gdb) print acts.losses
$1 = (float ) 0x7fbf652c8000
(gdb) n
780 for (int i = 0; i < BT; i++) {
(gdb) print model->cpu_losses
$2 = (float *) 0x7fffebe40000
(gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781
781 mean_loss += model->cpu_losses[i];
(gdb) print mean_loss
$3 = 0
(gdb) n
780 for (int i = 0; i < B*T; i++) {
(gdb) print i
$4 = 0
(gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781
781 mean_loss += model->cpu_losses[i];
(gdb) print model->cpu_losses[i]
$5 = -nan(0x7f0000)
(gdb) n
780 for (int i = 0; i < B*T; i++) {
(gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781
781 mean_loss += model->cpu_losses[i];
(gdb) print i
$6 = 2
(gdb) print model->cpu_losses[0]
$7 = -nan(0x7f0000)
(gdb) print mean_loss
$8 = -nan(0x7f0000)
(gdb)

[mszpindler]

I think it is happening here 779 cudaCheck(hipMemcpy(model->cpu_losses, acts.losses, B * T * sizeof(float), hipMemcpyDeviceToHost));

gdb trace

Thread 1 "train_gpt2amd" hit Breakpoint 1, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:779 779 cudaCheck(hipMemcpy(model->cpu_losses, acts.losses, B * T * sizeof(float), hipMemcpyDeviceToHost)); (gdb) print acts.losses $1 = (float _) 0x7fbf652c8000 (gdb) n 780 for (int i = 0; i < B_T; i++) { (gdb) print model->cpu_losses $2 = (float *) 0x7fffebe40000 (gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781 781 mean_loss += model->cpu_losses[i]; (gdb) print mean_loss $3 = 0 (gdb) n 780 for (int i = 0; i < B*T; i++) { (gdb) print i $4 = 0 (gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781 781 mean_loss += model->cpu_losses[i]; (gdb) print model->cpu_losses[i] $5 = -nan(0x7f0000) (gdb) n 780 for (int i = 0; i < B*T; i++) { (gdb) n

Thread 1 "train_gpt2amd" hit Breakpoint 2, gpt2_validate (model=0x7fffffffd8f0, inputs=0xbb0d70, targets=0xa63930, B=4, T=1024) at build/hip/train_gpt2.cu:781 781 mean_loss += model->cpu_losses[i]; (gdb) print i $6 = 2 (gdb) print model->cpu_losses[0] $7 = -nan(0x7f0000) (gdb) print mean_loss $8 = -nan(0x7f0000) (gdb)

I believe this issue can be addressed by careful use of hip_bfloat16 and __bfloat16 types which hipify translates not fully properly from cuda code (at least this is the case for rocm/6.2.2).