OpenCL SPIRV Unresolved external reference to _Z13get_global_idj

[Calandracas606]

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

This is only occurring when using the C++ for OpenCL mode

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);
}

When compiling with -cl-std=clc++

$ clang -c -target spirv64 -cl-std=clc++ -o sample_cpp.spv kernel.cl
$ spirv-link sample_cpp.spv 
error: 0: Unresolved external reference to "_Z13get_global_idj".

When I examine sample_cpp.spv with spirv-dis I see:

• OpFunctionCall %ulong %_Z13get_global_idj %uint_0
• OpFunctionCall %void %_Z12write_imagef14ocl_image2d_woDv2_iDv4_f %19 %29 %32

This appears to be incorrect.

When i repeat the same process but without the -cl-std=clc++ flag:

$ clang -c -target spirv64 -o sample.spv kernel.cl
$ spirv-link sample.spv 

I get no errors.

Examining the spirv output, I see the expected output in the spirv:

• %21 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
• OpImageWrite %19 %28 %18

I have tried this with multiple versions of clang, and it happens with clang>=19. It does not occur with clang18

clang version 20.0.0git 14a259f

[Calandracas606]

git bisect pointed me to this commit: 5fd5b8a

[Calandracas606]

this is unlikely to be an issue with clang, it looks like a bug with the spirv-translator not dealing with OpSource OpenCL_CPP 100000

[svenvh]

This is an issue with the SPIRV-LLVM-Translator project indeed; closing this issue as we have an issue in that project.