What are the required Spirv kernel attributes and metadata in LLVM IR?

[fwinter]

Creating a kernel (zeKernelCreate) failed on a kernel that I built with IRBuilder and then translated by LLVMSPIRLib (see below). The call failed returning number 2013265944 which does not match any of these: ZE_RESULT_ERROR_UNINITIALIZED, ZE_RESULT_ERROR_DEVICE_LOST, ZE_RESULT_ERROR_INVALID_ARGUMENT. I have no idea what went wrong. Best guess: I didn't provide the required attributes/metadata?!

Here's the source IR (self built):

target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
target triple = "spir64"

define spir_kernel void @eval_scalar0(i32* %arg0, float* %arg1, float* %arg2, float* %arg3) {
stack:
  br label %afterstack

afterstack:                                       ; preds = %stack
  %0 = call i64 @get_global_id(i32 0)
  %1 = getelementptr i32, i32* %arg0, i64 %0
  %2 = load i32, i32* %1, align 4
  %3 = add nsw i32 0, %2
  %4 = getelementptr float, float* %arg2, i32 %3
  %5 = load float, float* %4, align 4
  %6 = add nsw i32 0, %2
  %7 = getelementptr float, float* %arg3, i32 %6
  %8 = load float, float* %7, align 4
  %9 = fadd float %8, %5
  %10 = add nsw i32 0, %2
  %11 = getelementptr float, float* %arg1, i32 %10
  store float %9, float* %11, align 4
  ret void
}

; Function Attrs: readnone
declare spir_func i64 @get_global_id(i32) #0

attributes #0 = { readnone }

If I look at a simple kernel generated by clang from an OpenCL source I see a lot of decorations at the kernel:

define dso_local spir_kernel void @mxm(i32 addrspace(1)* nocapture readonly %a, i32 addrspace(1)* nocapture readonly %b, i32 addrspace(1)* nocapture %c, i32 %n) 
local_unnamed_addr #0 
!kernel_arg_addr_space !3 
!kernel_arg_access_qual !4 
!kernel_arg_type !5 
!kernel_arg_base_type !5 
!kernel_arg_type_qual !6 

attributes #0 = { convergent nofree norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
attributes #1 = { convergent mustprogress nofree nounwind readnone willreturn "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
attributes #2 = { convergent nounwind readnone willreturn }

!llvm.module.flags = !{!0}
!opencl.ocl.version = !{!1}
!opencl.spir.version = !{!1}
!llvm.ident = !{!2}

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 2}
!2 = !{!"clang version 13.0.0"}
!3 = !{i32 1, i32 1, i32 1, i32 0}
!4 = !{!"none", !"none", !"none", !"none"}
!5 = !{!"int*", !"int*", !"int*", !"int"}
!6 = !{!"", !"", !"", !""}
!7 = !{!8, !8, i64 0}
!8 = !{!"int", !9, i64 0}
!9 = !{!"omnipotent char", !10, i64 0}
!10 = !{!"Simple C/C++ TBAA"}

Which of these attributes and metadata are needed so that a Spirv kernel can be created with zeKernelCreate? There is no documentation on this, or is there?

[jandres742]

@fwinter please open the issue in https://github.com/intel/compute-runtime, as your error seems to be related with implementation.

In particular, your error is: ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED

So it seems that your kernel is not correctly linked, or you are missing some files/symbols.

[fwinter]

Okay, thanks! Will do.