Symbolicated dissassemly

[stolk]

I have followed this guide to get the disassembled OpenCL kernel.

I retrieved the .isabin file from my Intel AlderLake-S GT1 using:

$ cliloader --dump-output-binaries --dump-kernel-isa-binaries ./viewer fs=0

And then I disassembled with:

$ iga64 -d -p 12p1 CLI_0000_27223D7A_0000_00000000_GPU_gib_raytest.isabin

I get the disassembled code, which is great. However, I find it hard to read.

There are no variable names, no source-code line-numbers.

How can I get a disassembled OpenCL kernel where the symbol names or at least the line numbers are shown?

Update

I added -g as an option to clBuildProgram() but that does not seem to have made a difference? The resulting .bin dump is a lot larger with that option, but the .isabin files are roughly the same size.

[bashbaug]

Hello! Yes, the query used to obtain the "isabin" only contains the kernel ISA itself and does not include any debugging information. I think there's a way to get interleaved OpenCL C source and disassembled kernel ISA by passing the right program build options and then disassembling with the "ocloc" offline compiler, but I can't remember what the right build options so my experiments aren't working.

I'm asking around to see if there is a way to do this. In the meantime, here's a solution you can try. It's a little klunky but I've verified that it works:

1. Set the Intel Graphics Compiler (IGC) ShaderDumpEnable environment variable to enable IGC "shader dumps". See here for more information.
2. Pass the following build options when you build your OpenCL programs:
   • -g to enable debug information.
   • -s /full/path/to/your/kernel.cl to associate the debug information with your kernel source.
3. Run your program with the ShaderDumpEnable environment variable set.
4. Look for the interleaved kernel source and asm in the /tmp/IntelIGC directory. If you compile several kernels you should be able to grep the asm files for your kernel name to find the right file.

Example:

// File: /home/bashbaug/git/SimpleOpenCLSamples/install/Release/sample_kernel.cl

// Line 3:  uint index = get_global_id(0);
(W)     mul (1|M0)               acc0.0<1>:d   r8.2<0;1,0>:d     r3.2<0;1,0>:uw   {@1}               // $8

Give this a try and let me know if it works for you - thanks!

[stolk]

Thank you.

This works as described on Intel OpenCL. AMD OpenCL get's confused when passing '-g' or any option, really. But I guess that's not surprising. Thank you for the solution.