OpenCL: Excessive kernel arguments

[linehill]

clGetKernelInfo() incorrectly reports that the saxpy kernel, as shown in the following SPIR-V listing, has eight kernel arguments instead of four - the original amount.

; A test case that triggers an issues in Intel's OpenCL.
;
;  * With Intel Compute Runtime version 21.49.21786 the 'saxpy' kernel appears
;    in a different name (e.g. to "saxpy.1").
;
;  * With Intel Compute Runtime version 22.16.22992 the 'saxpy' appears to have
;    eight kernel arguments instead of four.
;
; The issue disappears in both versions if both the 'OpName %saxpy "saxpy"' and
; 'OpDecorate %saxpy LinkageAttributes "saxpy" Export' instructions are removed.
;
; Generated originally from OpenCL:
;
;   kernel void saxpy(global float *z, global const float *x,
;                     global const float *y, float a) {
;     size_t id = get_global_id(0);
;     z[id] = a * x[id] + y[id];
;   }
;
; Assemble with command:
;
;  spirv-as --target-env opencl2.0 saxpy-kernel.spt -o saxpy-kernel.spv
;
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %24 "saxpy" %__spirv_BuiltInGlobalInvocationId
               OpSource OpenCL_C 200000
               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
               OpName %saxpy "saxpy"
               OpName %z "z"
               OpName %x "x"
               OpName %y "y"
               OpName %a "a"
               OpName %entry "entry"
               OpName %call "call"
               OpName %arrayidx "arrayidx"
               OpName %arrayidx1 "arrayidx1"
               OpName %arrayidx2 "arrayidx2"
               OpName %z_0 "z"
               OpName %x_0 "x"
               OpName %y_0 "y"
               OpName %a_0 "a"
               OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
               OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
               OpDecorate %saxpy LinkageAttributes "saxpy" Export
               OpDecorate %z FuncParamAttr NoCapture
               OpDecorate %x FuncParamAttr NoCapture
               OpDecorate %x FuncParamAttr NoWrite
               OpDecorate %y FuncParamAttr NoCapture
               OpDecorate %y FuncParamAttr NoWrite
               OpDecorate %z_0 FuncParamAttr NoCapture
               OpDecorate %x_0 FuncParamAttr NoCapture
               OpDecorate %x_0 FuncParamAttr NoWrite
               OpDecorate %y_0 FuncParamAttr NoCapture
               OpDecorate %y_0 FuncParamAttr NoWrite
      %ulong = OpTypeInt 64 0
    %v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
       %void = OpTypeVoid
      %float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
          %9 = OpTypeFunction %void %_ptr_CrossWorkgroup_float %_ptr_CrossWorkgroup_float %_ptr_CrossWorkgroup_float %float
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
      %saxpy = OpFunction %void None %9
          %z = OpFunctionParameter %_ptr_CrossWorkgroup_float
          %x = OpFunctionParameter %_ptr_CrossWorkgroup_float
          %y = OpFunctionParameter %_ptr_CrossWorkgroup_float
          %a = OpFunctionParameter %float
      %entry = OpLabel
         %16 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
       %call = OpCompositeExtract %ulong %16 0
   %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %x %call
         %19 = OpLoad %float %arrayidx Aligned 4
  %arrayidx1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %y %call
         %21 = OpLoad %float %arrayidx1 Aligned 4
         %22 = OpExtInst %float %1 mad %a %19 %21
  %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %z %call
               OpStore %arrayidx2 %22 Aligned 4
               OpReturn
               OpFunctionEnd
         %24 = OpFunction %void None %9
        %z_0 = OpFunctionParameter %_ptr_CrossWorkgroup_float
        %x_0 = OpFunctionParameter %_ptr_CrossWorkgroup_float
        %y_0 = OpFunctionParameter %_ptr_CrossWorkgroup_float
        %a_0 = OpFunctionParameter %float
         %29 = OpLabel
         %30 = OpFunctionCall %void %saxpy %z_0 %x_0 %y_0 %a_0
               OpReturn
               OpFunctionEnd

Attempting to launch the kernel with the original amount of arguments fails too.

