SPIR support

SPIR support

Hi, i'm trying to load a SPIR64 binary with the (latest?) 14.1 runtime on linux. I compiled the OpenCL C code with Kernel Builder to SPIR 64. The readable SPIR is converted to BC with llvm-link and spir_verifier says all is good to go. The code to link the binary at runtime:

cl::Program::Binaries bin;
bin.push_back(std::make_pair((void*)llvm_start, llvm_size));
std::vector<int> status;
cl_int err;
cl::Program prog(_current_device->ctx, dev, bin, &status, &err);
prog.build(dev, "-x spir -spir-std=1.2");


The build log does not provide any useful output:

Linking started
Linking failed

Anyone here who has expirience with SPIR and Intel SDK and can provide some help?

Thanks in advance!


5 posts / 0 new
Last post
For more complete information about compiler optimizations, see our Optimization Notice.

Hi Michael,

Could you please share the kernel you are trying to compile?



Hi Boaz,

here is the kernel:


__kernel void foo()

and here is the SPIR version generated by the Kernel Builder:

; ModuleID = 'qt_temp.cr1524'
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024"
target triple = "spir64-unknown-unknown"

define cc76 void @foo() nounwind readnone {
  ret void

!opencl.kernels = !{!0}
!opencl.enable.FP_CONTRACT = !{}
!opencl.spir.version = !{!7}
!opencl.ocl.version = !{!7}
!opencl.used.extensions = !{!8}
!opencl.used.optional.core.features = !{!8}
!opencl.compiler.options = !{!8}

!0 = metadata !{void ()* @foo, metadata !1, metadata !2, metadata !3, metadata !4, metadata !5, metadata !6}
!1 = metadata !{metadata !"kernel_arg_addr_space"}
!2 = metadata !{metadata !"kernel_arg_access_qual"}
!3 = metadata !{metadata !"kernel_arg_type"}
!4 = metadata !{metadata !"kernel_arg_type_qual"}
!5 = metadata !{metadata !"kernel_arg_base_type"}
!6 = metadata !{metadata !"kernel_arg_name"}
!7 = metadata !{i32 1, i32 2}
!8 = metadata !{}


Hi Michael,

I couldn't reproduce this issue so far. SPIR LLVM (text) representation is the same. But I don't quite understand how did you get the SPIR binary. Could you please generate it using ioc64 instead of KernelBuilder with the following command:
ioc64 -input=sample.cl -spir64=sample.bc64

and check if the SPIR binaries are the same and if they are different - try run the program again.

In case nothing helps - I would like to ask for complete reproducer, including host program and generated SPIR binary.


Hi Yuri,

sometimes the problem is so small that it gets overlooked every time.

In detail: at first I created the bitcode with llvm-as from the current LLVM trunk, than after Boaz told me only LLVM 3.2 is supported, I switched to the 3.2 toolchain. Same error all the time. Your suggestion to check the binaries for equality was the right push. The binary built with llvm-as was exactly equal to the binary from ioc64, however my program loaded the binary with the wrong size, since it was old code which was written for plain OpenCL C and computed the size for the program with a final '\0'. That missmatch in size was the problem. Now all is fine. The Intel OpenCL runtime loaded the binary built with LLVM 3.5 toolchain as well without any problem.

Thank you very much for your help Yuri, and of course Boaz.

Leave a Comment

Please sign in to add a comment. Not a member? Join today