khronosgroup / spir

SPIR generator/Clang Installation Instructions

These instructions describe how to build, install and operate SPIR generator Clang.

---

Step 1: Organization

SPIR generator/Clang is designed to be built as part of an LLVM build.

SPIR generator/Clang is based on LLVM/Clang version 3.2.

The LLVM source code could be downloaded from http://www.llvm.org/releases/3.2/llvm-3.2.src.tar.gz.

It is also available directly from the LLVM svn server:

  svn co http://llvm.org/svn/llvm-project/llvm/tags/RELEASE_32/final llvm

Or could be cloned from LLVM git repository:

  git clone http://llvm.org/git/llvm.git llvm
  cd llvm
  git checkout --track -b release_32 remotes/origin/release_32

Assuming that the LLVM source code is located at $LLVM_SRC_ROOT, then the clang source code should be installed as: $LLVM_SRC_ROOT/tools/clang.

The directory is not required to be called clang, but doing so will allow the LLVM build system to automatically recognize it and build it along with LLVM.

  cd $LLVM_SRC_ROOT/tools
  git clone https://github.com/KhronosGroup/SPIR clang
  cd clang
  git checkout --track -b spir_12 remotes/origin/spir_12

---

Step 2: Configure and Build LLVM

Configure and build your copy of LLVM (see $LLVM_SRC_ROOT/GettingStarted.html for more information).

Assuming you installed clang at $LLVM_SRC_ROOT/tools/clang then Clang will automatically be built with LLVM. Otherwise, run make in the Clang source directory to build Clang.

• Note: currently there might be failures in check_clang project.

---

Step 3: (Optional) Verify Your Build

It is a good idea to run the Clang tests to make sure your build works correctly. From inside the Clang build directory, run make test to run the tests.

---

Step 4: Install Clang

If you wish to run Clang from the generated binary directory, you may skip this section.

From inside the Clang build directory, run make install to install the Clang compiler and header files into the prefix directory selected when LLVM was configured.

The Clang compiler is available as clang and clang++. It supports a gcc like command line interface. See the man page for clang (installed into $prefix/share/man/man1) for more information.

---

Step 5: Creating SPIR binaries

To create a SPIR binary from a valid OpenCL-C file (.cl), use the following command lines:

  clang -cc1 -emit-llvm-bc -triple <triple> <OpenCL compile options> -cl-spir-compile-options "<OpenCL compile options>" -include <opencl_spir.h> -o <output> <input>

• <triple>: for 32 bit SPIR use spir-unknown-unknown, for 64 bit SPIR use spir64-unknown-unknown.
• Note: <OpenCL compile options> appears twice. The command line option -cl-spir-compile-options "<OpenCL compile options>" specifies the compile options that occur in the SPIR metadata.
• <opencl_spir.h>: download opencl_spir.h from https://github.com/KhronosGroup/SPIR-Tools/blob/master/headers/opencl_spir.h
• -O: -O0 (default) is the only tested option value at the moment. It's assumed by design that all optimizations are executed by SPIR consumer.

---

Reporting issues

Bugs/feature requests can be filed via github or Khronos Bugzilla bug tracker.

spir's Issues

xyzw and rgba selectors for vectors of 8 and 16 components

OpenCL C++ specification has this example:

int8 v8;
...
int3 v8c1 = v8.xyz; // ill-formed: xyzw and rgba selectors
                    // are not allowed on vector expressions
                    // with more than 4 components

"Table 2.5: Selector values and their corresponding components in swizzle" allows xyzw and rgba selectors only for vectors of 2, 3, 4 components.

Currently clang compiles this code without errors or warnings.
Should it fail in this case?

clang trunk requires cl_khr_int64_base_atomics to be enabled before atomic_long can be used in the header file

Command Line:
clang.exe -cc1 -x cl -cl-std=CL2.0 -include C:/llvmspirv/llvm/tools/clang/lib/Headers/opencl-20.h -emit-llvm -fblocks -triple spir64-unknown-unknown -o _test.bc test.i

Error Messages:
test.i:2:10: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_base_atomics extension to be enabled
volatile atomic_long var;
^
test.i:2:10: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_extended_atomics extension to be enabled
test.i:3:17: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_base_atomics extension to be enabled
global volatile atomic_long g_var;
^
test.i:3:17: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_extended_atomics extension to be enabled
test.i:4:10: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_base_atomics extension to be enabled
volatile atomic_long a_var[2];
^
test.i:4:10: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_extended_atomics extension to be enabled
test.i:5:17: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_base_atomics extension to be enabled
global volatile atomic_long* p_var;
^
test.i:5:17: error: use of type 'atomic_long' (aka '_Atomic(long)') requires cl_khr_int64_extended_atomics extension to be enabled
8 errors generated.

Cause: clang trunk requires cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics to be enabled before atomic_long can be used in the header file

Testcase:
test.zip

`expected top-level entity`, when using `parseIRFile`, on output of `spirv_dis`

The error message:

