[A770] language/test_core.py::test_noinline[shared] failed

[pbchekin]

2024-05-09T19:34:10.8294846Z FAILED language/test_core.py::test_noinline[shared] - AssertionError: assert False
2024-05-09T19:34:10.8305399Z  +  where False = <built-in method equal of type object at 0x7f54f8a1b600>(tensor([[nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan, nan,\n         nan, nan],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.],\n        [99., 99., 99., 99., 99., 99., 99., 99., 99., nan, 99., 99., 99., nan,\n         99., 99.]], device='xpu:0'), ((tensor([[16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.],\n        [16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16., 16.,\n         16., 16.]], device='xpu:0') + tensor([1.], device='xpu:0')) + tensor([2.], device='xpu:0')))
2024-05-09T19:34:10.8313484Z  +    where <built-in method equal of type object at 0x7f54f8a1b600> = torch.equal

[anmyachev]

Using tl.device_print I see null pointers. @alexbaden any suggestions on where to look next?

@triton.jit(noinline=True)
def noinline_shared_fn(x, y, Z):
    offs = tl.arange(0, 16)[:, None] * 16 + tl.arange(0, 16)[None, :]
    z = tl.load(Z + offs)
    tl.device_print("Z1=", Z)  # <- output from
    z = tl.dot(z, z) + x + y
    tl.store(Z + offs, z)

Output:

pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)                                                                                                                              pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: (nil)
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
pid (0, 0, 0) idx () Z1=: 0xffffe38f1c000400
....

[anmyachev]

The test case works with threads_per_warp=8 ( kernel[(1, )](x, y, z, num_warps=1, threads_per_warp=8)).

[alexbaden]

Let's look at the IR from ~/.triton/cache - I think we need to understand more about the shared local memory capabilities of A770 and where it might be silently failing.

[anmyachev]

@alexbaden I was able to better localize the problem, with a simplified kernel the disassembled spirv code differs only in two lines:

               # doesn't work correctly
               OpExecutionMode %36 SubgroupSize 32
               OpExecutionMode %36 MaxWorkgroupSizeINTEL 32 1 1
vs
               OpExecutionMode %36 SubgroupSize 16
               OpExecutionMode %36 MaxWorkgroupSizeINTEL 16 1 1

It seems like the problem is at lower levels. Where should I try to report the problem? https://github.com/intel/intel-graphics-compiler or https://github.com/intel/compute-runtime?

Full working spirv dump

; SPIR-V
; Version: 1.3
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 59
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability Int8
               OpCapability SubgroupDispatch
               OpCapability KernelAttributesINTEL
               OpExtension "SPV_INTEL_kernel_attributes"
          %2 = OpExtInstImport "OpenCL.DebugInfo.100"
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %36 "kernel"
               OpExecutionMode %36 ContractionOff
               OpExecutionMode %36 SubgroupSize 16
               OpExecutionMode %36 MaxWorkgroupSizeINTEL 16 1 1
         %42 = OpString "/home/amyachev/intel-xpu-backend-for-triton/python/test/unit/language/test_core.py"
         %48 = OpString "_Z12get_local_idj"
         %50 = OpString "kernel"
         %53 = OpString "noinline_shared_fn__fp32_fp32_Pfp32__"
               OpSource OpenCL_CPP 100000
               OpName %_Z12get_local_idj "_Z12get_local_idj"
               OpName %kernel "kernel"
               OpName %noinline_shared_fn__fp32_fp32_Pfp32__ "noinline_shared_fn__fp32_fp32_Pfp32__"
               OpModuleProcessed "Debug info producer: triton"
               OpModuleProcessed "Debug info producer: triton"
               OpDecorate %_Z12get_local_idj LinkageAttributes "_Z12get_local_idj" Import
               OpDecorate %kernel LinkageAttributes "kernel" Export
               OpDecorate %15 FuncParamAttr NoCapture
               OpDecorate %15 FuncParamAttr NoWrite
               OpDecorate %16 FuncParamAttr NoCapture
               OpDecorate %16 FuncParamAttr NoWrite
               OpDecorate %17 FuncParamAttr NoCapture
               OpDecorate %21 FuncParamAttr NoCapture
               OpDecorate %37 FuncParamAttr NoCapture
               OpDecorate %37 FuncParamAttr NoWrite
               OpDecorate %38 FuncParamAttr NoCapture
               OpDecorate %38 FuncParamAttr NoWrite
               OpDecorate %39 FuncParamAttr NoCapture
      %ulong = OpTypeInt 64 0
       %uint = OpTypeInt 32 0
      %uchar = OpTypeInt 8 0
     %uint_0 = OpConstant %uint 0
