<div dir="ltr">Does anyone know what the status of SPIR generation for OpenCL is in clang? I'm able to generate a SPIR file with clang 3.6, but it seems to miss some of the necessary attributes.<div><br></div><div>If I take this simple file</div><div><br></div><div>//<a href="http://hello.cl">hello.cl</a></div><div><p style="margin:0px;font-size:11px;font-family:Menlo">__kernel void hello(int x) {</p>
<p style="margin:0px;font-size:11px;font-family:Menlo">}</p><p style="margin:0px;font-size:11px;font-family:Menlo"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo">then run </p><p style="margin:0px;font-size:11px;font-family:Menlo"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo">clang -x cl -fno-builtin -target spir -c -emit-llvm <a href="http://hello.cl">hello.cl</a></p><p style="margin:0px;font-size:11px;font-family:Menlo"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo">it generates a bc file, but it doesn't have the additional attributes that describe the function arguments.</p><p style="margin:0px;font-size:11px;font-family:Menlo"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo">Here's the comparison with the output from Khronos group's version of clang (<a href="https://github.com/KhronosGroup/SPIR">https://github.com/KhronosGroup/SPIR</a>)</p><p style="margin:0px;font-size:11px;font-family:Menlo"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo">########################## clang 3.6</p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">; ModuleID = 'hello.bc' </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:51<span style="background-color:rgb(48,48,48)">2</span>-v1024:1024"</p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">target triple = "spir" </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">; Function Attrs: nounwind readnone </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">define void @hello(i32 %x) #0 { </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">entry: </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)"> ret void </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">attributes #0 = { nounwind readnone "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no<span style="background-color:rgb(48,48,48)">-</span>frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }</p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.kernels = !{!0} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!llvm.ident = !{!1} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!0 = metadata !{void (i32)* @hello} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p><p style="margin:0px;font-size:11px;font-family:Menlo">
</p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!1 = metadata !{metadata !"clang version 3.6.0 "} </p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)"><br></p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)"><span style="color:rgb(34,34,34);background-color:rgb(255,255,255)">########################## Khronos </span></p><p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">; ModuleID = 'hello.bc' </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:<span style="background-color:rgb(48,48,48)">1</span>024"</p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">target triple = "spir64-unknown-unknown" </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">; Function Attrs: nounwind readnone </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">define spir_kernel void @hello(i32 %x) #0 { </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">entry: </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)"> ret void </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">attributes #0 = { nounwind readnone } </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.kernels = !{!0} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.enable.FP_CONTRACT = !{} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.spir.version = !{!6} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.ocl.version = !{!6} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.used.extensions = !{!7} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.used.optional.core.features = !{!7} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!opencl.compiler.options = !{!8} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18);min-height:13px"><span style="color:rgb(208,208,208)"> </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!0 = metadata !{void (i32)* @hello, metadata !1, metadata !2, metadata !3, metadata !4, metadata !5<span style="background-color:rgb(48,48,48)">}</span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!1 = metadata !{metadata !"kernel_arg_addr_space", i32 0} <span style="color:rgb(0,0,0)"> </span><span style="background-color:rgb(48,48,48)"> </span><span style="color:rgb(0,0,0)"> </span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!2 = metadata !{metadata !"kernel_arg_access_qual", metadata !"none"} <span style="color:rgb(0,0,0)"> </span><span style="background-color:rgb(48,48,48)"> </span><span style="color:rgb(0,0,0)"> </span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!3 = metadata !{metadata !"kernel_arg_type", metadata !"int"} <span style="color:rgb(0,0,0)"> </span><span style="background-color:rgb(48,48,48)"> </span><span style="color:rgb(0,0,0)"> </span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!4 = metadata !{metadata !"kernel_arg_type_qual", metadata !""} <span style="color:rgb(0,0,0)"> </span><span style="background-color:rgb(48,48,48)"> </span><span style="color:rgb(0,0,0)"> </span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!5 = metadata !{metadata !"kernel_arg_base_type", metadata !"int"} <span style="color:rgb(0,0,0)"> </span><span style="background-color:rgb(48,48,48)"> </span><span style="color:rgb(0,0,0)"> </span></p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!6 = metadata !{i32 1, i32 2} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;background-color:rgb(18,18,18)"><span style="color:rgb(208,208,208)">!7 = metadata !{} </span> <span style="color:rgb(208,208,208);background-color:rgb(48,48,48)"> </span> </p>
<p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)">!8 = metadata !{metadata !"-cl-std=CL1.2"} <span style="background-color:rgb(255,255,255);color:rgb(34,34,34)"> </span></p><p style="margin:0px;font-size:11px;font-family:Menlo;color:rgb(208,208,208);background-color:rgb(18,18,18)"><span style="background-color:rgb(255,255,255);color:rgb(34,34,34)"> </span><br></p></div></div>