The c++ code:
https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/test_spir_cpp.cpp#L27-L34

    llvm::LLVMContext context;
    llvm::SMDiagnostic smDiagnostic;
    std::string llFilename = "cl_kernel1.ll";
    std::unique_ptr<llvm::Module> M = parseIRFile(llFilename, smDiagnostic, context);
    if(!M) {
        smDiagnostic.print("irtoopencl", llvm::errs());
        throw std::runtime_error("failed to parse IR");
    }

The input SPIR-V code:

~/git-local/pub-prototyping/spirv/build (master|…5) $ cat cl_kernel1.ll
; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 22
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int8
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical32 OpenCL
               OpEntryPoint Kernel %10 "mykernel"
               OpSource OpenCL_C 102000
               OpName %5 "__spirv_BuiltInGlobalInvocationId"
               OpName %11 "cmem0"
               OpName %12 "offset"
               OpName %13 "entry"
               OpName %14 "add.ptr"
               OpName %18 "call"
               OpName %20 "add"
               OpName %21 "arrayidx"
               OpDecorate %5 BuiltIn GlobalInvocationId
               OpDecorate %5 Constant
               OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
          %2 = OpTypeInt 32 0
          %7 = OpTypeInt 8 0
         %19 = OpConstant %2 123
          %3 = OpTypeVector %2 3
          %4 = OpTypePointer UniformConstant %3
          %6 = OpTypeVoid
          %8 = OpTypePointer CrossWorkgroup %7
          %9 = OpTypeFunction %6 %8 %2
         %15 = OpTypePointer CrossWorkgroup %2
          %5 = OpVariable %4 UniformConstant
         %10 = OpFunction %6 None %9
         %11 = OpFunctionParameter %8
         %12 = OpFunctionParameter %2
         %13 = OpLabel
         %14 = OpInBoundsPtrAccessChain %8 %11 %12
         %16 = OpBitcast %15 %14
         %17 = OpLoad %3 %5
         %18 = OpCompositeExtract %2 %17 0
         %20 = OpIAdd %2 %18 %19
         %21 = OpInBoundsPtrAccessChain %15 %16 %18
               OpStore %21 %20 Aligned 4
               OpReturn
               OpFunctionEnd

Generated by running this script https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/cl-to-spirv.sh

clang -cc1 -emit-spirv -triple spir-unknown-unknown -cl-std=CL1.2 -include opencl.h -x cl -o cl_kernel1.spv ../cl_kernel1.cl
spirv-dis cl_kernel1.spv -o cl_kernel1.ll

against this minimal opencl kernel: https://github.com/hughperkins/pub-prototyping/blob/3310247911a6b26a6768baf165f435faaa25e24e/spirv/cl_kernel1.cl

kernel void mykernel(global char *cmem0, unsigned int offset) {
    global int *data0 = (global int *)(cmem0 + offset);
    int tid = get_global_id(0);
    data0[tid] = tid + 123;
}

Versoins:

• spir-v 1.1
• khronos clang: 3.6.1

Thoughts?

some functions in opencl-20.h failed to compile due to empty parameter list

I was trying to compile some cl kernels with Clang but I got these errors:

C:\p4\stg_win50\git-llvm\llvm\tools\clang\lib\Headers/opencl-20.h:11271:13: error: function with no prototype cannot use the spir_function calling convention
clk_event_t create_user_event();
^
C:\p4\stg_win50\git-llvm\llvm\tools\clang\lib\Headers/opencl-20.h:11279:9: error: function with no prototype cannot use the spir_function calling convention
queue_t get_default_queue();
^
C:\p4\stg_win50\git-llvm\llvm\tools\clang\lib\Headers/opencl-20.h:11585:38: error: function with no prototype cannot use the spir_function calling convention
size_t attribute((overloadable)) get_global_linear_id();
^
C:\p4\stg_win50\git-llvm\llvm\tools\clang\lib\Headers/opencl-20.h:11585:38: error: 'overloadable' function 'get_global_linear_id' must have a prototype
C:\p4\stg_win50\git-llvm\llvm\tools\clang\lib\Headers/opencl-20.h:11586:38: error: function with no prototype cannot use the spir_function calling convention
size_t attribute((overloadable)) get_local_linear_id();

seems clang does not like function declarations with empty parameter list. Add void to that should fix it.

spirv-1.2 branch?

so many versions and repos, and out of date readmes :-) . I'm currently following #44 , but shouldnt we be using branch spirv-1.2 instead of spirv-1.1 now? Per https://www.khronos.org/spir , seems like 1.2 is the current latest 'release' version of SPIR-V?

How to convert LLVM-IR to SPIR-V

The current examples look like this is a one way conversion tool from *.cl files to spirv files.

Is there a way to provide inputs in LLVM-IR? Either via the clang or llc front end?

[[cl::max_size]] does not work for all specified pointer/ref types

OpenCL C++ specification:

This attribute can be provided with a kernel argument of type constant_ptr<T>, constant<T>*, constant<T>&, local_ptr<T>, local<T>*, local<T>&.

