Code Monkey home page Code Monkey logo

spirv-llvm-translator's Introduction

LLVM/SPIR-V Bi-Directional Translator

Out-of-tree build & tests In-tree build & tests

This repository contains source code for the LLVM/SPIR-V Bi-Directional Translator, a library and tool for translation between LLVM IR and SPIR-V. This project currently only supports the OpenCL/compute "flavour" of SPIR-V: it consumes and produces SPIR-V modules that declare the Kernel capability.

The LLVM/SPIR-V Bi-Directional Translator is open source software. You may freely distribute it under the terms of the license agreement found in LICENSE.txt.

Directory Structure

The files/directories related to the translator:

Build Instructions

The main branch of this repo is aimed to be buildable with the latest LLVM main revision.

Build with pre-installed LLVM

The translator can be built with the latest(nightly) package of LLVM. For Ubuntu and Debian systems LLVM provides repositories with nightly builds at http://apt.llvm.org/. For example the latest package for Ubuntu 16.04 can be installed with the following commands:

wget -O - https://apt.llvm.org/llvm-snapshot.gpg.key | sudo apt-key add -
sudo add-apt-repository "deb http://apt.llvm.org/xenial/ llvm-toolchain-xenial main"
sudo apt-get update
sudo apt-get install llvm-19-dev llvm-19-tools clang-19 libclang-19-dev

The installed version of LLVM will be used by default for out-of-tree build of the translator.

git clone https://github.com/KhronosGroup/SPIRV-LLVM-Translator.git
mkdir SPIRV-LLVM-Translator/build && cd SPIRV-LLVM-Translator/build
cmake ..
make llvm-spirv -j`nproc`

Build with pre-built LLVM

If you have a custom build (based on the latest version) of LLVM libraries you can link the translator against it.

git clone https://github.com/KhronosGroup/SPIRV-LLVM-Translator.git
mkdir SPIRV-LLVM-Translator/build && cd SPIRV-LLVM-Translator/build
cmake .. -DLLVM_DIR=<llvm_build_dir>/lib/cmake/llvm/
make llvm-spirv -j`nproc`

If the translator is used as part of another CMake project, you will need to define LLVM_SPIRV_BUILD_EXTERNAL:

cmake .. -DLLVM_DIR=<llvm_build_dir>/lib/cmake/llvm/ -DLLVM_SPIRV_BUILD_EXTERNAL=YES

Where llvm_build_dir is the LLVM build directory.

LLVM in-tree build

The translator can be built as a regular LLVM subproject. To do that you need to clone it into the llvm/projects or llvm/tools directory.

git clone https://github.com/llvm/llvm-project.git
cd llvm-project/llvm/projects
git clone https://github.com/KhronosGroup/SPIRV-LLVM-Translator.git

Run (or re-run) cmake as usual for LLVM. After that you should have llvm-spirv and check-llvm-spirv targets available.

mkdir llvm-project/build && cd llvm-project/build
cmake ../llvm -DLLVM_ENABLE_PROJECTS="clang"
make llvm-spirv -j`nproc`

Note on enabling the clang project: there are tests in the translator that depend on clang binary, which makes clang a required dependency (search for LLVM_SPIRV_TEST_DEPS in test/CMakeLists.txt) for check-llvm-spirv target.

Building clang from sources takes time and resources and it can be avoided:

  • if you are not interested in launching unit-tests for the translator after build, you can disable generation of test targets by passing -DLLVM_SPIRV_INCLUDE_TESTS=OFF option.
  • if you are interested in launching unit-tests, but don't want to build clang you can pass -DSPIRV_SKIP_CLANG_BUILD cmake option to avoid adding clang as dependency for check-llvm-spirv target. However, LIT will search for clang binary when tests are launched and it should be available at this point.
  • building and testing completely without clang is not supported at the moment, see KhronosGroup/SPIRV-LLVM-Translator#477 to track progress, discuss and contribute.

Build with SPIRV-Tools

The translator can use SPIRV-Tools to generate assembly with widely adopted syntax. If SPIRV-Tools have been installed prior to the build it will be detected and used automatically. However it is also possible to enable use of SPIRV-Tools from a custom location using the following instructions:

  1. Checkout, build and install SPIRV-Tools using the following instructions. Example using CMake with Ninja:
cmake -G Ninja <SPIRV-Tools source location> -DCMAKE_INSTALL_PREFIX=<SPIRV-Tools installation location>
ninja install
  1. Point pkg-config to the SPIR-V tools installation when configuring the translator by setting PKG_CONFIG_PATH=<SPIRV-Tools installation location>/lib/pkgconfig/ variable before the cmake line invocation. Example:
PKG_CONFIG_PATH=<SPIRV-Tools installation location>/lib/pkgconfig/ cmake <other options>

To verify the SPIR-V Tools integration in the translator build, run the following line

llvm-spirv --spirv-tools-dis input.bc -o -

The output should be printed in the standard assembly syntax.

Configuring SPIR-V Headers

The translator build is dependent on the official Khronos header file spirv.hpp that maps SPIR-V extensions, decorations, instructions, etc. onto numeric tokens. The official header version is available at KhronosGroup/SPIRV-Headers. There are several options for accessing the header file:

  • By default, the header file repository will be downloaded from Khronos Group GitHub and put into <build_dir>/SPIRV-Headers.
  • If you are building the translator in-tree, you can manually download the SPIR-V Headers repo into llvm/projects - this location will be automatically picked up by the LLVM build scripts. Make sure the folder retains its default naming in that of SPIRV-Headers.
  • Any build type can also use an external installation of SPIR-V Headers - if you have the headers downloaded somewhere in your system and want to use that version, simply extend your CMake command with -DLLVM_EXTERNAL_PROJECTS="SPIRV-Headers" -DLLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR=</path/to/headers_dir>.

Test instructions

