generation of SPIR from OpenCL

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.

If I take this simple file

//hello.cl

__kernel void hello(int x) {

}

then run

clang -x cl -fno-builtin -target spir -c -emit-llvm hello.cl

it generates a bc file, but it doesn’t have the additional attributes that describe the function arguments.

Here’s the comparison with the output from Khronos group’s version of clang (https://github.com/KhronosGroup/SPIR)

########################## clang 3.6

; ModuleID = ‘hello.bc’

target datalayout = “e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024”

target triple = “spir”

; Function Attrs: nounwind readnone

define void @hello(i32 %x) #0 {

entry:

ret void

}

attributes #0 = { nounwind readnone “less-precise-fpmad”=“false” “no-frame-pointer-elim”=“true” “no-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” }

!opencl.kernels = !{!0}

!llvm.ident = !{!1}

!0 = metadata !{void (i32)* @hello}

!1 = metadata !{metadata !"clang version 3.6.0 "}

########################## Khronos

; ModuleID = ‘hello.bc’

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-unknown-unknown”

; Function Attrs: nounwind readnone

define spir_kernel void @hello(i32 %x) #0 {

entry:

ret void

}

attributes #0 = { nounwind readnone }

!opencl.kernels = !{!0}

!opencl.enable.FP_CONTRACT = !{}

!opencl.spir.version = !{!6}

!opencl.ocl.version = !{!6}

!opencl.used.extensions = !{!7}

!opencl.used.optional.core.features = !{!7}

!opencl.compiler.options = !{!8}

!0 = metadata !{void (i32)* @hello, metadata !1, metadata !2, metadata !3, metadata !4, metadata !5}

!1 = metadata !{metadata !“kernel_arg_addr_space”, i32 0}

!2 = metadata !{metadata !“kernel_arg_access_qual”, metadata !“none”}

!3 = metadata !{metadata !“kernel_arg_type”, metadata !“int”}

!4 = metadata !{metadata !“kernel_arg_type_qual”, metadata !“”}

!5 = metadata !{metadata !“kernel_arg_base_type”, metadata !“int”}

!6 = metadata !{i32 1, i32 2}

!7 = metadata !{}

!8 = metadata !{metadata !“-cl-std=CL1.2”}