#include <opencl_memory>

template <typename T>
struct local_ptr_test
{
    __local T* ptr;
};

typedef local_ptr_test<int> local_int_test;

kernel void foo0([[cl::max_size(1)]] local_ptr_test<int> arg) { }


kernel void foo1([[cl::max_size(1)]] cl::local_ptr<int> arg) { }
kernel void foo2([[cl::max_size(1)]] cl::local<int>* arg) { }
kernel void foo3([[cl::max_size(1)]] cl::local<int>& arg) { }

kernel void foo4([[cl::max_size(1)]] cl::constant_ptr<int> arg) { }
kernel void foo5([[cl::max_size(1)]] cl::constant<int>* arg) { }
kernel void foo6([[cl::max_size(1)]] cl::constant<int>& arg) { }

(foo0 is a kernel from https://github.com/KhronosGroup/SPIR/blob/spirv-1.1/test/OpenCL/OpenCL22/attributes/TestBasicAttributes_max_size.cl)

max_size.cl:14:57: error: max_size attribute only applies to kernel paremeters which are in local or constant address space
kernel void foo1([[cl::max_size(1)]] cl::local_ptr<int> arg) { }
                                                        ^
max_size.cl:16:54: error: max_size attribute only applies to kernel paremeters which are in local or constant address space
kernel void foo3([[cl::max_size(1)]] cl::local<int>& arg) { }
                                                     ^
max_size.cl:18:60: error: max_size attribute only applies to kernel paremeters which are in local or constant address space
kernel void foo4([[cl::max_size(1)]] cl::constant_ptr<int> arg) { }
                                                           ^
max_size.cl:20:57: error: max_size attribute only applies to kernel paremeters which are in local or constant address space
kernel void foo6([[cl::max_size(1)]] cl::constant<int>& arg) { }
                                                        ^
4 errors generated.

I.e. the attribute works with cl::local<int>* and cl::constant<int>*.

(Also there is a typo (paremeters) in error messages)

How to compile an OpenCL2.0 kernel cl file into spir-v 1.0?

I clone the spir-v 1.0 branch and build the clang binary. I want to use it to compile some OpenCL2.0 cl file into the spir-v 1.0 format:

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=CL2.0 -include opencl_spir.h test_pipe_workgroup_readwrite_int.cl

And it shows some syntax error

test_pipe_workgroup_readwrite_int.cl:7:6: error: implicit declaration of function 'is_valid_reserve_id' is invalid in OpenCL
if(is_valid_reserve_id(res_id))
^
1 error generated.

I suppose it should include some header file defining those OpenCL2.0 kernel built-ins. Is there a header file available? Where can I find it?

The only one I can find is for the SPIR (https://github.com/KhronosGroup/SPIR-Tools/blob/master/headers/opencl_spir.h), I guess it only defines OpenCL 1.2 built-ins.

Thanks.

Is it a bug? "'half' cannot be used as the type of a kernel parameter"?

Given a 1.2 kernel as t.cl

pragma OPENCL EXTENSION cl_khr_fp16 : enable

kernel void half_scalar_d(half halfd,
const half consthalfd,
private half privatehalfd,
private const half privateconsthalfd)
{}

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=CL1.2 -include opencl.h t.cl

t.cl:3:32: error: 'half' cannot be used as the type of a kernel parameter
kernel void half_scalar_d(half halfd,
^
t.cl:4:38: error: 'const half' cannot be used as the type of a kernel parameter
const half consthalfd,
^
t.cl:5:40: error: 'half' cannot be used as the type of a kernel parameter
private half privatehalfd,
^
t.cl:6:46: error: 'const half' cannot be used as the type of a kernel parameter
private const half privateconsthalfd)
^
4 errors generated.

[[cl::unroll_hint]] and [[cl::ivdep]] hints are not passed to SPIR-V

If I recall correctly, it is not specified in any document if OpenCL C++-to-SPIR-V compiler should always pass [[cl::unroll_hint]] and [[cl::ivdep]] (ignore vector dependencies) hints to SPIR-V, or if OpenCL C++-to-SPIR-V compiler can decide not unroll the loop and ignore unroll hint.

However, in my opinion, since SPIR-V is an intermediate language between human-readable OpenCL (and other languages) and hardware-specific byte code, OpenCL C++-to-SPIR-V compiler should compile loops with [[cl::unroll_hint]] and [[cl::ivdep]] attributes. That is, it should compile those loops to structured loops (see StructuredControlFlow) with OpLoopMerge instruction with information about the hints, so that later SPIR-V-to-hardware-specific-byte-code compiler can make a decision whether to unroll or vectorize the loop.

Currently, [[cl::unroll_hint]] and [[cl::ivdep]] hints are ignored and are not passed SPIR-V.

clang trunk requires function prototype for atom_add and atom_or

Command Line:
clang -cc1 -x cl -cl-std=CL2.0 -include C:/llvmspirv/llvm/tools/clang/lib/Headers/opencl-20.h -emit-llvm -fblocks -triple spir64-unknown-unknown -o test.bc test.i
Error Message:
test.i:6:5: warning: implicit declaration of function 'atom_add' is invalid in C99
atom_add(&dst[t_address], 1u);
^
test.i:6:5: error: function with no prototype cannot use the spir_function calling convention
test.i:7:5: warning: implicit declaration of function 'atom_or' is invalid in C99
atom_or(&dst[t_address], 1u);
^
test.i:7:5: error: function with no prototype cannot use the spir_function calling convention
2 warnings and 2 errors generated.
Cause:
Requires function prototype for atom_add and atom_or.
Testcase:
test.zip

Error fetching SPIRV-1.1 branch?

Hi,
just following the instructions in step:
git checkout --track -b khronos/spirv-3.6.1 remotes/origin/khronos/spirv-3.6.1
I get
fatal: A branch named 'khronos/spirv-3.6.1' already exists.
is that OK?

There is a bug of "get_image_width"

Given a 1.2 kernel as image.cl

__kernel void testWriteui(__global uchar *src, write_only image2d_t dstimg)
{
int tid_x = get_global_id(0);
int tid_y = get_global_id(1);
int indx = tid_y * get_image_width(dstimg) + tid_x;
uint4 color;

indx *= 4;
     color.x = (uint)src[indx+0];
     color.y = (uint)src[indx+1];
     color.z = (uint)src[indx+2];
     color.w = (uint)src[indx+3];
write_imageui(dstimg, (int2)(tid_x, tid_y), color);

}

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=CL1.2 -include opencl.h image.cl

image.cl:5:27: error: no matching function for call to 'get_image_width'
int indx = tid_y * get_image_width(dstimg) + tid_x;
^~~~~~~~~~~~~~~
opencl.h:15380:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image1d_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image1d_t image);
^
opencl.h:15381:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image1d_buffer_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image1d_buffer_t image);
^
opencl.h:15382:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image2d_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image2d_t image);
^
opencl.h:15383:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image3d_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image3d_t image);
^
opencl.h:15384:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image1d_array_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image1d_array_t image);
^
opencl.h:15385:48: note: candidate function not viable: no known conversion from 'write_only image2d_t' to '__read_only image2d_array_t' for 1st argument
int __const_func __attribute((overloadable)) get_image_width(image2d_array_t image);
^
1 error generated.

