Bug in NIR to PTX translation?

[yavuz650]

I was tinkering with the RTV shaders in Lumibench, and I ran into a strange issue. You can reproduce it by making the following minimal changes to the TraceAO.rgen shader,

    vec3 pixelColor = vec3(0);
+   uint dummy = gl_LaunchIDEXT.x == 42 ? 10 : 20; // Line 29
    // Accumulate all the rays for this pixels. (SPP)
    for (uint s = 0; s < Camera.NumberOfSamples; ++s)
    {

......

    //  debugPrintfEXT("[%d, %d] rgba(%f, %f, %f)\n", gl_LaunchIDEXT.x, gl_LaunchIDEXT.y, pixelColor.x, pixelColor.y, pixelColor.z);
+   imageStore(OutputImage, ivec2(dummy, gl_LaunchIDEXT.y), vec4(pixelColor, 0)); // Line 104
}

When you compile this shader and run it on vulkan-sim, it throws the following error while generating PTXINFO,

gpgpusim: invoking gpgpusim
GPGPU-Sim PTX: instruction assembly for function 'MESA_SHADER_RAYGEN_func0_main'...   done.
GPGPU-Sim PTX: finished parsing EMBEDDED .ptx file /home/root/vulkan-sim-root/vulkan-sim/../mesa-vulkan-sim/gpgpusimShaders/MESA_SHADER_RAYGEN_0.ptx
GPGPUSIM: Generating PTXINFO for/home/root/vulkan-sim-root/vulkan-sim/../mesa-vulkan-sim/gpgpusimShaders/MESA_SHADER_RAYGEN_0.ptxinfo
ptxas temp_ptxas_shader.ptx, line 3320; error   : Arguments mismatch for instruction 'mov'
ptxas fatal   : Ptx assembly aborted due to errors

I dug into the PTX code(temp_ptxas_shader.ptx) and found that the error is caused by a mov.f32 instruction whose argument is a u32 instead of f32

.reg .u32 %ssa_27_1;

mov.f32 %ssa_802_1, %ssa_27_1;

So the source operand appears to be u32 (which I believe is the dummy variable we declared above), but I'm not sure why ptxas considers this illegal.

System Info

• Cuda 11.7
• Ubuntu 20.04

[lucylufei]

Hi! Thanks for bringing this to our attention. There are a few assumptions that we make in the PTX translator that breaks with more custom shaders. We are working to improve the translator to handle these kind of cases. In the meantime, you can try to manually adjust the variable types in the PTX file to pass through ptxinfo. You should be able to skip the PTX translation stage (to avoid overwriting your modified PTX files) by disabling the nir_translate_shader_to_ptx function and the run_rt_translation_passes function.