C++ for OpenCL: Unresolved external reference to _Z13get_global_idj

[Calandracas606]

The translator is generating spirv with function calls which should be treated as builtins.

This is only occurring when clang is run with the C++ for OpenCL mode.

When run in clc++ mode, clang generates a !opencl.cxx.version = !{!3} entry, which seems to be causing some issues.

Here is an easy way to reproduce:

kernel.cl:

// this sample doesn't use any C++ for OpenCL features
__kernel void sample(write_only image2d_t img) {

  const int col = get_global_id(0);
  const int row = get_global_id(1);
  int2 coord;
  coord.x = col;
  coord.y = row;

  float4 px = {0, 0, 0, 1};

  write_imagef(img, coord, px);
}

run:
clang -target spirv64 -cl-std=clc++ kernel.cl
error: 0: Unresolved external reference to "_Z13get_global_idj".

Examine the generated spirv:
clang -c -target spirv64 -cl-std=clc++ -emit-llvm -o - kernel.cl | llvm-spirv -o - | spirv-dis
Notice how the spirv contains some odd OpFunctionCall:

         %22 = OpFunctionCall %ulong %_Z13get_global_idj %uint_0
         %23 = OpUConvert %uint %22
         %25 = OpFunctionCall %ulong %_Z13get_global_idj %uint_1
         %26 = OpUConvert %uint %25
         %28 = OpCompositeInsert %v2uint %23 %27 0
         %29 = OpCompositeInsert %v2uint %26 %28 1
         %33 = OpFunctionCall %void %_Z12write_imagef14ocl_image2d_woDv2_iDv4_f %19 %29 %32

Looking at the llvm generated by clang:
clang -c -target spirv64 -cl-std=clc++ -emit-llvm -o - kernel.cl | llvm-dis
Notice the !opencl.cxx.version = !{!3}

If I edit the llvm, and pass it back through llvm-spirv:

clang -c -target spirv64 -cl-std=clc++ -emit-llvm -o - kernel.cl \
     | llvm-dis -o - \
     | sed '/!opencl.cxx.version/d' \
     | llvm-as -o - \
     | llvm-spirv -o - \
     | spirv-dis

The output appears normal:

         %12 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
         %13 = OpCompositeExtract %ulong %12 0
         %15 = OpUConvert %uint %13
         %16 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
         %17 = OpCompositeExtract %ulong %16 1
         %18 = OpUConvert %uint %17
         %21 = OpCompositeInsert %v2uint %15 %20 0
         %22 = OpCompositeInsert %v2uint %18 %21 1
                     OpImageWrite %10 %22 %27

We can verify that this links correctly by running spirv-link.

Repeating the same processes, but running clang without -cl-std=clc++ results in no errors. The only meaningful difference in the generated llvm is the !opencl.cxx.version = !{!3} tag

related issue: llvm/llvm-project#118576

[svenvh]

Thanks for reporting; this highlights a long-standing problem that doesn't have a trivial solution unfortunately. See e.g. the attempt reverted by #2508.

[davidrohr]

Any progress on this, or on llvm/llvm-project#122453 ?
This is really preventing us from debugging with printf at the moment.
Or is there any workaround we could try?

[svenvh]

Or is there any workaround we could try?

Did you try the workaround suggested by the reporter (i.e. stripping the !opencl.cxx.version metadata from the LLVM IR before passing it to llvm-spirv)?

Alternatively, in llvm/llvm-project#122453 you mentioned that it works using llvm's SPIR-V target (which was promoted out of its "experimental" status last month). I wonder if you still need llvm-spirv for producing SPIR-V, or if you can switch to the llvm SPIR-V target?

[MrSidims]

I believe @maarquitos14 is working on this.

[maarquitos14]

I believe @maarquitos14 is working on this.

I am indeed working on this. We're prototyping a new solution, but it's not ready yet.