All tests related to the translator are placed in the test directory. A number of the tests require spirv-as (part of SPIR-V Tools) to run, but the remainder of the tests can still be run without this. Optionally the tests can make use of spirv-val (part of SPIRV-Tools) in order to validate the generated SPIR-V against the official SPIR-V specification.

In case tests are failing due to SPIRV-Tools not supporting certain SPIR-V features, please get an updated package. The PKG_CONFIG_PATH environmental variable can be used to let cmake point to a custom installation.

Execute the following command inside the build directory to run translator tests:

make test

This requires that the -DLLVM_SPIRV_INCLUDE_TESTS=ON argument is passed to CMake during the build step. Additionally, -DLLVM_EXTERNAL_LIT="/usr/lib/llvm-19/build/utils/lit/lit.py" is needed when building with a pre-installed version of LLVM.

The translator test suite can be disabled by passing -DLLVM_SPIRV_INCLUDE_TESTS=OFF to CMake.

Run Instructions for llvm-spirv

To translate between LLVM IR and SPIR-V:

  1. Execute the following command to translate input.bc to input.spv

    llvm-spirv input.bc
    
  2. Execute the following command to translate input.spv to input.bc

    llvm-spirv -r input.spv
    

    Recommended options:

    • -spirv-target-env - to specify target version of OpenCL builtins to translate to (default CL1.2)
  3. Other options accepted by llvm-spirv

    • -o file_name - to specify output name
    • -spirv-debug - output debugging information
    • -spirv-text - read/write SPIR-V in an internal textual format for debugging purpose. The textual format is not defined by SPIR-V spec.
    • --spirv-tools-dis - print SPIR-V assembly in SPIRV-Tools format. Only available on builds with SPIRV-Tools.
    • -help - to see full list of options

Translation from LLVM IR to SPIR-V and then back to LLVM IR is not guaranteed to produce the original LLVM IR. In particular, LLVM intrinsic call instructions may get replaced by function calls to OpenCL builtins and metadata may be dropped.

Handling SPIR-V versions generated by the translator

There is one option to control the behavior of the translator with respect to the version of the SPIR-V file which is being generated/consumed.

  • -spirv-max-version= - this option allows restricting the SPIRV-LLVM-Translator not to generate a SPIR-V with a version which is higher than the one specified via this option.

    If the -r option was also specified, the SPIRV-LLVM-Translator will reject the input file and emit an error if the SPIR-V version in it is higher than one specified via this option.

Allowed values are 1.0, 1.1, 1.2, 1.3, 1.4, and 1.5.

More information can be found in SPIR-V versions and extensions handling

Handling SPIR-V extensions generated by the translator

By default, during SPIR-V generation, the translator doesn't use any extensions. However, during SPIR-V consumption, the translator accepts input files that use any known extensions.

If certain extensions are required to be enabled or disabled, the following command line option can be used:

  • --spirv-ext= - this options allows controlling which extensions are allowed/disallowed

Valid value for this option is comma-separated list of extension names prefixed with + or - - plus means allow to use extension, minus means disallow to use extension. There is one more special value which can be used as extension name in this option: all - it affects all extension which are known to the translator.

If --spirv-ext contains the name of an extension which is not known for the translator, it will emit an error.

More information can be found in SPIR-V versions and extensions handling

Branching strategy

Code on the main branch in this repository is intended to be compatible with the main branch of the llvm project. That is, for an OpenCL kernel compiled to llvm bitcode by the latest git revision of Clang it should be possible to translate it to SPIR-V with the llvm-spirv tool.

All new development should be done on the main branch.

To have versions compatible with released versions of LLVM and Clang, corresponding tags are available in this repository. For example, to build the translator with LLVM 7.0.0 one should use the v7.0.0-1 tag. The 7.x releases are maintained on the llvm_release_70 branch. As a general rule, commits from the main branch may be backported to the release branches as long as they do not depend on features from a later LLVM/Clang release and there are no objections from the maintainer(s). There is no guarantee that older release branches are proactively kept up to date with main, but you can request specific commits on older release branches by creating a pull request or raising an issue on GitHub.

Releasing strategy

As mentioned earlier there are branches llvm_release_* that get backported changes. Those changes if exists are released automatically by github CI on monthly basis in a format <llvm_major>.<llvm_minor>.<latest patch +1>.

spirv-llvm-translator's People

Contributors

agindinson avatar alexeysachkov avatar alexeysotkin avatar aratajew avatar asudarsa avatar bader avatar bashbaug avatar bwlodarcz avatar dmitrybushev avatar fznamznon avatar jcranmer-intel avatar karolherbst avatar kkyzylova avatar kornevnikita avatar ldudziak avatar lu-john avatar mlychkov avatar mochalovaan avatar mrsidims avatar nikitarudenkointel avatar pauzinl avatar pierremoreau avatar sarnex avatar scottp101 avatar stuartdbrady avatar svenvh avatar vmaksimo avatar vyacheslavlevytskyy avatar wenju-he avatar yxsamliu avatar

Stargazers

 avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar

Watchers

 avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar

spirv-llvm-translator's Issues

Translation of barrier bulitins from OpenCL to SPIR-V

OpenCL barrier and work_group_barrier bulitins have argument of cl_mem_fence_flags type which specifies local or global address space. While in SPIR-V OpControlBarrier has ‘Memory Semantics’ operand which specifies both memory-order constraints, and on what storage classes those constraints apply to. OpenCL address space maps fine to SPIRV storage class, but what about memory-order? According to the description here https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/work_group_barrier.html it seems like OpenCL barrier implies acquire-release memory-order semantics, doesn’t it?
If yes, should we always set AcquireRelease bit in the ‘Memory Semantics’ operand translating OpenCL [working_group_]barrier to OpControlBarrier in SPIR-V ?

handling extensions and spirv versions

I haven't dug too deeply on this yet, but I was wondering what would the plan be for handling SPIR-V extensions.