%ulong_4294967295 = OpConstant %ulong 4294967295
    %ulong_0 = OpConstant %ulong 0
          %5 = OpTypeFunction %ulong %uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
      %float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
         %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_float
         %19 = OpTypeFunction %void %_ptr_CrossWorkgroup_float
       %bool = OpTypeBool
    %float_1 = OpConstant %float 1
         %43 = OpExtInst %void %2 DebugSource %42
         %44 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %43 Unknown
         %45 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %43 Unknown
         %46 = OpExtInst %void %2 DebugTypeFunction None %void
         %47 = OpExtInst %void %2 DebugTypeFunction None %void
         %49 = OpExtInst %void %2 DebugFunctionDeclaration %48 %46 %43 1359 0 %45 %48 FlagIsOptimized
         %51 = OpExtInst %void %2 DebugInfoNone
         %52 = OpExtInst %void %2 DebugFunction %50 %46 %43 1359 0 %44 %50 FlagIsDefinition|FlagIsOptimized 1359 %kernel %51
         %54 = OpExtInst %void %2 DebugFunction %53 %47 %43 1323 0 %45 %53 FlagIsDefinition|FlagIsOptimized 1323 %noinline_shared_fn__fp32_fp32_Pfp32__ %51
%_Z12get_local_idj = OpFunction %ulong None %5
          %7 = OpFunctionParameter %uint
               OpFunctionEnd
     %kernel = OpFunction %void None %13
         %15 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %16 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %17 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %18 = OpLabel
         %55 = OpExtInst %void %2 DebugScope %52
               OpLine %42 1362 29
         %22 = OpFunctionCall %void %noinline_shared_fn__fp32_fp32_Pfp32__ %17
               OpLine %42 1362 4
               OpReturn
               OpFunctionEnd
%noinline_shared_fn__fp32_fp32_Pfp32__ = OpFunction %void DontInline %19
         %21 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %23 = OpLabel
         %27 = OpFunctionCall %ulong %_Z12get_local_idj %uint_0
         %56 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1325 16
         %29 = OpBitwiseAnd %ulong %27 %ulong_4294967295
         %32 = OpIEqual %bool %29 %ulong_0
               OpBranchConditional %32 %24 %25
         %24 = OpLabel
         %57 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1324 16
         %33 = OpLoad %float %21 Aligned 4
               OpLine %42 1325 20
         %35 = OpFAdd %float %33 %float_1
               OpLine %42 1325 16
               OpStore %21 %35 Aligned 4
               OpBranch %25
         %25 = OpLabel
         %58 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1325 4
               OpReturn
               OpFunctionEnd
         %36 = OpFunction %void None %13
         %37 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %38 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %39 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %40 = OpLabel
         %41 = OpFunctionCall %void %kernel %37 %38 %39
               OpReturn
               OpFunctionEnd

Full non-working correctly spirv dump

; SPIR-V
; Version: 1.3
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 59
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability Int8
               OpCapability SubgroupDispatch
               OpCapability KernelAttributesINTEL
               OpExtension "SPV_INTEL_kernel_attributes"
          %2 = OpExtInstImport "OpenCL.DebugInfo.100"
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %36 "kernel"
               OpExecutionMode %36 ContractionOff
               OpExecutionMode %36 SubgroupSize 32
               OpExecutionMode %36 MaxWorkgroupSizeINTEL 32 1 1
         %42 = OpString "/home/amyachev/intel-xpu-backend-for-triton/python/test/unit/language/test_core.py"
         %48 = OpString "_Z12get_local_idj"
         %50 = OpString "kernel"
         %53 = OpString "noinline_shared_fn__fp32_fp32_Pfp32__"
               OpSource OpenCL_CPP 100000
               OpName %_Z12get_local_idj "_Z12get_local_idj"
               OpName %kernel "kernel"
               OpName %noinline_shared_fn__fp32_fp32_Pfp32__ "noinline_shared_fn__fp32_fp32_Pfp32__"
               OpModuleProcessed "Debug info producer: triton"
               OpModuleProcessed "Debug info producer: triton"
               OpDecorate %_Z12get_local_idj LinkageAttributes "_Z12get_local_idj" Import
               OpDecorate %kernel LinkageAttributes "kernel" Export
               OpDecorate %15 FuncParamAttr NoCapture
               OpDecorate %15 FuncParamAttr NoWrite
               OpDecorate %16 FuncParamAttr NoCapture
               OpDecorate %16 FuncParamAttr NoWrite
               OpDecorate %17 FuncParamAttr NoCapture
               OpDecorate %21 FuncParamAttr NoCapture
               OpDecorate %37 FuncParamAttr NoCapture
               OpDecorate %37 FuncParamAttr NoWrite
               OpDecorate %38 FuncParamAttr NoCapture
               OpDecorate %38 FuncParamAttr NoWrite
               OpDecorate %39 FuncParamAttr NoCapture
      %ulong = OpTypeInt 64 0
       %uint = OpTypeInt 32 0
      %uchar = OpTypeInt 8 0
     %uint_0 = OpConstant %uint 0