"get_image_width" with a write_only image is valid, isn't it?

default branch should be `master` or `spirv-1.1`, not `spir_12`?

I think it is confusing that the default branch is spir_12, rather than master or spirv-1.2

clang trunk requires cl_khr_gl_msaa_sharing to be enabled before image2d_array_msaa_t can be used in the header file

Command Line:
clang -cc1 -x cl -cl-std=CL2.0 -include C:\llvmspirv\llvm\tools\clang\lib\Headers/opencl-20.h -emit-llvm -triple spir64-unknown-unknown -o _temp_1_Tonga.clang.bc _temp_1_Tonga.i

Error Messages:
In file included from :303:
In file included from :1:
C:\llvmspirv\llvm\tools\clang\lib\Headers/opencl-20.h:3945:60: error: use of type 'image2d_msaa_t' requires cl_khr_gl_msaa_sharing extension to be enabled
float4 attribute((overloadable)) read_imagef(read_only image2d_msaa_t image, int2 coord, int sample);

Cause: clang trunk requires cl_khr_gl_msaa_sharing to be enabled before image2d_array_masaa_t can be used, in the header file.

Workaround: Enable this extension at the beginning of header file, and disable at the end.

errors like "no known conversion from '__global half *' to 'const double *'"

When I run "clang -cc1 -emit-llvm-bc -triple spir64-unknown-unknown -cl-std=CL1.2 -cl-mad-enable -cl-spir-compile-options "-cl-std=CL1.2 -cl-mad-enable" -include opencl_spir.h -o result kernels.cl", errors like below appear, can anyone help, thanks.

...
opencl_spir.h:5001:37: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant char
char3 __attribute__((overloadable)) vload3(size_t offset, const __constant char *p);
                                    ^
opencl_spir.h:5002:38: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant uchar
uchar3 __attribute__((overloadable)) vload3(size_t offset, const __constant uchar *p);
                                     ^
opencl_spir.h:5003:38: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant short
short3 __attribute__((overloadable)) vload3(size_t offset, const __constant short *p);
                                     ^
opencl_spir.h:5004:39: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant ushor
ushort3 __attribute__((overloadable)) vload3(size_t offset, const __constant ushort *p);
                                      ^
opencl_spir.h:5005:36: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant int *
int3 __attribute__((overloadable)) vload3(size_t offset, const __constant int *p);
                                   ^
opencl_spir.h:5006:37: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant uint
uint3 __attribute__((overloadable)) vload3(size_t offset, const __constant uint *p);
                                    ^
opencl_spir.h:5007:37: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant long
long3 __attribute__((overloadable)) vload3(size_t offset, const __constant long *p);
                                    ^