Say we added a new CL SPIR-V extension for a new instruction and I wanted to generate it or we wanted to generate SPIR-V compatible with spirv 1.1 or 1.3 etc.

wrt triples would this need to be encoded in the triple?

Apologies if this is a bit vague :-)

Mac OS X: composite_construct_vector.spt fails

detected with travis-ci

Script:
--
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-spirv /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/composite_construct_vector.spt -to-binary -o /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/composite_construct_vector.spt.tmp.spv
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-spirv -r /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/composite_construct_vector.spt.tmp.spv -o /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/composite_construct_vector.spt.tmp.bc
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-dis < /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/composite_construct_vector.spt.tmp.bc | /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/FileCheck /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/composite_construct_vector.spt --check-prefix=CHECK-LLVM
--
Exit Code: 1
Command Output (stderr):
--
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/composite_construct_vector.spt:48:15: error: expected string not found in input
; CHECK-LLVM: %[[vector:[0-9]+]] = getelementptr inbounds <4 x i32>, <4 x i32> addrspace(1)* %in, i64 %{{[0-9]*}}
              ^
<stdin>:1:1: note: scanning from here
; ModuleID = '<stdin>'
^
--

Evaluate support of optimised IR translation

Current use case of translator assumes that non-optimised IR is given.

A number of questions that were asked recently are as follows:

  • Do we have any idea what issues do we have with optimised IR? Would it be worth running a some tests and collect the issues we find?
  • Is this a valuable use case?
  • Would support for this be easier to implement in a LLVM backend rather then translation format?

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: ece2937)..
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"}

Build Failure since llvm@328315

SPIRV-LLVM-Translator/lib/SPIRV/SPIRVLowerMemmove.cpp:82:20: error: ‘class llvm::MemMoveInst’ has no member named ‘getAlignment’

Mac OS X: copy_object.spt fails

on travis-ci

Script:
--
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-spirv /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/copy_object.spt -to-binary -o /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/copy_object.spt.tmp.spv
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-spirv -r /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/copy_object.spt.tmp.spv -o /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/copy_object.spt.tmp.bc
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/llvm-dis < /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/tools/llvm-spirv/test/test_output/Output/copy_object.spt.tmp.bc | /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/build/bin/FileCheck /Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/copy_object.spt --check-prefix=CHECK-LLVM
--
Exit Code: 1
Command Output (stderr):
--
/Users/travis/build/karolherbst/SPIRV-LLVM-Translator/llvm/tools/llvm-spirv/test/copy_object.spt:48:15: error: expected string not found in input
; CHECK-LLVM: %[[char:[0-9]+]] = alloca i8
              ^
<stdin>:1:1: note: scanning from here
; ModuleID = '<stdin>'
^
<stdin>:1:3: note: possible intended match here
; ModuleID = '<stdin>'
  ^
--

SPIRV storage classes Private and Function

We have two storage classes defined in SPIR-V which seem to map to the concept of Private address space in OpenCL, StorageClass::Private and StorageClass::Function. Private is shader/ vulkan specific and represents thread local storage, accessible from any function within the invocation, so at program scope. The question is, if we want a true bidrectional translator, we'd need to maintain this, shader_private, as a separate address space within LLVM, which may result in some passes not optimising as efficiently in face of it.

My question is, is it reasonable to map StorageClass::Private to the same address space as CL private, (0), if we lose bidirectionality in this case, or should we keep it as two and deal with having to specialise the optimisation passes?

Remove _SPIRV_LLVM_API

When setting SPIRV_USE_LLVM_API to OFF, the code does not compile due to some parts of the code not using that macro and always using the LLVM API (see further down). The code could be fixed to ensure that the entrypoints properly convert the LLVM streams to the std ones, if needed, but is this option really that useful/used?

/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/SPIRVWriter.cpp:1858:6: error: invalid operands to binary expression ('llvm::raw_ostream' and 'SPIRV::SPIRVModule')
  OS << *BM;
  ~~ ^  ~~~
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:201:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const void *' for 1st argument; take the address of the argument with &
  raw_ostream &operator<<(const void *P);
               ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVModule.h:325:24: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'std::ostream &' (aka 'basic_ostream<char> &') for 1st argument
  friend spv_ostream & operator<<(spv_ostream &O, SPIRVModule& M);
                       ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:150:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