%ulong_4294967295 = OpConstant %ulong 4294967295
    %ulong_0 = OpConstant %ulong 0
          %5 = OpTypeFunction %ulong %uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
      %float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
         %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_uchar %_ptr_CrossWorkgroup_float
         %19 = OpTypeFunction %void %_ptr_CrossWorkgroup_float
       %bool = OpTypeBool
    %float_1 = OpConstant %float 1
         %43 = OpExtInst %void %2 DebugSource %42
         %44 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %43 Unknown
         %45 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %43 Unknown
         %46 = OpExtInst %void %2 DebugTypeFunction None %void
         %47 = OpExtInst %void %2 DebugTypeFunction None %void
         %49 = OpExtInst %void %2 DebugFunctionDeclaration %48 %46 %43 1359 0 %45 %48 FlagIsOptimized
         %51 = OpExtInst %void %2 DebugInfoNone
         %52 = OpExtInst %void %2 DebugFunction %50 %46 %43 1359 0 %44 %50 FlagIsDefinition|FlagIsOptimized 1359 %kernel %51
         %54 = OpExtInst %void %2 DebugFunction %53 %47 %43 1323 0 %45 %53 FlagIsDefinition|FlagIsOptimized 1323 %noinline_shared_fn__fp32_fp32_Pfp32__ %51
%_Z12get_local_idj = OpFunction %ulong None %5
          %7 = OpFunctionParameter %uint
               OpFunctionEnd
     %kernel = OpFunction %void None %13
         %15 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %16 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %17 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %18 = OpLabel
         %55 = OpExtInst %void %2 DebugScope %52
               OpLine %42 1362 29
         %22 = OpFunctionCall %void %noinline_shared_fn__fp32_fp32_Pfp32__ %17
               OpLine %42 1362 4
               OpReturn
               OpFunctionEnd
%noinline_shared_fn__fp32_fp32_Pfp32__ = OpFunction %void DontInline %19
         %21 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %23 = OpLabel
         %27 = OpFunctionCall %ulong %_Z12get_local_idj %uint_0
         %56 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1325 16
         %29 = OpBitwiseAnd %ulong %27 %ulong_4294967295
         %32 = OpIEqual %bool %29 %ulong_0
               OpBranchConditional %32 %24 %25
         %24 = OpLabel
         %57 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1324 16
         %33 = OpLoad %float %21 Aligned 4
               OpLine %42 1325 20
         %35 = OpFAdd %float %33 %float_1
               OpLine %42 1325 16
               OpStore %21 %35 Aligned 4
               OpBranch %25
         %25 = OpLabel
         %58 = OpExtInst %void %2 DebugScope %54
               OpLine %42 1325 4
               OpReturn
               OpFunctionEnd
         %36 = OpFunction %void None %13
         %37 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %38 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %39 = OpFunctionParameter %_ptr_CrossWorkgroup_float
         %40 = OpLabel
         %41 = OpFunctionCall %void %kernel %37 %38 %39
               OpReturn
               OpFunctionEnd

[anmyachev]

Reproducer without pytest:

import torch
import intel_extension_for_pytorch  # type: ignore # noqa: F401

import triton
import triton.language as tl



@triton.jit(noinline=True)
def noinline_shared_fn(Z):
    # this will print:
    #   * 64 lines in output with (nil) for threads_per_warp=32
    #     * however with noinline=False: 32 lines without (nil)
    #   * 16 lines in output without (nil) for threads_per_warp=16
    tl.device_print("address: ", Z)



def main():

    @triton.jit
    def kernel(Z):
        noinline_shared_fn(Z)

    z = torch.ones((16, 16), device="xpu:0", dtype=torch.float32)
    kernel[(1, )](z, num_warps=1, threads_per_warp=16)


if __name__ == "__main__":
    main()

Output when threads_per_warp=16:

pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000

Output when threads_per_warp=32:

pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: (nil)
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000
pid (0, 0, 0) idx () address: 0xffffe38f1c000000

[anmyachev]

@alexbaden more dumps for the previous reproducer:

kernel.llir noinline=false

; ModuleID = 'LLVMDialectModule'
source_filename = "LLVMDialectModule"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
target triple = "spirv64v1.3-unknown-unknown"

@printfFormat_0 = internal addrspace(2) constant [29 x i8] c"pid (%u, %u, %u) idx ()%s%p\0A\00"
@printfPrefix_0 = internal addrspace(2) constant [11 x i8] c" address: \00"

; Function Attrs: mustprogress nofree nosync nounwind willreturn memory(none)
declare !dbg !8 spir_func i64 @_Z12get_group_idj(i32) local_unnamed_addr #0

declare !dbg !11 spir_func i32 @_Z18__spirv_ocl_printf(ptr addrspace(2), ...) local_unnamed_addr

define spir_kernel void @kernel(ptr addrspace(1) %0) local_unnamed_addr !dbg !13 !intel_reqd_sub_group_size !14 !max_work_group_size !15 {
  %2 = tail call spir_func i64 @_Z12get_group_idj(i32 0)
  %3 = trunc i64 %2 to i32, !dbg !16
  %4 = tail call spir_func i64 @_Z12get_group_idj(i32 1)
  %5 = trunc i64 %4 to i32, !dbg !16
  %6 = tail call spir_func i64 @_Z12get_group_idj(i32 2)
  %7 = trunc i64 %6 to i32, !dbg !16
  %8 = tail call i32 (ptr addrspace(2), ...) @_Z18__spirv_ocl_printf(ptr addrspace(2) @printfFormat_0, i32 %3, i32 %5, i32 %7, ptr addrspace(2) @printfPrefix_0, ptr addrspace(1) %0)
  ret void, !dbg !19
}

attributes #0 = { mustprogress nofree nosync nounwind willreturn memory(none) }

!llvm.module.flags = !{!0, !1, !2}
!llvm.dbg.cu = !{!3}
!opencl.spir.version = !{!5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5, !5}
!spirv.Source = !{!6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6}
!llvm.ident = !{!7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 7, !"frame-pointer", i32 2}
!3 = distinct !DICompileUnit(language: DW_LANG_C, file: !4, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!4 = !DIFile(filename: "noinline_reproducer.py", directory: "/home/amyachev/intel-xpu-backend-for-triton")
!5 = !{i32 1, i32 2}
!6 = !{i32 4, i32 100000}
!7 = !{!"Intel(R) oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240216)"}
!8 = !DISubprogram(name: "_Z12get_group_idj", linkageName: "_Z12get_group_idj", scope: !4, file: !4, line: 22, type: !9, scopeLine: 22, spFlags: DISPFlagOptimized)
!9 = !DISubroutineType(cc: DW_CC_normal, types: !10)
!10 = !{}
!11 = !DISubprogram(name: "_Z18__spirv_ocl_printf", linkageName: "_Z18__spirv_ocl_printf", scope: !12, file: !12, type: !9, spFlags: DISPFlagOptimized)
!12 = !DIFile(filename: "<unknown>", directory: "")
!13 = distinct !DISubprogram(name: "kernel", linkageName: "kernel", scope: !4, file: !4, line: 22, type: !9, scopeLine: 22, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !3)
!14 = !{i64 32}
!15 = !{i64 32, i64 1, i64 1}
!16 = !DILocation(line: 15, column: 33, scope: !17, inlinedAt: !18)
!17 = distinct !DILexicalBlockFile(scope: !13, file: !4, discriminator: 0)
!18 = !DILocation(line: 23, column: 23, scope: !13)
!19 = !DILocation(line: 23, column: 4, scope: !13)

kernel.llir noinline=true

; ModuleID = 'LLVMDialectModule'
source_filename = "LLVMDialectModule"
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
target triple = "spirv64v1.3-unknown-unknown"

@printfFormat_0 = internal addrspace(2) constant [29 x i8] c"pid (%u, %u, %u) idx ()%s%p\0A\00"
@printfPrefix_0 = internal addrspace(2) constant [11 x i8] c" address: \00"

; Function Attrs: mustprogress nofree nosync nounwind willreturn memory(none)
declare !dbg !9 spir_func i64 @_Z12get_group_idj(i32) local_unnamed_addr #0

declare !dbg !12 spir_func i32 @_Z18__spirv_ocl_printf(ptr addrspace(2), ...) local_unnamed_addr

define spir_kernel void @kernel(ptr addrspace(1) %0) local_unnamed_addr !dbg !14 !intel_reqd_sub_group_size !15 !max_work_group_size !16 {
  tail call fastcc void @noinline_shared_fn__Pfp32__(ptr addrspace(1) %0), !dbg !17
  ret void, !dbg !18
}

; Function Attrs: noinline
define internal fastcc void @noinline_shared_fn__Pfp32__(ptr addrspace(1) %0) unnamed_addr #1 !dbg !19 {
  %2 = tail call spir_func i64 @_Z12get_group_idj(i32 0)
  %3 = trunc i64 %2 to i32, !dbg !20
  %4 = tail call spir_func i64 @_Z12get_group_idj(i32 1)
  %5 = trunc i64 %4 to i32, !dbg !20
  %6 = tail call spir_func i64 @_Z12get_group_idj(i32 2)
  %7 = trunc i64 %6 to i32, !dbg !20
  %8 = tail call i32 (ptr addrspace(2), ...) @_Z18__spirv_ocl_printf(ptr addrspace(2) @printfFormat_0, i32 %3, i32 %5, i32 %7, ptr addrspace(2) @printfPrefix_0, ptr addrspace(1) %0)
  ret void, !dbg !21
}