opencl_spir.h:5008:38: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant ulong
ulong3 __attribute__((overloadable)) vload3(size_t offset, const __constant ulong *p);
                                     ^
opencl_spir.h:5009:38: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant float
float3 __attribute__((overloadable)) vload3(size_t offset, const __constant float *p);
                                     ^
opencl_spir.h:5038:39: note: candidate function not viable: no known conversion from '__global half *' to 'const __constant doubl
double3 __attribute__((overloadable)) vload3(size_t offset, const __constant double *p);

5 warnings and 2100 errors generated.

#spirv-1.0 Undeclared identifier 'createSPIRVWriterPass'

I encountered the following error while trying to compile branch spirv-1.0:

[1424/1526] Building CXX object tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/BackendUtil.cpp.o
FAILED: /usr/bin/clang++   -DGTEST_HAS_RTTI=0 -D_DEBUG -D_GNU_SOURCE -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS -Itools/clang/lib/CodeGen -I../tools/clang/lib/CodeGen -I../tools/clang/include -Itools/clang/include -Iinclude -I../include -D_GLIBCXX_USE_CXX11_ABI=0  -fPIC -fvisibility-inlines-hidden -Wall -W -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -pedantic -Wno-long-long -Wcovered-switch-default -Wnon-virtual-dtor -std=c++11 -fcolor-diagnostics -ffunction-sections -fdata-sections -fno-common -Woverloaded-virtual -fno-strict-aliasing -Wno-nested-anon-types -g -fPIC   -fno-exceptions -fno-rtti -MMD -MT tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/BackendUtil.cpp.o -MF tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/BackendUtil.cpp.o.d -o tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/BackendUtil.cpp.o -c ../tools/clang/lib/CodeGen/BackendUtil.cpp
In file included from ../tools/clang/lib/CodeGen/BackendUtil.cpp:21:
In file included from ../include/llvm/IR/DataLayout.h:25:
In file included from ../include/llvm/IR/DerivedTypes.h:21:
In file included from ../include/llvm/IR/Type.h:19:
../include/llvm/ADT/APFloat.h:352:12: warning: redundant move in return statement [-Wredundant-move]
    return std::move(Value);
           ^
../include/llvm/ADT/APFloat.h:352:12: note: remove std::move call here
    return std::move(Value);
           ^~~~~~~~~~     ~
../tools/clang/lib/CodeGen/BackendUtil.cpp:601:31: error: use of undeclared identifier 'createSPIRVWriterPass'
    getPerModulePasses()->add(createSPIRVWriterPass(*OS));

How to force to generate "TypeFloat 32" instead of "TypeFloat 64" when calling "printf("%f")

Given a simple 'printf' test kernel as:

__kernel void test(const float x)
{
printf("%f\n", x);
}

My question is:

Suppose the platform doesn't support 'double'. Is there any way for the tool to generate "TypeFloat 32" instead of "TypeFloat 64" and not to introduce "FConvert" when calling "printf"? Is there any compilation options to control it?

The tool generates the following SPV code:

119734787 65536 393230 25 0
2 Capability Addresses
2 Capability Kernel
2 Capability Float64
5 ExtInstImport 1 "OpenCL.std"
3 MemoryModel 1 2
5 EntryPoint 6 16 "test"
3 Source 3 102000
4 Name 12 ".str"
3 Name 17 "x"
4 Name 18 "entry"
4 Name 20 "conv"
4 Name 24 "call"
3 Decorate 12 Constant
4 Decorate 12 Alignment 1
4 TypeInt 2 8 0
4 TypeInt 7 32 0
4 Constant 2 3 37
4 Constant 2 4 102
4 Constant 2 5 10
4 Constant 2 6 0
4 Constant 7 8 4
4 Constant 7 21 0
4 TypeArray 9 2 8
4 TypePointer 11 0 9
2 TypeVoid 13
3 TypeFloat 14 32
4 TypeFunction 15 13 14
3 TypeFloat 19 64
4 TypePointer 22 0 2
7 ConstantComposite 9 10 3 4 5 6
5 Variable 11 12 0 10

5 Function 13 16 0 15
3 FunctionParameter 14 17

2 Label 18
4 FConvert 19 20 17
6 InBoundsPtrAccessChain 22 23 12 21 21
7 ExtInst 7 24 1 printf 23 20
1 Return

1 FunctionEnd

clang trunk pre-define ndrange_t in the compiler

Command Line:
clang -cc1 -x cl -cl-std=CL2.0 -include C:\llvmspirv\llvm\tools\clang\lib\Headers/opencl-20.h -emit-llvm -fblocks -triple spir64-unknown-unknown -o _temp_1_Tonga.clang.bc _temp_1_Tonga.i

Error Message:
In file included from :305:
In file included from :1:
C:\llvmspirv\llvm\tools\clang\lib\Headers/opencl-20.h:11230:3: error: typedef redefinition with different types ('struct ndrange_t' vs 'ndrange_t')
} ndrange_t;

Likely cause: clang trunk pre-defines ndrange_t in the compiler