operator<<(const SPIRVEncoder& O, T V) {
^
/home/pmoreau/projects/llvm/include/llvm/ADT/APInt.h:2018:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::APInt' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &OS, const APInt &I) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/IR/Value.h:670:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Value' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &OS, const Value &V) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/IR/Type.h:455:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Type' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &OS, const Type &T) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/ADT/Twine.h:533:23: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Twine' for 2nd argument
  inline raw_ostream &operator<<(raw_ostream &OS, const Twine &RHS) {
                      ^
/home/pmoreau/projects/llvm/include/llvm/IR/Metadata.h:160:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Metadata' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &OS, const Metadata &MD) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/IR/Comdat.h:64:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Comdat' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &OS, const Comdat &C) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/IR/Module.h:859:21: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::Module' for 2nd argument
inline raw_ostream &operator<<(raw_ostream &O, const Module &M) {
                    ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:145:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'char' for 1st argument
  raw_ostream &operator<<(char C) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:152:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'unsigned char' for 1st argument
  raw_ostream &operator<<(unsigned char C) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:159:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'signed char' for 1st argument
  raw_ostream &operator<<(signed char C) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:166:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'llvm::StringRef' for 1st argument
  raw_ostream &operator<<(StringRef Str) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:181:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const char *' for 1st argument
  raw_ostream &operator<<(const char *Str) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:188:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const std::string' (aka 'const basic_string<char>') for 1st argument
  raw_ostream &operator<<(const std::string &Str) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:193:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const SmallVectorImpl<char>' for 1st argument
  raw_ostream &operator<<(const SmallVectorImpl<char> &Str) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:197:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'unsigned long' for 1st argument
  raw_ostream &operator<<(unsigned long N);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:198:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'long' for 1st argument
  raw_ostream &operator<<(long N);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:199:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'unsigned long long' for 1st argument
  raw_ostream &operator<<(unsigned long long N);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:200:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'long long' for 1st argument
  raw_ostream &operator<<(long long N);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:203:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'unsigned int' for 1st argument
  raw_ostream &operator<<(unsigned int N) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:207:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'int' for 1st argument
  raw_ostream &operator<<(int N) {
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:211:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'double' for 1st argument
  raw_ostream &operator<<(double N);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:228:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::format_object_base' for 1st argument
  raw_ostream &operator<<(const format_object_base &Fmt);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:231:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::FormattedString' for 1st argument
  raw_ostream &operator<<(const FormattedString &);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:234:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::FormattedNumber' for 1st argument
  raw_ostream &operator<<(const FormattedNumber &);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:237:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::formatv_object_base' for 1st argument
  raw_ostream &operator<<(const formatv_object_base &);
               ^
/home/pmoreau/projects/llvm/include/llvm/Support/raw_ostream.h:240:16: note: candidate function not viable: no known conversion from 'SPIRV::SPIRVModule' to 'const llvm::FormattedBytes' for 1st argument
  raw_ostream &operator<<(const FormattedBytes &);
               ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:190:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
SPIRV_DEC_ENCDEC(Op)
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:185:21: note: expanded from macro 'SPIRV_DEC_ENCDEC'
const SPIRVEncoder& \
                    ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:191:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
SPIRV_DEC_ENCDEC(Capability)
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:185:21: note: expanded from macro 'SPIRV_DEC_ENCDEC'
const SPIRVEncoder& \
                    ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:192:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
SPIRV_DEC_ENCDEC(Decoration)
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:185:21: note: expanded from macro 'SPIRV_DEC_ENCDEC'
const SPIRVEncoder& \
                    ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:193:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
SPIRV_DEC_ENCDEC(OCLExtOpKind)
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:185:21: note: expanded from macro 'SPIRV_DEC_ENCDEC'
const SPIRVEncoder& \
                    ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:194:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
SPIRV_DEC_ENCDEC(LinkageType)
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:185:21: note: expanded from macro 'SPIRV_DEC_ENCDEC'
const SPIRVEncoder& \
                    ^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:197:1: note: candidate function not viable: no known conversion from 'llvm::raw_ostream' to 'const SPIRV::SPIRVEncoder' for 1st argument
operator<<(const SPIRVEncoder&O, const std::string& Str);
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:164:1: note: candidate template ignored: could not match 'T *' against 'SPIRV::SPIRVModule'
operator<<(const SPIRVEncoder& O, T* P) {
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:170:1: note: candidate template ignored: could not match 'vector<type-parameter-0-0, allocator<type-parameter-0-0> >' against 'SPIRV::SPIRVModule'
operator<<(const SPIRVEncoder& O, const std::vector<T>& V) {
^
/home/pmoreau/projects/llvm/tools/SPIRV-LLVM-Translator/lib/SPIRV/libSPIRV/SPIRVStream.h:178:1: note: candidate template ignored: could not match 'pair<type-parameter-0-0, type-parameter-0-0>' against 'SPIRV::SPIRVModule'
operator<<(const SPIRVEncoder& Encoder, const std::pair<IterTy,IterTy> &Range) {
^
/home/pmoreau/projects/llvm/include/llvm/Support/ScaledNumber.h:731:23: note: candidate template ignored: could not match 'ScaledNumber<type-parameter-0-0>' against 'llvm::raw_ostream'
ScaledNumber<DigitsT> operator<<(const ScaledNumber<DigitsT> &L,
                      ^
/home/pmoreau/projects/llvm/include/llvm/Support/ScaledNumber.h:743:14: note: candidate template ignored: could not match 'ScaledNumber<type-parameter-0-0>' against 'SPIRV::SPIRVModule'
raw_ostream &operator<<(raw_ostream &OS, const ScaledNumber<DigitsT> &X) {
             ^
19 warnings and 1 error generated.
ninja: build stopped: subcommand failed.

Interface variables aren't declared in OpEntryPoint

Interface variables in the Input and Output storage classes are required to be declared in OpEntryPoint as per section 2.6 of the SPIR-V spec.

The following module allows to reproduce the issue:

; ModuleID = 'simple.cl'
source_filename = "simple.cl"
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: convergent nounwind writeonly
define spir_kernel void @test(i32 addrspace(1)* nocapture %out) 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 {
entry:
  %call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2
  %conv = trunc i64 %call to i32
  %idxprom = and i64 %call, 4294967295
  %arrayidx = getelementptr inbounds i32, i32 addrspace(1)* %out, i64 %idxprom
  store i32 %conv, i32 addrspace(1)* %arrayidx, align 4, !tbaa !8
  ret void
}

; Function Attrs: convergent nounwind readnone
declare spir_func i64 @_Z13get_global_idj(i32) local_unnamed_addr #1

attributes #0 = { convergent nounwind writeonly "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"="false" "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 nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "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 readnone }

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

!0 = !{i32 1, !"wchar_size", i32 4}
!1 = !{i32 1, i32 0}
!2 = !{i32 1, i32 2}
!3 = !{!"clang version 9.0.0 (/path/to/clang a364d599ab76802eff58ed6dfdcbc8fc104405ca)"}
!4 = !{i32 1}
!5 = !{!"none"}
!6 = !{!"int*"}
!7 = !{!""}
!8 = !{!9, !9, i64 0}
!9 = !{!"int", !10, i64 0}
!10 = !{!"omnipotent char", !11, i64 0}
!11 = !{!"Simple C/C++ TBAA"}

Should we disallow passing blocks as a function parameter ?

The subject seems desirable from the compiler point of view to avoid function pointers. And that requires changes in the OpenCL spec, i.e. adding restrictions to block usage. But that may break already existing code with such usage of blocks.

memory alignment assertion failure on intel opencl compute runtime (neo)

We get this assertion failure:
/home/dvrogozh/git/github/SPIRV-LLVM-Translator/lib/SPIRV/SPIRVWriter.cpp:1174: SPIRV::SPIRVValue* SPIRV::LLVMToSPIRV::transIntrinsicInst(llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*): Assertion cast(II)->getSourceAlignment() == cast(II)->getDestAlignment() && "Alignment mismatch!"' failed.
`
on one of our projects which uses https://github.com/intel/compute-runtime. The issue can be reproduced via ocloc compiler (internal target built within neo build procedure):

# cat test-alignment.cl
struct DATA {
  float data[3];
};

struct FOO
{
  global char* barPtr;
};

struct BAR
{
  unsigned long offset;
  struct DATA payload;
};

void kernel test(global __const struct FOO* foo)
{
  const uint i = get_global_id(0);
  global struct BAR* bar = (global struct BAR*) foo[i].barPtr;

  struct DATA data = bar->payload;
}

# ./build/bin/ocloc -file test-alignment.cl -device skl
ocloc: /home/dvrogozh/git/github/SPIRV-LLVM-Translator/lib/SPIRV/SPIRVWriter.cpp:1174: SPIRV::SPIRVValue* SPIRV::LLVMToSPIRV::transIntrinsicInst(llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*): Assertion `cast<MemCpyInst>(II)->getSourceAlignment() == cast<MemCpyInst>(II)->getDestAlignment() && "Alignment mismatch!"' failed.

We use https://github.com/KhronosGroup/SPIRV-LLVM-Translator/tree/llvm_release_80 branch.

LLVM version support policy

The code seems to currently only build against LLVM 7.0, which hasn’t been released yet. I do not know whether a policy has already been decided on which versions of LLVM to support,but I think we should at the very least support the latest release version. And I would definitely agree with tracking LLVM HEAD as closely as possible.

Obsolete README

The README does not seem to have been updated yet to reflect the new architecture changes.

OCL20ToSPIRV::visitCallScalToVec seems to produce incorrect step code

define linkonce_odr dso_local spir_func <16 x double> @_Z10smoothstepffDv16_d(float, float, <16 x double>) local_unnamed_addr #0 {
%4 = shufflevector <16 x double> %2, <16 x double> undef, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7>
%5 = tail call spir_func <8 x double> @_Z10smoothstepffDv8_d(float %0, float %1, <8 x double> %4) #5

is the LLVM IR which produces:
%998 = OpFunction %v16double None %997
%999 = OpFunctionParameter %float
%1000 = OpFunctionParameter %float
%1001 = OpFunctionParameter %v16double
%1002 = OpLabel
%1003 = OpVectorShuffle %v8double %1001 %757 0 1 2 3 4 5 6 7
%1004 = OpCompositeInsert %v8double %999 %762 0
%1005 = OpVectorShuffle %v8double %1004 %762 0 0 0 0 0 0 0 0
%1006 = OpCompositeInsert %v8double %1000 %762 0
%1007 = OpVectorShuffle %v8double %1006 %762 0 0 0 0 0 0 0 0
%1008 = OpExtInst %v8double %1 smoothstep %1005 %1007 %1003

This fails validation as it tries to insert 999 a float into a v8double.

Lack of unit-testing for passes which are used in SPIRVWriter

As I can see, there are bunch of passes which perform some preprocessing of LLVM IR before translation into SPIR-V:

  PassMgr.add(createTransOCLMD());
  PassMgr.add(createOCL21ToSPIRV());
  PassMgr.add(createSPIRVLowerSPIRBlocks());
  PassMgr.add(createOCLTypeToSPIRV());
  PassMgr.add(createSPIRVLowerOCLBlocks());
  PassMgr.add(createOCL20ToSPIRV());
  PassMgr.add(createSPIRVRegularizeLLVM());
  PassMgr.add(createSPIRVLowerConstExpr());
  PassMgr.add(createSPIRVLowerBool());
  PassMgr.add(createSPIRVLowerMemmove());

But we don't have possibility to unit-test them: we cannot launch opt -ocl21-to-spirv to test behavior of a certain pass.

What do you think? Is it important or not? How can we achieve that?

Unify the version numbers used throughout the project

  • The project is exposed as version 0.2.1 via pkg-config (see in CMakeLists.txt);
  • New releases are tagged as $LLVM_VERSION-$iteration; latest release is 8.0.0-1
  • llvm-spirv --version returns the LLVM version (same as the release tags, but without the -1 suffix);
  • The version of the tool stored in the SPIR-V binary (in the lower 16 bits of word 2 of the SPIR-V header), is stuck to 14. Maybe this one doesn’t need to match the other versions though.

Separate SPIR-V parsing and translation

Currently the translator parses and translates SPIR-V to LLVM in one step from the API perspective. It would be useful to separate out the parsing and translation into two separate API calls. This would allow the correct implementation of clCreateProgramWithIL as this API function needs to be able to return CL_INVALID_VALUE if the passed IL is not well-formed. Currently passing malformed IL causes a crash in the parser or translation step. This also allows a separate call to clCompileProgram/clBuildProgram to perform the translation at that stage, giving an appropriate error, CL_INVALID_OPERATION, if the program has not been created from source or IL and also if the target device doesn't support a capability requested in the SPIR-V binary, see https://www.khronos.org/registry/OpenCL/sdk/2.1/docs/man/xhtml/clCompileProgram.html.
clCompileProgram/clBuildProgram also supports the passing of compiler options, which we need to think about how to support at this stage, and we'd want to be able to extract the build log using clGetProgramBuildInfo which should be populated by the last call to clCompileProgram or clBuildProgram.

The additional benefit of doing this separation would be to allow us to reuse the common SPIR-V parsing and in-memory representation found in the SPIRV-Tools repository, cutting down on the amount of duplication and effort required to support newer versions of SPIR-V.

Specialization constant representation in LLVM

How do we represent, and carry forward, specialization constants in LLVM during the translation step?

Tim Renouf raised this issue in an email.

"Currently the AMD code passes the constant values into readSpirv() and
does the specialization inside the spir-v reader.

That's the most pragmatic approach, and we should probably go with it on
the basis that it meets our need and obviously no-one else has yet found
a need to support specialization constants at all :-).

But you could argue that a goal of the spir-v reader is to represent
everything from spir-v in llvm IR, and thus the specialization should
not be done in readSpirv(). In that case, there would need to be a
representation in IR of specialization constants. That seems difficult
to me, in that you can use a specialization constant anywhere you can
use a normal constant, including the initializer of a global variable."

A follow up discussion has concluded that we should allow the specialization constants to be passed in on the command line, however, it feels that if we could represent the spec constants in llvm using some mechanism, this would be really useful.

valgrind: invalid read in test "LLVM_SPIRV :: selection_merge.spt"

******************** TEST 'LLVM_SPIRV :: selection_merge.spt' FAILED ********************
Script:
--
/home/kherbst/git/llvm/build/bin/llvm-spirv /home/kherbst/git/SPIRV-LLVM-Translator/test/selection_merge.spt -to-binary -o /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.spv
/home/kherbst/git/llvm/build/bin/llvm-spirv /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.spv -to-text -o /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.spt
/home/kherbst/git/llvm/build/bin/FileCheck < /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.spt /home/kherbst/git/SPIRV-LLVM-Translator/test/selection_merge.spt --check-prefix=CHECK-SPIRV
/home/kherbst/git/llvm/build/bin/llvm-spirv -r /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.spv -o /home/kherbst/git/llvm/build/tools/llvm-spirv/test/test_output/Output/selection_merge.spt.tmp.bc
--
Exit Code: 123

Command Output (stderr):
--
==15889== Invalid read of size 4
==15889==    at 0x53336AC: SPIRV::SPIRVLoopMerge::getLoopControl() (in /home/kherbst/git/llvm/build/lib/libLLVMSPIRVLib.so.7svn)
==15889==    by 0x531D133: SPIRV::SPIRVToLLVM::setLLVMLoopMetadata(SPIRV::SPIRVLoopMerge*, llvm::BranchInst*) (SPIRVReader.cpp:913)
==15889==    by 0x531EA22: SPIRV::SPIRVToLLVM::transValueWithoutDecoration(SPIRV::SPIRVValue*, llvm::Function*, llvm::BasicBlock*, bool) (SPIRVReader.cpp:1536)
==15889==    by 0x531CBC9: SPIRV::SPIRVToLLVM::transValue(SPIRV::SPIRVValue*, llvm::Function*, llvm::BasicBlock*, bool) (SPIRVReader.cpp:952)
==15889==    by 0x532423E: SPIRV::SPIRVToLLVM::transFunction(SPIRV::SPIRVFunction*) (SPIRVReader.cpp:2008)
==15889==    by 0x53289C2: SPIRV::SPIRVToLLVM::translate() (SPIRVReader.cpp:2473)
==15889==    by 0x532BA4C: llvm::ReadSPIRV(llvm::LLVMContext&, std::istream&, llvm::Module*&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) (SPIRVReader.cpp:3103)
==15889==    by 0x405281: convertSPIRVToLLVM() (llvm-spirv.cpp:158)
==15889==    by 0x404B66: main (llvm-spirv.cpp:295)
==15889==  Address 0x7e976f8 is 0 bytes after a block of size 200 alloc'd
==15889==    at 0x4C2D0AA: operator new(unsigned long) (vg_replace_malloc.c:333)
==15889==    by 0x51FE333: SPIRV::SPIRVEntry* SPIRV::create<SPIRV::SPIRVSelectionMerge>() (in /home/kherbst/git/llvm/build/lib/libLLVMSPIRVLib.so.7svn)
==15889==    by 0x51F7DA1: SPIRV::SPIRVEntry::create(spv::Op) (SPIRVEntry.cpp:88)
==15889==    by 0x527B5F4: SPIRV::SPIRVDecoder::getEntry() (SPIRVStream.cpp:237)
==15889==    by 0x523FFB2: SPIRV::SPIRVFunction::decodeBB(SPIRV::SPIRVDecoder&) (SPIRVFunction.cpp:151)
==15889==    by 0x523FE3A: SPIRV::SPIRVFunction::decode(std::istream&) (SPIRVFunction.cpp:121)
==15889==    by 0x51F9731: SPIRV::operator>>(std::istream&, SPIRV::SPIRVEntry&) (SPIRVEntry.cpp:432)
==15889==    by 0x527B6DA: SPIRV::SPIRVDecoder::getEntry() (SPIRVStream.cpp:246)
==15889==    by 0x524A8DF: SPIRV::operator>>(std::istream&, SPIRV::SPIRVModule&) (SPIRVModule.cpp:1510)
==15889==    by 0x532BA05: llvm::ReadSPIRV(llvm::LLVMContext&, std::istream&, llvm::Module*&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&) (SPIRVReader.cpp:3099)
==15889==    by 0x405281: convertSPIRVToLLVM() (llvm-spirv.cpp:158)
==15889==    by 0x404B66: main (llvm-spirv.cpp:295)
==15889==

OpIEqual can't be used for pointers

if a pointer is compared against NULL, an OpIEqual instruction is generated with the sources being pointers.

The spec states for OpIEqual: "The type of Operand 1 and Operand 2 must be a scalar or vector of integer type. They must have the same component width, and they must have the same number of components as Result Type."

So I think OpConvertPtrToU have to be inserted before using OpIEqual.

example code:

kernel void
test(ulong global* in)
{
        if (!in)
                return;
        *in += 4;
}
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 19
; Schema: 0
               OpCapability Addresses
               OpCapability Kernel
               OpCapability Int64
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %6 "test"
               OpSource OpenCL_C 102000
      %ulong = OpTypeInt 64 0
    %ulong_4 = OpConstant %ulong 4
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_ulong = OpTypePointer CrossWorkgroup %ulong
          %5 = OpTypeFunction %void %_ptr_CrossWorkgroup_ulong
       %bool = OpTypeBool
         %11 = OpConstantNull %_ptr_CrossWorkgroup_ulong
       %true = OpConstantTrue %bool
          %6 = OpFunction %void None %5
          %7 = OpFunctionParameter %_ptr_CrossWorkgroup_ulong
          %8 = OpLabel
         %13 = OpIEqual %bool %7 %11
         %15 = OpLogicalNotEqual %bool %13 %true
               OpSelectionMerge %10 None
               OpBranchConditional %15 %9 %10
          %9 = OpLabel
         %16 = OpLoad %ulong %7 Aligned 8
         %18 = OpIAdd %ulong %16 %ulong_4
               OpStore %7 %18 Aligned 8
               OpBranch %10
         %10 = OpLabel
               OpReturn
               OpFunctionEnd

spirv-val output:

error: 19: Expected operands to be scalar or vector int: IEqual

Intergration w LLVM

How to integrate this tool with the rest of LLVM to provide seamless flow for the developers.

TODO: A number of upstream threads to be summarized here.

OCL21toSPIRV.cpp relies on using the OpenCL CXX runtime

This seems like a design decision but I'm not sure it's a correct one.

It appears the OCL21 code relies on the libclcxx having being used to lower the C++ APIs down to internal symbols in the cl::__spirv namespace.

SYCL however would like to declare itself as C++ language but doesn't use the libclcxx to generate the symbols, is there any reason this can't just lower the original mangled C++ symbols instead?

drop '-N' from the tag name

That digit usually belongs to downstream packagers, and shouldn't be necessary for you anyway, as you can just bump the minor number when needed.

Duplicate non-aggregate type declarations are not allowed

The current code generates two identical OpTypePointer declarations when using linked lists in OpenCL C. spir-val complains with this: "error: 14: Duplicate non-aggregate type declarations are not allowed. Opcode: TypePointer id: 5"

struct Node {
        global struct Node *node;
};

kernel void test(global struct Node* out, global struct Node* in)
{
        out = in->node;
}
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 12
; Schema: 0
               OpCapability Addresses
               OpCapability Kernel
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %7 "test"
               OpSource OpenCL_C 102000
               OpDecorate %11 FuncParamAttr NoCapture
         %11 = OpDecorationGroup
               OpGroupDecorate %11 %8 %9
               OpTypeForwardPointer %_ptr_CrossWorkgroup__struct_3 CrossWorkgroup
       %void = OpTypeVoid
  %_struct_3 = OpTypeStruct %_ptr_CrossWorkgroup__struct_3
%_ptr_CrossWorkgroup__struct_3 = OpTypePointer CrossWorkgroup %_struct_3
%_ptr_CrossWorkgroup__struct_3_0 = OpTypePointer CrossWorkgroup %_struct_3
          %6 = OpTypeFunction %void %_ptr_CrossWorkgroup__struct_3_0 %_ptr_CrossWorkgroup__struct_3_0
          %7 = OpFunction %void Pure %6
          %8 = OpFunctionParameter %_ptr_CrossWorkgroup__struct_3_0
          %9 = OpFunctionParameter %_ptr_CrossWorkgroup__struct_3_0
         %10 = OpLabel
               OpReturn
               OpFunctionEnd

SPIR ABI injects OpenCL intrinsics names in global namespace

While working on the triSYCL compiler, I have just realized that at some point to generate some SPIR-df code we have to use the SPIR OpenCL intrinsics functions which are just normal C names according to SPIR 2.0 section "1.2 Name mangling" on page 6.
That works for the initial minimum goal of SPIR: supporting OpenCL, but then, what happens if some non-OpenCL code (let us say SYCL or even OpenCL C++ kernel language) is using some user functions with the exact same signature as one of the numerous OpenCL C instrinsic function? :-(
Using SPIR-V from LLVM does not seem to solve this issue since it is basically a translator form SPIR-df LLVM IR to SPIR-V and the clash will happen inside LLVM...

Is it a misinterpretation of the specification by me or we are in trouble...

UConvert from llvm ints < 8 generates incorrect SPIR-V code.

The following IR (generated from libclc):

; Function Attrs: norecurse nounwind readnone
define linkonce_odr dso_local spir_func <2 x i8> @_Z7shuffleDv4_cDv2_h(<4 x i8>, <2 x i8>) local_unnamed_addr #0 {
%3 = and <2 x i8> %1, <i8 3, i8 3>
%4 = extractelement <2 x i8> %3, i64 0
%5 = trunc i8 %4 to i2
switch i2 %5, label %9 [
i2 0, label %10
i2 1, label %6
i2 -2, label %7
i2 -1, label %8
]

; :6: ; preds = %2
br label %10

; :7: ; preds = %2
br label %10

; :8: ; preds = %2
br label %10

; :9: ; preds = %2
unreachable

; :10: ; preds = %8, %7, %6, %2
%11 = phi i64 [ 1, %6 ], [ 2, %7 ], [ 3, %8 ], [ 0, %2 ]
%12 = extractelement <4 x i8> %0, i64 %11
%13 = insertelement <2 x i8> undef, i8 %12, i64 0
%14 = extractelement <2 x i8> %3, i64 1
%15 = trunc i8 %14 to i2
switch i2 %15, label %19 [
i2 0, label %20
i2 1, label %16
i2 -2, label %17
i2 -1, label %18
]

; :16: ; preds = %10
br label %20

; :17: ; preds = %10
br label %20

; :18: ; preds = %10
br label %20

; :19: ; preds = %10
unreachable

; :20: ; preds = %18, %17, %16, %10
%21 = phi i64 [ 1, %16 ], [ 2, %17 ], [ 3, %18 ], [ 0, %10 ]
%22 = extractelement <4 x i8> %0, i64 %21
%23 = insertelement <2 x i8> %13, i8 %22, i64 1
ret <2 x i8> %23
}

creates SPIR-V that fails validation as it creates a 2-bit integer type which is illegal in unextended SPIR-V.

    %u2 = OpTypeInt 2 0

  %41190 = OpFunction %v2uchar Pure %41189
  %41191 = OpFunctionParameter %v4uchar
  %41192 = OpFunctionParameter %v2uchar
  %41193 = OpLabel
  %41219 = OpBitwiseAnd %v2uchar %41192 %41218
  %41220 = OpCompositeExtract %uchar %41219 0
  %41221 = OpUConvert %u2 %41220

These should probably be changed to u8 via masking.

Changes to the installed header file

  1. Should we install it in a prefix or keep it directly under “include”?
    I was initially pushing for installing it in a prefix, but since it’s a single header file, it doesn’t seem that useful.

  2. Should we rename the installed header file? “SPIRV.h” a) seems too generic, and b) could easily be confused for one of the SPIR-V header files, which it is not. We should probably rename it to “LLVMSPIRVLib.h”, to match the library name.

opencl-neo build fails on compiling source with spirv translator invocation

llvm_unreachable("unknown mangling!");

I was building https://github.com/intel/compute-runtime and its dependencies against LLVM-8 and met this build failure from NEO:

unknown mangling!
UNREACHABLE executed at /home/dvrogozh/git/github/SPIRV-LLVM-Translator/lib/SPIRV/OCLUtil.cpp:182!
/bin/sh: line 1: 30434 Aborted                 (core dumped) LD_LIBRARY_PATH=/home/dvrogozh/git/github/compute-runtime/build/bin /home/dvrogozh/git/github/compute-runtime/build/bin/ocloc -q -file media_kernels_backend.cl -device cnl -64 -out_dir /home/dvrogozh/git/github/compute-runtime/build/bin/Gen10core/test_files/x64/

This commit in SPIRV translator master branch fixed the issue: 3fef016. Can it, please, be cherry-picked to https://github.com/KhronosGroup/SPIRV-LLVM-Translator/tree/llvm_release_80 and new release produced?

@tjaalton, @tripzero: FYI

oclIsBuiltin needs rework

The function oclIsBuiltin is suppose to check if the given function name is an OpenCL built-in.

The function is obviously wrong as it will return true for any function that is mangled (starts with _Z).

This is causing issues when converting C++ based applications for OpenCL (like SYCL).

Support for global constructors and destructors

All kernel functions in @llvm.global_ctors and @llvm.global_dtors should have their execution mode marked as, respectively, Initializer or Finalizer. Anything else?

This depends on SPIRV 1.1.

OCL21ToSPIRV.cpp uses language version wrong

I've asked internally at Khronos and was told the source language version in SPIRV for C++ is meant to be 1.0 and isn't the CL version (2.1).

OCL21ToSPIRV::runOnModule checks for CLVer = std::get<1>(Src) which is incorrect.

Should it just check == 10000 or should the check just be removed?

Incorrect Op*AccessChain instruction for vec ptr to scalar ptr casts

I was debugging an issue w/ vstore_local cl cts tests, which does something like:

    vstore2( srcValues[ tid ], offsets[ tid ], ( (__local char *)sSharedStorage ) + alignmentOffset );

I've simplified this to a more trivial example:

__kernel void test_fn(__global int *res, __global int2 *v)
{
    *res = ((__global int *)v)[1];
}

This generates an AccessChain instruction like:

         %10 = OpFunctionParameter %_ptr_CrossWorkgroup_v2uint
         %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %uint_0 %uint_1

I believe this should have been an OpInBoundsAccessChain instead, or equivalently:

         %15 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %10 %uint_0 %uint_0 %uint_1

given that for the Ptr access chain variants:

The type of Base after being dereferenced with Element is still the same as the original type of Base.

Note also spirv-val doesn't seem to be checking the resulting type of the AccessChain instruction, since it doesn't seem to complain about this. Perhaps a separate bug?

(Originally filed at KhronosGroup/SPIRV-Tools#1498 but that seems to be the wrong place)

Should the mangler library be isolated?

In #11 we discussed that the mangler (found in “lib/SPIRV/Mangler”) is not part of the SPIR-V code per se, and could already be found in https://github.com/KhronosGroup/SPIR-Tools.

The mangler should probably be moved out of “lib/SPIRV” to make it clearer that it is not part of the SPIR-V codebase. It could be moved to “lib”, or maybe spun into its separate repository which would then be used in this one as a library or a submodule.

Recommend Projects

  • React photo React

    A declarative, efficient, and flexible JavaScript library for building user interfaces.

  • Vue.js photo Vue.js

    🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.

  • Typescript photo Typescript

    TypeScript is a superset of JavaScript that compiles to clean JavaScript output.

  • TensorFlow photo TensorFlow

    An Open Source Machine Learning Framework for Everyone

  • Django photo Django

    The Web framework for perfectionists with deadlines.

  • D3 photo D3

    Bring data to life with SVG, Canvas and HTML. 📊📈🎉

Recommend Topics

  • javascript

    JavaScript (JS) is a lightweight interpreted programming language with first-class functions.

  • web

    Some thing interesting about web. New door for the world.

  • server

    A server is a program made to process requests and deliver data to clients.

  • Machine learning

    Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.

  • Game

    Some thing interesting about game, make everyone happy.

Recommend Org

  • Facebook photo Facebook

    We are working to build community through open source technology. NB: members must have two-factor auth.

  • Microsoft photo Microsoft

    Open source projects and samples from Microsoft.

  • Google photo Google

    Google ❤️ Open Source for everyone.

  • D3 photo D3

    Data-Driven Documents codes.