attributes #0 = { mustprogress nofree nosync nounwind willreturn memory(none) }
attributes #1 = { noinline }

!llvm.module.flags = !{!0, !1, !2}
!llvm.dbg.cu = !{!3, !5}
!opencl.spir.version = !{!6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6, !6}
!spirv.Source = !{!7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7, !7}
!llvm.ident = !{!8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8, !8}

!0 = !{i32 2, !"Debug Info Version", i32 3}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 7, !"frame-pointer", i32 2}
!3 = distinct !DICompileUnit(language: DW_LANG_C, file: !4, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!4 = !DIFile(filename: "noinline_reproducer.py", directory: "/home/amyachev/intel-xpu-backend-for-triton")
!5 = distinct !DICompileUnit(language: DW_LANG_C, file: !4, producer: "triton", isOptimized: true, runtimeVersion: 0, emissionKind: LineTablesOnly)
!6 = !{i32 1, i32 2}
!7 = !{i32 4, i32 100000}
!8 = !{!"Intel(R) oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240216)"}
!9 = !DISubprogram(name: "_Z12get_group_idj", linkageName: "_Z12get_group_idj", scope: !4, file: !4, line: 22, type: !10, scopeLine: 22, spFlags: DISPFlagOptimized)
!10 = !DISubroutineType(cc: DW_CC_normal, types: !11)
!11 = !{}
!12 = !DISubprogram(name: "_Z18__spirv_ocl_printf", linkageName: "_Z18__spirv_ocl_printf", scope: !13, file: !13, type: !10, spFlags: DISPFlagOptimized)
!13 = !DIFile(filename: "<unknown>", directory: "")
!14 = distinct !DISubprogram(name: "kernel", linkageName: "kernel", scope: !4, file: !4, line: 22, type: !10, scopeLine: 22, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !3)
!15 = !{i64 32}
!16 = !{i64 32, i64 1, i64 1}
!17 = !DILocation(line: 23, column: 23, scope: !14)
!18 = !DILocation(line: 23, column: 4, scope: !14)
!19 = distinct !DISubprogram(name: "noinline_shared_fn__Pfp32__", linkageName: "noinline_shared_fn__Pfp32__", scope: !4, file: !4, line: 10, type: !10, scopeLine: 10, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !5)
!20 = !DILocation(line: 15, column: 33, scope: !19)
!21 = !DILocation(line: 15, column: 4, scope: !19)

==================================

kernel.spv noinline=false

; SPIR-V
; Version: 1.3
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 81
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability Int8
               OpCapability SubgroupDispatch
               OpCapability KernelAttributesINTEL
               OpExtension "SPV_INTEL_kernel_attributes"
          %2 = OpExtInstImport "OpenCL.DebugInfo.100"
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %55 "kernel"
               OpExecutionMode %55 ContractionOff
               OpExecutionMode %55 SubgroupSize 32
               OpExecutionMode %55 MaxWorkgroupSizeINTEL 32 1 1
         %59 = OpString "/home/amyachev/intel-xpu-backend-for-triton/noinline_reproducer.py"
         %63 = OpString "_Z12get_group_idj"
         %65 = OpString "_Z18__spirv_ocl_printf"
         %66 = OpString "<unknown>"
         %69 = OpString "kernel"
               OpSource OpenCL_CPP 100000
               OpName %printfFormat_0 "printfFormat_0"
               OpName %printfPrefix_0 "printfPrefix_0"
               OpName %_Z12get_group_idj "_Z12get_group_idj"
               OpName %kernel "kernel"
               OpModuleProcessed "Debug info producer: triton"
               OpDecorate %printfFormat_0 Constant
               OpDecorate %printfPrefix_0 Constant
               OpDecorate %_Z12get_group_idj LinkageAttributes "_Z12get_group_idj" Import
               OpDecorate %kernel LinkageAttributes "kernel" Export
      %ulong = OpTypeInt 64 0
      %uchar = OpTypeInt 8 0
       %uint = OpTypeInt 32 0
   %ulong_29 = OpConstant %ulong 29
  %uchar_112 = OpConstant %uchar 112
  %uchar_105 = OpConstant %uchar 105
  %uchar_100 = OpConstant %uchar 100
   %uchar_32 = OpConstant %uchar 32
   %uchar_40 = OpConstant %uchar 40
   %uchar_37 = OpConstant %uchar 37
  %uchar_117 = OpConstant %uchar 117
   %uchar_44 = OpConstant %uchar 44
   %uchar_41 = OpConstant %uchar 41
  %uchar_120 = OpConstant %uchar 120
  %uchar_115 = OpConstant %uchar 115
   %uchar_10 = OpConstant %uchar 10
    %uchar_0 = OpConstant %uchar 0
   %ulong_11 = OpConstant %ulong 11
   %uchar_97 = OpConstant %uchar 97
  %uchar_114 = OpConstant %uchar 114
  %uchar_101 = OpConstant %uchar 101
   %uchar_58 = OpConstant %uchar 58
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
     %uint_2 = OpConstant %uint 2
%_arr_uchar_ulong_29 = OpTypeArray %uchar %ulong_29
%_ptr_UniformConstant__arr_uchar_ulong_29 = OpTypePointer UniformConstant %_arr_uchar_ulong_29
%_arr_uchar_ulong_11 = OpTypeArray %uchar %ulong_11
%_ptr_UniformConstant__arr_uchar_ulong_11 = OpTypePointer UniformConstant %_arr_uchar_ulong_11
         %33 = OpTypeFunction %ulong %uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
         %38 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
         %21 = OpConstantComposite %_arr_uchar_ulong_29 %uchar_112 %uchar_105 %uchar_100 %uchar_32 %uchar_40 %uchar_37 %uchar_117 %uchar_44 %uchar_32 %uchar_37 %uchar_117 %uchar_44 %uchar_32 %uchar_37 %uchar_117 %uchar_41 %uchar_32 %uchar_105 %uchar_100 %uchar_120 %uchar_32 %uchar_40 %uchar_41 %uchar_37 %uchar_115 %uchar_37 %uchar_112 %uchar_10 %uchar_0
%printfFormat_0 = OpVariable %_ptr_UniformConstant__arr_uchar_ulong_29 UniformConstant %21
         %30 = OpConstantComposite %_arr_uchar_ulong_11 %uchar_32 %uchar_97 %uchar_100 %uchar_100 %uchar_114 %uchar_101 %uchar_115 %uchar_115 %uchar_58 %uchar_32 %uchar_0
%printfPrefix_0 = OpVariable %_ptr_UniformConstant__arr_uchar_ulong_11 UniformConstant %30
         %60 = OpExtInst %void %2 DebugSource %59
         %61 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %60 Unknown
         %62 = OpExtInst %void %2 DebugTypeFunction None %void
         %64 = OpExtInst %void %2 DebugFunctionDeclaration %63 %62 %60 22 0 %61 %63 FlagIsOptimized
         %67 = OpExtInst %void %2 DebugSource %66
         %68 = OpExtInst %void %2 DebugFunctionDeclaration %65 %62 %67 0 0 %61 %65 FlagIsOptimized
         %70 = OpExtInst %void %2 DebugInfoNone
         %71 = OpExtInst %void %2 DebugFunction %69 %62 %60 22 0 %61 %69 FlagIsDefinition|FlagIsOptimized 22 %kernel %70
         %72 = OpExtInst %void %2 DebugLexicalBlockDiscriminator %60 0 %71
         %73 = OpExtInst %void %2 DebugInlinedAt 23 %71
%_Z12get_group_idj = OpFunction %ulong None %33
         %35 = OpFunctionParameter %uint
               OpFunctionEnd
     %kernel = OpFunction %void None %38
         %40 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %41 = OpLabel
         %43 = OpFunctionCall %ulong %_Z12get_group_idj %uint_0
         %74 = OpExtInst %void %2 DebugScope %72 %73
               OpLine %59 15 33
         %44 = OpUConvert %uint %43
         %75 = OpExtInst %void %2 DebugNoScope
         %46 = OpFunctionCall %ulong %_Z12get_group_idj %uint_1
         %76 = OpExtInst %void %2 DebugScope %72 %73
         %47 = OpUConvert %uint %46
         %77 = OpExtInst %void %2 DebugNoScope
         %49 = OpFunctionCall %ulong %_Z12get_group_idj %uint_2
         %78 = OpExtInst %void %2 DebugScope %72 %73
         %50 = OpUConvert %uint %49
         %79 = OpExtInst %void %2 DebugNoScope
         %52 = OpBitcast %_ptr_UniformConstant_uchar %printfFormat_0
         %53 = OpBitcast %_ptr_UniformConstant_uchar %printfPrefix_0
         %54 = OpExtInst %uint %1 printf %52 %44 %47 %50 %53 %40
         %80 = OpExtInst %void %2 DebugScope %71
               OpLine %59 23 4
               OpReturn
               OpFunctionEnd
         %55 = OpFunction %void None %38
         %56 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %57 = OpLabel
         %58 = OpFunctionCall %void %kernel %56
               OpReturn
               OpFunctionEnd

kernel.spv noinline=true

; SPIR-V
; Version: 1.3
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 87
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability Int8
               OpCapability SubgroupDispatch
               OpCapability KernelAttributesINTEL
               OpExtension "SPV_INTEL_kernel_attributes"
          %2 = OpExtInstImport "OpenCL.DebugInfo.100"
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %59 "kernel"
               OpExecutionMode %59 ContractionOff
               OpExecutionMode %59 SubgroupSize 32
               OpExecutionMode %59 MaxWorkgroupSizeINTEL 32 1 1
         %63 = OpString "/home/amyachev/intel-xpu-backend-for-triton/noinline_reproducer.py"
         %68 = OpString "_Z12get_group_idj"
         %70 = OpString "_Z18__spirv_ocl_printf"
         %71 = OpString "<unknown>"
         %74 = OpString "kernel"
         %77 = OpString "noinline_shared_fn__Pfp32__"
               OpSource OpenCL_CPP 100000
               OpName %printfFormat_0 "printfFormat_0"
               OpName %printfPrefix_0 "printfPrefix_0"
               OpName %_Z12get_group_idj "_Z12get_group_idj"
               OpName %kernel "kernel"
               OpName %noinline_shared_fn__Pfp32__ "noinline_shared_fn__Pfp32__"
               OpModuleProcessed "Debug info producer: triton"
               OpModuleProcessed "Debug info producer: triton"
               OpDecorate %printfFormat_0 Constant
               OpDecorate %printfPrefix_0 Constant
               OpDecorate %_Z12get_group_idj LinkageAttributes "_Z12get_group_idj" Import
               OpDecorate %kernel LinkageAttributes "kernel" Export
      %ulong = OpTypeInt 64 0
      %uchar = OpTypeInt 8 0
       %uint = OpTypeInt 32 0
   %ulong_29 = OpConstant %ulong 29
  %uchar_112 = OpConstant %uchar 112
  %uchar_105 = OpConstant %uchar 105
  %uchar_100 = OpConstant %uchar 100
   %uchar_32 = OpConstant %uchar 32
   %uchar_40 = OpConstant %uchar 40
   %uchar_37 = OpConstant %uchar 37
  %uchar_117 = OpConstant %uchar 117
   %uchar_44 = OpConstant %uchar 44
   %uchar_41 = OpConstant %uchar 41
  %uchar_120 = OpConstant %uchar 120
  %uchar_115 = OpConstant %uchar 115
   %uchar_10 = OpConstant %uchar 10
    %uchar_0 = OpConstant %uchar 0
   %ulong_11 = OpConstant %ulong 11
   %uchar_97 = OpConstant %uchar 97
  %uchar_114 = OpConstant %uchar 114
  %uchar_101 = OpConstant %uchar 101
   %uchar_58 = OpConstant %uchar 58
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
     %uint_2 = OpConstant %uint 2
%_arr_uchar_ulong_29 = OpTypeArray %uchar %ulong_29
%_ptr_UniformConstant__arr_uchar_ulong_29 = OpTypePointer UniformConstant %_arr_uchar_ulong_29
%_arr_uchar_ulong_11 = OpTypeArray %uchar %ulong_11
%_ptr_UniformConstant__arr_uchar_ulong_11 = OpTypePointer UniformConstant %_arr_uchar_ulong_11
         %33 = OpTypeFunction %ulong %uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
         %38 = OpTypeFunction %void %_ptr_CrossWorkgroup_uchar
%_ptr_UniformConstant_uchar = OpTypePointer UniformConstant %uchar
         %21 = OpConstantComposite %_arr_uchar_ulong_29 %uchar_112 %uchar_105 %uchar_100 %uchar_32 %uchar_40 %uchar_37 %uchar_117 %uchar_44 %uchar_32 %uchar_37 %uchar_117 %uchar_44 %uchar_32 %uchar_37 %uchar_117 %uchar_41 %uchar_32 %uchar_105 %uchar_100 %uchar_120 %uchar_32 %uchar_40 %uchar_41 %uchar_37 %uchar_115 %uchar_37 %uchar_112 %uchar_10 %uchar_0
%printfFormat_0 = OpVariable %_ptr_UniformConstant__arr_uchar_ulong_29 UniformConstant %21
         %30 = OpConstantComposite %_arr_uchar_ulong_11 %uchar_32 %uchar_97 %uchar_100 %uchar_100 %uchar_114 %uchar_101 %uchar_115 %uchar_115 %uchar_58 %uchar_32 %uchar_0
%printfPrefix_0 = OpVariable %_ptr_UniformConstant__arr_uchar_ulong_11 UniformConstant %30
         %64 = OpExtInst %void %2 DebugSource %63
         %65 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %64 Unknown
         %66 = OpExtInst %void %2 DebugCompilationUnit 65536 0 %64 Unknown
         %67 = OpExtInst %void %2 DebugTypeFunction None %void
         %69 = OpExtInst %void %2 DebugFunctionDeclaration %68 %67 %64 22 0 %66 %68 FlagIsOptimized
         %72 = OpExtInst %void %2 DebugSource %71
         %73 = OpExtInst %void %2 DebugFunctionDeclaration %70 %67 %72 0 0 %66 %70 FlagIsOptimized
         %75 = OpExtInst %void %2 DebugInfoNone
         %76 = OpExtInst %void %2 DebugFunction %74 %67 %64 22 0 %65 %74 FlagIsDefinition|FlagIsOptimized 22 %kernel %75
         %78 = OpExtInst %void %2 DebugFunction %77 %67 %64 10 0 %66 %77 FlagIsDefinition|FlagIsOptimized 10 %noinline_shared_fn__Pfp32__ %75
%_Z12get_group_idj = OpFunction %ulong None %33
         %35 = OpFunctionParameter %uint
               OpFunctionEnd
     %kernel = OpFunction %void None %38
         %40 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %41 = OpLabel
         %79 = OpExtInst %void %2 DebugScope %76
               OpLine %63 23 23
         %44 = OpFunctionCall %void %noinline_shared_fn__Pfp32__ %40
               OpLine %63 23 4
               OpReturn
               OpFunctionEnd
%noinline_shared_fn__Pfp32__ = OpFunction %void DontInline %38
         %43 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %45 = OpLabel
         %47 = OpFunctionCall %ulong %_Z12get_group_idj %uint_0
         %80 = OpExtInst %void %2 DebugScope %78
               OpLine %63 15 33
         %48 = OpUConvert %uint %47
         %81 = OpExtInst %void %2 DebugNoScope
         %50 = OpFunctionCall %ulong %_Z12get_group_idj %uint_1
         %82 = OpExtInst %void %2 DebugScope %78
         %51 = OpUConvert %uint %50
         %83 = OpExtInst %void %2 DebugNoScope
         %53 = OpFunctionCall %ulong %_Z12get_group_idj %uint_2
         %84 = OpExtInst %void %2 DebugScope %78
         %54 = OpUConvert %uint %53
         %85 = OpExtInst %void %2 DebugNoScope
         %56 = OpBitcast %_ptr_UniformConstant_uchar %printfFormat_0
         %57 = OpBitcast %_ptr_UniformConstant_uchar %printfPrefix_0
         %58 = OpExtInst %uint %1 printf %56 %48 %51 %54 %57 %43
         %86 = OpExtInst %void %2 DebugScope %78
               OpLine %63 15 4
               OpReturn
               OpFunctionEnd
         %59 = OpFunction %void None %38
         %60 = OpFunctionParameter %_ptr_CrossWorkgroup_uchar
         %61 = OpLabel
         %62 = OpFunctionCall %void %kernel %60
               OpReturn
               OpFunctionEnd

[pbchekin]

Blocked by the IGC issue.