Possible Workaround: Remove the ndrange_t type definition in the header file

OpenCL C++ [[cl::packed]] attribute does not work for enum class

It works for enum but ignored for enum class.

Here is an example:

enum [[cl::packed]] enum1
{
    e11 = 0,
    e12 = 100
};

enum [[cl::packed]] enum2
{
    e21 = 0,
    e22 = 1000
};

enum class [[cl::packed]] enum3
{
    e31 = 0,
    e32 = 100
};

enum class [[cl::packed]] enum4
{
    e41 = 0,
    e42 = 1000
};

void f(enum1, enum2, enum3, enum4);


kernel void k()
{
    f(enum1::e11, enum2::e22, enum3::e32, enum4::e42);
}

static_assert(sizeof(enum1) == 1, "enum1");
static_assert(sizeof(enum2) == 2, "enum2");
static_assert(sizeof(enum3) == 1, "enum3");
static_assert(sizeof(enum4) == 2, "enum4");

SPIR-V

4 TypeInt 3 8 0 
4 TypeInt 4 16 0 
4 TypeInt 5 32 0 
4 Constant 3 15 0 
4 Constant 4 16 1000 
4 Constant 5 17 100 
4 Constant 5 18 1000 
2 TypeVoid 2 
7 TypeFunction 6 2 3 4 5 5 <-- should be 3 4 3 4

Output:

packed-enum.cl:35:1: error: static_assert failed "enum3"
static_assert(sizeof(enum3) == 1, "enum3");
^             ~~~~~~~~~~~~~~~~~~
packed-enum.cl:36:1: error: static_assert failed "enum4"
static_assert(sizeof(enum4) == 2, "enum4");
^             ~~~~~~~~~~~~~~~~~~
2 errors generated.

Cannot build spirv compiler. What i am doing wrong?

Hello

I am trying to build SPIRV compiler.
I think i do all according to instructions:

1. git clone -b khronos/spirv-3.6.1 https://github.com/KhronosGroup/SPIRV-LLVM.git llvm
2. cd llvm/tools
3. git clone -b spirv-1.1 https://github.com/KhronosGroup/SPIR clang
4. cd ../..
5. mkdir build
6. cd build
7. ../llvm/configure --prefix=$PWD/done --enable-optimized --enable-targets=x86_64
8. make

I get error during compiling:

make[2]: *** No rule to make target '/home/nick/opencl/build/Release+Asserts/lib/libLLVMSPIRVLib.a', needed by '/home/nick/opencl/build/Release+Asserts/bin/llvm-c-test'. Stop.

What i am doing wrong?

Besides that i was able to build older version of SPIR compiler from branch spir_12
But as i understand modern compiler is not a SPIR but SPIRV.. I think i should not study solution from older things..

OpenCL C++ attributes do not work for array variables/fields

(tested on spirv-1.1 branch)

Here is a test file based on examples from OpenCL C++ specification:

struct foo1
{
    char a;
    int x[2] [[cl::packed]];
    int y [[cl::packed]];
};

struct foo2
{
    int x[2] [[cl::aligned(8)]];
    int y [[cl::aligned(8)]];
};

int x [[cl::aligned(16)]] = 0;
short array[3] [[cl::aligned]];

Output for clang -cc1 array-attribute-cpp.cl -triple spir-unknown-unknown -cl-std=c++:

array-attribute-cpp.cl:4:16: error: 'packed' attribute cannot be applied to types
    int x[2] [[cl::packed]];
               ^
array-attribute-cpp.cl:10:16: error: 'aligned' attribute cannot be applied to types
    int x[2] [[cl::aligned(8)]];
               ^
array-attribute-cpp.cl:15:18: error: 'aligned' attribute cannot be applied to types
short array[3] [[cl::aligned]];
                 ^
3 errors generated.

The same OpenCL C code (but using __attribute__) is compiled without errors:

struct foo1
{
    char a;
    int x[2] __attribute__ ((packed));
    int y __attribute__ ((packed));
};

struct foo2
{
    int x[2] __attribute__ ((aligned (8)));
    int y __attribute__ ((aligned (8)));
};

int x __attribute__((aligned (16))) = 0;
short array[3] __attribute__ ((aligned));

(Interesting that int z [[cl::aligned(8)]] [2]; is compiled, but int z __attribute__ ((aligned (8))) [2]; is not)

Improper read_image built-ins attribute.

A bunch of read_imageXX bult-ins declared in opencl_1.2.h and opencl.h has attribute((const)) (translated to readnone in LLVM IR) which instructs the compiler what these functions do not access memory which is wrong by definition. Proper attribute is attribute((pure)) that means the functions do not modify memory. It seem both (SPIR 1.2 and SPIR-V) generators have this bug.

Clang/LLVM master integration

As per the discussion here: #43

The conversation has been moved to this thread and original PR closed. Ideally we'd like to integrate SPIR and SPIRV-Clang with Clang and LLVM master. Work has already been done to bring SPIRV-LLVM up to LLVM 3.8, but the corresponding work for SPIR is nontrivial. Any ideas of how we should do this?

