KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

C++ for OpenCL: Unresolved external reference to _Z13get_global_idj

Open
#2,908 5 comments 0 reactions 0 assignees View on GitHub
Dominant language
LLVM
Stars
625
Forks
279
Avg merge
3d 5h
Merged PRs (30d)
34

Description

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:
```C
// 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: https://github.com/llvm/llvm-project/issues/118576

Contributor guide

Open the contributing guide

Research direction

Start by reproducing the issue with the provided clang -cl-std=clc++ pipeline, then inspect how the !opencl.cxx.version metadata is handled by llvm-spirv. Compare the generated SPIR-V with and without that metadata; done means get_global_id and write_imagef are emitted as builtins and the result links successfully with spirv-link.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
42/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.