Add a native compiler, using SPIR-V IL.

[maleadt]

Initial sketch of a native compiler, relying on GPUCompiler's SPIR-V back-end and the ability to load IL into (some) OpenCL drivers.

Demo:

julia> using OpenCL, pocl_jll

julia> cl.platform!("pocl")
OpenCL.Platform('Portable Computing Language' @0x00007b31bede77e8)

julia> function vadd(a, b, c)
           gid = get_global_id(1)
           @inbounds c[gid] = a[gid] + b[gid]
           return
       end
vadd (generic function with 1 method)

julia> a = rand(Float32, 50_000);

julia> b = rand(Float32, 50_000);

julia> d_a = CLArray(a; access=:r);

julia> d_b = CLArray(b; access=:r);

julia> d_c = similar(d_a; access=:w);

julia> @device_code_native @opencl global_size=size(a) vadd(d_a, d_b, d_c)
// GPUCompiler.CompilerJob{GPUCompiler.SPIRVCompilerTarget, OpenCL.OpenCLCompilerParams}(MethodInstance for vadd(::CLDeviceVector{Float32, 1}, ::CLDeviceVector{Float32, 1}, ::CLDeviceVector{Float32, 1}), CompilerConfig for GPUCompiler.SPIRVCompilerTarget, 0x0000000000007b03)

; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 43
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability Int8
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %37 "_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE" %__spirv_BuiltInGlobalInvocationId
               OpExecutionMode %37 ContractionOff
               OpSource OpenCL_C 200000
               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
               OpName %_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE "_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE"
               OpName %conversion "conversion"
               OpName %unbox21 "unbox21"
               OpName %unbox72 "unbox72"
               OpName %unbox123 "unbox123"
               OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
               OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
               OpDecorate %_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE LinkageAttributes "_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE" Export
               OpDecorate %16 FuncParamAttr ByVal
               OpDecorate %17 FuncParamAttr ByVal
               OpDecorate %18 FuncParamAttr ByVal
               OpDecorate %38 FuncParamAttr ByVal
               OpDecorate %39 FuncParamAttr ByVal
               OpDecorate %40 FuncParamAttr ByVal
      %ulong = OpTypeInt 64 0
      %uchar = OpTypeInt 8 0
    %ulong_1 = OpConstant %ulong 1
    %v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
%_arr_ulong_ulong_1 = OpTypeArray %ulong %ulong_1
  %_struct_8 = OpTypeStruct %_ptr_CrossWorkgroup_uchar %ulong %_arr_ulong_ulong_1 %ulong
  %_struct_7 = OpTypeStruct %_struct_8
%_ptr_Function__struct_7 = OpTypePointer Function %_struct_7
         %14 = OpTypeFunction %void %_ptr_Function__struct_7 %_ptr_Function__struct_7 %_ptr_Function__struct_7
      %float = OpTypeFloat 32
%_ptr_CrossWorkgroup_float = OpTypePointer CrossWorkgroup %float
%_ptr_Function__ptr_CrossWorkgroup_float = OpTypePointer Function %_ptr_CrossWorkgroup_float
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE = OpFunction %void None %14
         %16 = OpFunctionParameter %_ptr_Function__struct_7
         %17 = OpFunctionParameter %_ptr_Function__struct_7
         %18 = OpFunctionParameter %_ptr_Function__struct_7
 %conversion = OpLabel
         %20 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
         %21 = OpCompositeExtract %ulong %20 0
         %25 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_float %16
    %unbox21 = OpLoad %_ptr_CrossWorkgroup_float %25 Aligned 8
         %27 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %unbox21 %21
         %28 = OpLoad %float %27 Aligned 4
         %29 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_float %17
    %unbox72 = OpLoad %_ptr_CrossWorkgroup_float %29 Aligned 8
         %31 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %unbox72 %21
         %32 = OpLoad %float %31 Aligned 4
         %33 = OpFAdd %float %28 %32
         %34 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_float %18
   %unbox123 = OpLoad %_ptr_CrossWorkgroup_float %34 Aligned 8
         %36 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_float %unbox123 %21
               OpStore %36 %33 Aligned 4
               OpReturn
               OpFunctionEnd
         %37 = OpFunction %void None %14
         %38 = OpFunctionParameter %_ptr_Function__struct_7
         %39 = OpFunctionParameter %_ptr_Function__struct_7
         %40 = OpFunctionParameter %_ptr_Function__struct_7
         %41 = OpLabel
         %42 = OpFunctionCall %void %_Z4vadd13CLDeviceArrayI7Float32Ll1ELl1EES_IS0_Ll1ELl1EES_IS0_Ll1ELl1EE %38 %39 %40
               OpReturn
               OpFunctionEnd

julia> a + b ≈ Array(d_c)
true

[codecov]

Codecov Report

Attention: Patch coverage is 82.55034% with 26 lines in your changes missing coverage. Please review.

Project coverage is 79.74%. Comparing base (aba4fce) to head (9656356).
Report is 1 commits behind head on master.

Files with missing lines	Patch %	Lines
src/compiler/execution.jl	78.40%	19 Missing ⚠️
src/compiler/reflection.jl	82.60%	4 Missing ⚠️
src/compiler/compilation.jl	91.42%	3 Missing ⚠️

Additional details and impacted files

@@            Coverage Diff             @@
##           master     #222      +/-   ##
==========================================
+ Coverage   77.24%   79.74%   +2.50%     
==========================================
  Files           3        6       +3     
  Lines         167      316     +149     
==========================================
+ Hits          129      252     +123     
- Misses         38       64      +26     

☔ View full report in Codecov by Sentry.
📢 Have feedback on the report? Share it here.