The code was originally generated by the latest llvm-spirv tool from the llvm_release_140. The issue disappears if both the OpName %saxpy "saxpy" and “OpDecorate %saxpy LinkageAttributes "saxpy" Export” lines are removed from the code.

Test environment: OS: Ubuntu 20.04.4 LTS Intel Compute Runtime: 22.16.22992. Device: Intel(R) HD Graphics 530 [0x1912].

[eero-t]

The code was originally generated by the latest llvm-spirv tool from the llvm_release_140.

intel-graphics-compiler (used by compute-runtime to compile SPIRV / shaders), does not support LLVM 14 quite yet: https://github.com/intel/intel-graphics-compiler/pull/242

And there's an issue with LLVM 13: https://github.com/intel/intel-graphics-compiler/issues/236

LLVM 11 is best supported, but LLVM 12 should work too. Do you get that issue also with either of those, or is it LLVM 14 specific?

(Everything in the compute stack obviously needs to be using same LLVM version because LLVM changes its API in every major release.)

[linehill]

Do you get that issue also with either of those, or is it LLVM 14 specific?

I don’t think the issue is strictly specific to a LLVM version but to the llvm-spirv tool revision we use in the CHIP-SPV. Nevertheless, the LLVM version should not matter as the pasted code is valid SPIR-V, right?

[eero-t]

It matters whether you've compiled everything in compute stack with same LLVM and llv-spirv version or not, and which LLVM version you use for this. Mixing different LLVM versions won't work (because LLVM breaks API between major versions), neither does mixing llvm-spirv versions, and LLVM 14 is not yet supported by IGC, so using "llvm_release_140" won't work (yet) either.

For details, see the tickets I listed for LLVM version support, and this one for SPIRV side of it: https://github.com/intel/intel-graphics-compiler/issues/224

[linehill]

I’m using released binaries of the compute runtime from https://github.com/intel/compute-runtime/releases so the versions of the components should be matching. Not sure if you are trying to tell that SPIR-V binary produced by llvm-spirv-14 is incompatible for LLVM 11 based compute runtime. Is that the case?

[eero-t]

Not sure if you are trying to tell that SPIR-V binary produced by llvm-spirv-14 is incompatible for LLVM 11 based compute runtime. Is that the case?

From https://github.com/intel/intel-graphics-compiler/issues/224

"libLLVMSPIRVLib used by OpenCL-clang and Vector Compiler must be the same ... If you want to build your own SPIRV_TRANSLATOR, you need to rebuild openCL-clang."

And OpenCL-clang needs to be built by same LLVM as rest of the compute stack, for LLVM API change reasons.

=> You should at least try a stack where everything is built using same versions of both LLVM and SPIRV, to see whether your problem goes away

[linehill]

Sorry, perhaps there are missing critical details to the issue: We (CHIP-SPV) are not compiling kernels from OpenCL sources but from HIP sources and we pass the produced SPIR-V modules through OpenCL API via clCreateProgramWithIL().

AFAIK, the OpenCL, SPIR-V and OpenCL SPIR-V Environment specification does not impose restrictions on how the SPIR-V binaries are created as long as the binaries are valid SPIR-V (the pasted code in the above should be valid SPIR-V). Let me know if this is not the case.

[kmazurki]

Hey folks, this looks like a problem addressed with this commit: https://github.com/intel/intel-graphics-compiler/commit/6a13fa903f380e17378286a7cd43995b0ae162ad. This change should land in a binary release expected next week.

[pengtu]

Hey folks, this looks like a problem addressed with this commit: intel/intel-graphics-compiler@6a13fa9. This change should land in a binary release expected next week.

Thank you @kmazurki !

[MathiasMagnus]

@kmazurki When can these changes be expected to land for users of the oclcpuexp package? Daily builds on intel/llvm don't sport that package, only full releases have that asset. We'd like to use the streamlined CLI for SPIR-V compilation which only Clang 14 has.

[bashbaug]

There is a user on the Intel OpenCL forums who is also encountering this issue also after updating drivers: link.

Is there a workaround?

CC @biggysmith (I think this is the same submitter)

[biggysmith]

yes I'm the same submitter. I also get the additional ".1"