Twice as many kernel arguments as specified in SPIR

[jonasbhjulstad]

I have defined a simple kernel function that takes one global input argument, which is compiled from OpenCL for C++ (-cl-std=clc++2021) in clang.

__kernel void foo(global float* x0)
{
    int a = 0;
}

Which gives the following readable output (using -emit-llvm and llvm-spirv):

119734787 65536 393230 19 0 
2 Capability Addresses 
2 Capability Linkage 
2 Capability Kernel 
5 ExtInstImport 1 "OpenCL.std" 
3 MemoryModel 1 2 
4 EntryPoint 6 15 "foo" 
3 Source 3 300000 
3 Name 6 "foo" 
3 Name 16 "x0" 
5 Decorate 6 LinkageAttributes "foo" Export 
4 Decorate 7 Alignment 4 
4 Decorate 10 Alignment 4 
4 Decorate 13 Alignment 4 
4 Decorate 16 Alignment 4 
4 TypeInt 11 32 0 
4 Constant 11 14 0 
2 TypeVoid 2 
3 TypeFloat 3 32 
4 TypePointer 4 5 3 
4 TypeFunction 5 2 4 
4 TypePointer 9 7 4 
4 TypePointer 12 7 11 


5 Function 2 6 2 5 
3 FunctionParameter 4 7 

2 Label 8 
4 Variable 9 10 7 
4 Variable 12 13 7 
5 Store 10 7 2 4 
5 Store 13 14 2 4 
1 Return 

1 FunctionEnd 

5 Function 2 15 2 5 
3 FunctionParameter 4 16 

2 Label 17 
5 FunctionCall 2 18 6 16 
1 Return 

1 FunctionEnd 

The binary version of spirv is loaded into OpenCL using clCreateProgramWithIL, which builds fine makes it possible to create a kernel for foo.

However, when specifying kernel input arguments I get status error CL_INVALID_KERNEL_ARGS.

Retrieving input argument names with clGetKernelArgInfo(.) yields two results: ‘x0’ and ‘’, both having CL_KERNEL_ARG_TYPE_NAME float*.

setting more than 1 kernel arg with clSetKernelArg results in errors, and increasing the number of kernel input arguments (e.g foo(x0, x1)) results in two new results each.

Is the second input argument a bug, or is it supposed to specify something else when creating programs from SPIR-V?

Edit:
Changing global float* to other types (e.g float) and adjusting clSetKernelArg results in the same error.