Updating to newer clang version / merging with SPIRV-LLVM

Hello,

What is the current status for this repository? Are there any plan to port the code to a newer version of clang, which would make it possible to use clang with SPIR support alongside LLVM with SPIR-V support, without having to use two different versions of LLVM? Or maybe, even try to upstream SPIR support?

error in ValueMap.h converting to a bool

Following the SPIRV generator build instructions from the readme, my make invocation fails at 10% with an error in ValueMap.h regarding a failure to convert to a bool.

In file included from /home/pkeir/code/ext/khronos-opencl-spirv/llvm/lib/CodeGen/CodeGenPrepare.cpp:36:
/home/pkeir/code/ext/khronos-opencl-spirv/llvm/include/llvm/IR/ValueMap.h: In member function ‘bool llvm::ValueMap<KeyT, ValueT, Config>::hasMD() const’:
/home/pkeir/code/ext/khronos-opencl-spirv/llvm/include/llvm/IR/ValueMap.h:104:31: error: cannot convert ‘const std::unique_ptr<llvm::DenseMap<const llvm::Metadata*, llvm::TrackingMDRef> >’ to ‘bool’ in return
   bool hasMD() const { return MDMap; }

I'm using the build commands below. Any suggestions?

git clone -b khronos/spirv-3.6.1 https://github.com/KhronosGroup/SPIRV-LLVM.git llvm
cd llvm/tools
git clone -b spirv-1.1 https://github.com/KhronosGroup/SPIR clang
cd ../..
mkdir build    && cd build 
cmake ../llvm -DLLVM_USE_LINKER=gold -DCMAKE_INSTALL_PREFIX=$PWD/../install -DCMAKE_BUILD_TYPE=Release
make -j 4 && make install

Thanks.

Calls to Barrier are being incorrectly duplicated

The OpenCL builtin barrier must be guaranteed to be executed by all instances of a workgroup with the same arguments. Currently, the SPIR frontend is applying Control-Flow optimisations which break this guarantee on input programs. Consider:

__kernel void entry_point(
__write_only image2d_t dst_image,
__read_only image2d_t src_image,
__global const float2 *tileULs_g,
__global const float4 *tileDxDy_g
) {
__local float integral[10][10];

sampler_t samplerA = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;

int lid = get_local_id(0);
int bid = get_group_id(0);
float2 ids;
ids.x = lid;
ids.y = lid;

float2 tileUL = tileULs_g[bid];
float4 tileDxDy = tileDxDy_g[bid];
int2 tileOut = (int2){(bid % ARRAYTILES) * TILESIZE, (bid / ARRAYTILES) * TILESIZE};

if (lid<10) {
    tileUL += ids*tileDxDy.xy;
    float sum = 0;
    for(int i=0; i<TILESIZE; ++i){
        sum += read_imagef( src_image, samplerA, tileUL).x;
        integral[i][lid] = sum;
        tileUL += tileDxDy.zw;
    }
}

// The Barrier here is unconditional
barrier(CLK_LOCAL_MEM_FENCE);

if (lid<10){
    float sum=0;
    for(int i=0; i<TILESIZE; ++i){
        sum+=integral[lid][i];
        write_imagef(dst_image, tileOut + (int2){i, lid}, sum);
    }
}

}

The SPIR output from this file is changed so that the semantics of the input program are like this:

