[Openmp-dev] Experimental SPIR-V back-end using OpenCL 2.1

[Jonas Hahnfeld via Openmp-dev]

Hi Daniel,

interesting work!

Two question:
  - With the latest commit, you perform a strncmp with "OpenCL 2.0 AMD". 
Does this mean that all other OpenCL implementations are effectively 
blocked out?
  - Is this the proprietary AMD OpenCL SDK or the "new" ROCm stack? 
https://rocm.github.io/

One remark: For __tgt_rtl_is_valid_binary: Does SPIR-V have its own 
machine id? That's how the CUDA plugin detects compatible binaries...

Cheers,
Jonas

Am 2017-10-20 11:24, schrieb Daniel Schürmann via Openmp-dev:
> Thanks for all the help!
> 
> These benchmarks, especially lulesh, showed me that my approach of
> setting the address spaces manually according to their scope doesn't
> work(*) and that I have to use the generic address space
> like the nvptx backend does.
> 
> Now with this much more robust version, I decided to make my project
> public:
> 
> My clang fork is available at
> https://github.com/daniel-schuermann/clang
> 
> and the openmp runtime fork at
> https://github.com/daniel-schuermann/openmp
> 
> and the necessary llvm fork (for generating SPIR-V) is from
> https://github.com/thewilsonator/llvm/tree/compute
> 
> All 3 of them might need a pull from upstream as they are not always
> synced.
> 
> libomptarget-spir needs an OpenCL runtime which supports SPIRV kernels
> (with OpenCL 2.1 headers).
> Unfortunately, the Intel OpenCL runtime started to segfault with the
> change to generic address space.
> Therefore, the only working OpenCL runtime I know of is AMDGPU-Pro.
> 
> The follwing pragmas should work for now:
> #target (enter/exit data)
> #teams
> #distribute / parallel for
> #master
> #barrier
> 
> as well as the clauses:
> map, shared, private, firstprivate, lastprivate, schedule
> 
> A device runtime is not planned for the moment (I also don't know, how
> to compile and link OpenCL sources into libomptarget), but the
> generated code works independently from runtime functions.
> 
> For those more adventurous, you can (try to) build your sources with
> -fopenmp -fopenmp-targets=spir64-unknown-unknown
> For C++, I recommend to add -fno-exceptions and -O0 as there seem to
> be optimizer passes enabled
> which don't work for spir.
> 
> I would be very thankful for some feedback (I hope, it won't get too
> depressing).
> Although I am not able to accept pull requests at the moment for legal
> purposes,
> I would welcome any hint to make the implementation more robust and
> complete
> as well as statements about (not) working programs.
> (no comments on code style please, this can be fixed later °°)
> 
> Kind regards,
> Daniel
> 
> (*) While this is a valid program snipped,
> 
> #pragma omp target map(to:a[0:n])
> { int * b = a; }
> 
> it doesn't work if |a| is a pointer to addrSpace(1) (cl_global) and
> |b| gets allocated as pointer to cl_private.