Issue generating SPIR-V from "BC from generated OpenCL C++ kernel"..

Hi, seeing IWOCL keynote slides, interesting is the slide 20: C++ for OpenCL in Clang project it has a sample: https://godbolt.org/z/nGvxAC and a I see it uses SPIR target instead of SPIR-V I wanted to test if can have success generating SPIR-V output with SPIRV-LLVM-translator from that sample: so I name this kernel clcpp.cl:

// Need to declare the prototype for get_global_id().
int get_global_id( int dim );

template<class T>
T add( T x, T y )
{
    return x + y;
}

__kernel void test( __global float* a, __global float* b)
{
    // Need to use unsigned; uint doesn't work.
    auto index = get_global_id(0);
    a[ index ] = add( b[ index ], b[ index + 1 ] );
}

and get latest clang-9 nightly form apt.llvm.org and latest SPIRV-LLVM-Translator-dev build available here (version.txt mentions commit: ece29378764c7b16c273276e62a61efad801c633).. then using:

clang++-9 -cl-std=c++ clpp.cl -emit-llvm -target spir -c -o clppspir.bc

as seems clang++-9 from apt.llvm.org doesn't accept -target spir-v anyway seems llvm-spirv can convert BC's using SPIR target also.. but then using:

./llvm-spirv clpp.bc clppspir.spv

I get:

llvm-spirv: Too many positional arguments specified!
Can specify at most 1 positional arguments: See: ./llvm-spirv --help

So questions are:

any additional flags needed to pass to llvm-spirv or clang++-9 to fix this issue? SPIRV-LLVM translator is ready to translate C++ OCL kernels like this? finally also in slide I see:

Offline compilation into SPIR-V or device binary
- Generates SPIR-V 1.0 for most features
- Uses SPIR-V 1.2 where necessary

so question is: can share when SPIR-V 1.2 instead of SPIR-V 1.0 is needed? also some additional argument to llmv-spirv or clang compilation steps is needed in such cases? I say because Intel Neo driver already has SPIR-V 1.2 support and wanted to test working support!

I attach bitcode in text form ( llvm-dis-9 <clppspir.bc) in case you want to explore also:

; ModuleID = '<stdin>'
source_filename = "clpp.cl"
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: convergent nounwind
define dso_local spir_kernel void @test(float addrspace(1)* nocapture, float addrspace(1)* nocapture readonly) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
  %3 = tail call spir_func i32 @_Z13get_global_idi(i32 0) #2
  %4 = getelementptr inbounds float, float addrspace(1)* %1, i32 %3
  %5 = load float, float addrspace(1)* %4, align 4, !tbaa !8
  %6 = add nsw i32 %3, 1
  %7 = getelementptr inbounds float, float addrspace(1)* %1, i32 %6
  %8 = load float, float addrspace(1)* %7, align 4, !tbaa !8
  %9 = fadd float %5, %8
  %10 = getelementptr inbounds float, float addrspace(1)* %0, i32 %3
  store float %9, float addrspace(1)* %10, align 4, !tbaa !8
  ret void
}

; Function Attrs: convergent
declare dso_local spir_func i32 @_Z13get_global_idi(i32) local_unnamed_addr #1

attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "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" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #2 = { convergent nounwind }

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

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 0, i32 0}
!2 = !{i32 0, i32 2}
!3 = !{!"clang version 9.0.0-svn361218-1~exp1+0~20190521040613.149~1.gbp9d9550 (trunk)"}
!4 = !{i32 1, i32 1}
!5 = !{!"none", !"none"}
!6 = !{!"float*", !"float*"}
!7 = !{!"", !""}
!8 = !{!9, !9, i64 0}
!9 = !{!"float", !10, i64 0}
!10 = !{!"omnipotent char", !11, i64 0}
!11 = !{!"Simple C++ TBAA"}

Hi @oscarbg,

./llvm-spirv clpp.bc clppspir.spv

You need to use:

./llvm-spirv clcpp.bc -o clppspir.spv
# or
./llvm-spirv clcpp.bc # to get clcpp.spv file

Offline compilation into SPIR-V or device binary

Generates SPIR-V 1.0 for most features

Uses SPIR-V 1.2 where necessary

From personal experience, I would say it generates not only SPIR-V 1.0 in some cases. For example, if you have functions which accept arguments by reference in your OpenCL C++ code, SPIRV-LLVM-Translator will generate some additional decorations and will generate SPIR-V 1.1 binary.

can share when SPIR-V 1.2 instead of SPIR-V 1.0 is needed?

I'm not sure that there is some well-documented list of OpenCL C++ features that requires SPIR-V 1.2 instead of SPIR-V 1.0, but at least you can take a look at spec changelog to find some information.

also some additional argument to llmv-spirv or clang compilation steps is needed in such cases?

If I understand correctly, SPIRV-LLVM-Translator automatically selects resulting version based on features used in input file.

Unfortunately, there are no options to force min/max or required version of the resulting SPIR-V file. And that leads to a "hacks" like this one - this options allows you to generate SPIR-V 1.0 even if source code contains functions which accept arguments by reference.

KhronosGroup / SPIRV-LLVM-Translator

Issue generating SPIR-V from "BC from generated OpenCL C++ kernel".. #201