(code trimmed for brevity)
if (lid<10) {
tileUL += ids*tileDxDy.xy;
float sum = 0;
for (int i=0; i<TILESIZE; ++i) {
sum += read_imagef( src_image, samplerA, tileUL).x;
integral[i][lid] = sum;
tileUL += tileDxDy.zw;

    // Duplicated call
    barrier(CLK_LOCAL_MEM_FENCE);

float sum=0;
for(int i=0; i<TILESIZE; ++i){//add up rows
    sum+=integral[lid][i];
    write_imagef(dst_image, tileOut + (int2){i, lid}, sum);
}

}
else
{
// Duplicated call
barrier(CLK_LOCAL_MEM_FENCE);
}

In more modern versions of clang/llvm, the barrier function can be applied the attribute((noduplicate)) which prevents this sort of optimisations from being performed.

Adding -O0 to the build options appears to fix the issue, so that's a potential work-around.

Capability Addresses is not allowed by Vulkan 1.0 specification

Using LunarG's VulkanSDK=1.1.73.0, LLVM=spirv-3.6.1 and clang=spirv-1.0, I generated a SPIR-V binary from an OpenCL kernel and the disassembly looks like this:

; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 18
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical32 OpenCL
               OpEntryPoint Kernel %9 "CopyBuffer"
               OpSource OpenCL_C 102000
               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
               OpName %src "src"
               OpName %dst "dst"
               OpName %entry "entry"
               OpName %call "call"
               OpName %arrayidx "arrayidx"
               OpName %arrayidx1 "arrayidx1"
               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
               OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
               OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
       %uint = OpTypeInt 32 0
     %v3uint = OpTypeVector %uint 3
%_ptr_UniformConstant_v3uint = OpTypePointer UniformConstant %v3uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
          %8 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_UniformConstant_v3uint UniformConstant
          %9 = OpFunction %void None %8
        %src = OpFunctionParameter %_ptr_CrossWorkgroup_uint
        %dst = OpFunctionParameter %_ptr_CrossWorkgroup_uint
      %entry = OpLabel
         %13 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId
       %call = OpCompositeExtract %uint %13 0
   %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %src %call
         %16 = OpLoad %uint %arrayidx Aligned 4
  %arrayidx1 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %dst %call
               OpStore %arrayidx1 %16 Aligned 4
               OpReturn
               OpFunctionEnd

But trying to create a shader module using vkShaderCreateModule(), this shows up:

Validation(ERROR): msg_code: 5: Object: VK_NULL_HANDLE (Type = 0) | SPIR-V module not valid: Capability Addresses is not allowed by Vulkan 1.0 specification (or requires extension)

This seems strange as I can see Addresses listed under the Capabilities column in SPIR-V 1.0 (https://www.khronos.org/registry/spir-v/specs/1.0/SPIRV.html#Capability).

The vkEnumerateInstanceVersion() returns 1.1.73 so I am definitely above Vulkan 1.0 version as well.

Also, the function returns a VK_ERROR_INITIALIZATION_FAILED code which doesn't seem to be inline with possible return return codes listed here : https://www.khronos.org/registry/vulkan/specs/1.1-extensions/man/html/vkCreateShaderModule.html.
This seems more of an issue with AMD's Vulkan implementation though.

Should it be possible to use long long?

File test/libclcxx/pipes/TestBasicPipeStructData.cl (OpenCL 2.2 C++) from libclcxx tests contains a structure which contains a long long member. If you compile that file as describled in the readme there are no errors. If later you translate binary .spir file to text SPIR-V .spt file (using llvm-spirv or spirv-dis), it turns out that there is a TypeInt <id> 128 0 SPIR-V instruction, which defines a 128-bit integer.

That means that not only OpenCL C++ compiler accepted long long type, which is not defined in OpenCL C++ spec (however, long long is a reserved for future use), but also compiled it to 128-bit integer (I would somehow understand if it compiles long long to 64-bit integer`).

// in .cl file:
struct structType
{
    float a;
    int b;
    long long c;
};

// in .spt file:
4 TypeInt <ID> 128 0

Is it correct behaviour? Because I would expect a compilation error.

llvm-spirv does not output optional literals in OpExecutionMode

When Execution Mode is SubgroupsPerWorkgroup (36) it outputs
4 ExecutionMode 4 36
It should output: 4 ExecutionMode 4 36 <integer literal>.
I haven't try all other Execution Modes, but it works correctly for LocalSize, LocalSizeHint, VecTypeHint.

(spirv-dis from SPIRV-Tools shows “OpExecutionMode %4 SubgroupsPerWorkgroup 12” as expected)

Issue about handling the 'select' built-in

OpenCL 1.2 SPec Table 6.14:

gentype select (gentype a, gentype b, igentype c)

For each component of a vector type,
result[i] = if MSB of c[i] is set ? b[i] : a[i].

I write the following simple kernel:

typedef int4 T;

__kernel void test(__global T* result, __global T* a, __global T* b, __global T* c)
{
result[0] = select( a[0], b[0], c[0] );
}

And the tool generates the following SPV file:

5 Function 2 7 0 6
3 FunctionParameter 5 8
3 FunctionParameter 5 9
3 FunctionParameter 5 10
3 FunctionParameter 5 11

2 Label 12
5 InBoundsPtrAccessChain 5 14 9 13
6 Load 4 15 14 2 16
5 InBoundsPtrAccessChain 5 16 10 13
6 Load 4 17 16 2 16
5 InBoundsPtrAccessChain 5 18 11 13
6 Load 4 19 18 2 16
8 ExtInst 4 20 1 select 15 17 19
5 InBoundsPtrAccessChain 5 21 8 13
5 Store 21 20 2 16
1 Return

1 FunctionEnd

Basically, it just directly translates the OpenCL built-in 'select' into the SPIR-V OpCode "ExtInst Select".

However, in OpenCL.ExtendedInstructionSe.100.pdf from Khronos website: page 51

2.5 Relational instructions
This section describes the list of relational instructions that take scalar or vector arguments. The vector versions of the integer functions operate component-wise. The description is per-component.

select

Each bit of the result is the corresponding bit of a if the corresponding bit of c is 0. Otherwise it is the corresponding
bit of b.

Apparently, the OpenCL built-in "select" has a different meaning from the SPIR-V's "ExtInst select".

OpenCL select:
result[i] = if MSB of c[i] is set ? b[i] : a[i].

SPIR-V "Ext select": bit by bit
result[i] = c[i] ? b[i] : a[i].

It is not correct by simply converting 'select' built-in into 'Ext select'. Is it right?