[windows] Fix assembler error in pal trap handler.

[stellaraccident]

This has been showing up in real use as an error printed to the console complaining that the ".not_s_trap" label cannot be found on device initialization.

Tracked back to this commit: 7b72c1b

In the referenced original source, the .not_s_trap label contained conditional code for gfx94x which was removed. Outside of that case, it jumps to .no_skip_debug_trap, which we use here.

Note that I think this was broken since the above commit was landed. However, there does not seem to be error handling around the comgr call to assemble and install the trap handler, so this only manifests as an LLVM error printed to the console. I'm not sure what testing is done on this, but this means that if it was just an unattended pipeline, the bug may have been missed. Recommend fixing error handling to fail on error.

Tested by running hipblaslt tests on Windows/gfx1151 in TheRock and verifying that the error was not printed and tests run correctly.

[amd-justchen]

Explanation makes sense to me

[lancesix]

Hi, thanks for catching that. I thought I had a similar patch somewhere in the pipeline to fix this, but this one is equivalent.

FWIW, this LGTM.

[stellaraccident]

Do you all have a special procedure for landing patches? I'm getting "Cannot change this locked branch".

@jayhawk-commits

The base branch was changed.

[stellaraccident]

Hi, thanks for catching that. I thought I had a similar patch somewhere in the pipeline to fix this, but this one is equivalent.

FWIW, this LGTM.

Indeed - looks fixed on develop. I got confused and thought this was one of the oddball repos that was developing on a staging branch. Anyway, closing.

[lamb-j]

@lancesix does https://github.com/AMD-ROCm-Internal/clr/commit/c35e9643ecaba41a69da3fc4ac01203586454176 fix this issue?

It looks like we're still using .not_s_trap instead of .no_skip_debug_trap

[lancesix]

It should fix the assembling issue, doesn't it? .not_s_trap and .no_skip_debug_trap do refer to the same point in the code. I just used another label to make it slightly clearer. Some pseudo code for what this does is:

if (ttmp1.trap_id != 0)
  {
    // we entered the trap handler because we executed `s_trap X`, X being available in ttmp1
    ttmp_0_1.PC += 4; // the wave's saved PC points to the s_trap instruction itself, advance past it.

    if (ttmp1.trap_id == debug_trap && !ttmp11.debug_enabled)
      {
        // this is __builtin_debugtrap (), but the debugger is not attached, just skip this trap.
       // Logically return from the trap handler, we need to jump to .exit to do that.
       return;
      }
    .no_skip_debug_trap
  }
.not_s_trap

// continue processing the trap

[lamb-j]

Great, thanks for clarifying @lancesix!