khronosgroup / spirv-llvm Goto Github PK

ATM the reader can only consume little-endian byte order on a little-endian machine. It needs to detect byte-order of a SPIR-V BC and a host machine and do the byte order conversion when necessary.

Even with 0dfb100, LegalIntWidths remains empty since none of the Desc fed to DataLayout contains an 'n' attribute (which is the only case where LegalIntWidths will be appended new elements).

Releases would be very useful for https://github.com/ldc-developers/ldc to do testing as we are in the process of adding a SPIR-V backend.

Seems a new regression has appeared with the translation of atomic_inc.

I have included the testcase to reproduce the issue. Also I have attached the translated llvm before regression and after regression. The translated llvm is now incorrect.

test.zip

The main differences:
@_Z10atomic_incPVU3AS1i(i32 addrspace(1)* %count)
has become
@_Z10atomic_incPVU3AS1iii(i32 addrspace(1)* %count, i32 1, i32 16)

So declaration changed from:
declare spir_func i32 @Z10atomic_incPVU3AS1i(i32 addrspace(1)) #0
into
declare spir_func i32 @Z10atomic_incPVU3AS1iii(i32 addrspace(1), i32, i32) #0

llvm-spirv tool moves OpUndef instruction from function to program scope. It is bug in the llvm-spirv tool.
Step to reproduce:

1. Convert attached spv file to binary (llvm-spirv.exe -to-binary reproducer.spt.txt)
2. Convert spv file back to text.
3. Undef will not be in body of function.
   After fixing this bug, please implement LIT test for OpUndef in function.

reproducer.spt.txt

given the following OpenCL C code:

typedef struct {
    int val;
} test_struct;

kernel void struct_test(global int* buf, test_struct param) {
    buf[get_global_id(0)] = param.val;
}

kernel void int_test(global int* buf, int param) {
    buf[get_global_id(0)] = param;
}

resulting in the following IR (shortened for brevity):

%struct.test_struct = type { i32 }

define spir_kernel void @struct_test(i32 addrspace(1)* %buf, %struct.test_struct* %param) nounwind {
  %1 = getelementptr inbounds %struct.test_struct* %param, i64 0, i32 0
  %2 = load i32* %1, align 4, !tbaa !12
  %3 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %4 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %3
  store i32 %2, i32 addrspace(1)* %4, align 4, !tbaa !12
  ret void
}

define spir_kernel void @int_test(i32 addrspace(1)* %buf, i32 %param) nounwind {
  %1 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %2 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %1
  store i32 %param, i32 addrspace(1)* %2, align 4, !tbaa !12
  ret void
}

resulting in the following SPIR-V (shortened for brevity):

               OpEntryPoint Kernel %12 "struct_test"
               OpEntryPoint Kernel %25 "int_test"
               OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
          %2 = OpTypeInt 64 0
          %7 = OpTypeInt 32 0
         %16 = OpConstant %2 0
         %17 = OpConstant %7 0
          %3 = OpTypeVector %2 3
          %4 = OpTypePointer UniformConstant %3
          %6 = OpTypeVoid
          %8 = OpTypePointer CrossWorkgroup %7
          %9 = OpTypeStruct %7
         %10 = OpTypePointer Function %9
         %11 = OpTypeFunction %6 %8 %10
         %18 = OpTypePointer Function %7
         %24 = OpTypeFunction %6 %8 %7
          %5 = OpVariable %4 UniformConstant
         %12 = OpFunction %6 None %11
         %13 = OpFunctionParameter %8
         %14 = OpFunctionParameter %10
         %15 = OpLabel
         %19 = OpInBoundsPtrAccessChain %18 %14 %16 %17
         %20 = OpLoad %7 %19 Aligned 4
         %21 = OpLoad %3 %5 Aligned 0
         %22 = OpCompositeExtract %2 %21 0
         %23 = OpInBoundsPtrAccessChain %8 %13 %22
               OpStore %23 %20 Aligned 4
               OpReturn
               OpFunctionEnd
         %25 = OpFunction %6 None %24
         %26 = OpFunctionParameter %8
         %27 = OpFunctionParameter %7
         %28 = OpLabel
         %29 = OpLoad %3 %5 Aligned 0
         %30 = OpCompositeExtract %2 %29 0
         %31 = OpInBoundsPtrAccessChain %8 %26 %30
               OpStore %31 %27 Aligned 4
               OpReturn
               OpFunctionEnd

Is the way kernel struct parameters are handled really the correct/intended behavior?
Considering that scalar types are directly used in OpFunctionParameter/OpTypeFunction, shouldn't structs be handled the same way instead of going through an "OpTypePointer Function" indirection? Even more, doesn't this indirection say that only a pointer argument will be set/used (4 or 8 bytes), not so much a struct object (which could be any size)?
I know that the issue here is that LLVM/SPIR can only handle struct parameters as pointers, but something like that isn't specified for SPIR-V.

How to solve this?

Option 1 (preferable):
Keep it the way it is right now, but explicitly specify that kernel pointer parameters to Function/private memory actually perform some kind of allocation of the element/pointee type on the device side, and are set as this element/pointee type on the host side (not as the pointer type). Note that private address space pointer kernel arguments are otherwise invalid.

Option 2 (impossible?):
Directly use OpTypeStruct in OpFunctionParameter/OpTypeFunction. This will however require IR/SPIR-V translator changes, since OpTypeStruct is no longer a pointer type (making all GEPs/Op*AccessChain instructions using it invalid). This might be impossible to do though, since there is no way of getting a pointer to this struct then in SPIR-V (afaik).

edit:
Option 3:
Require a OpVariable in OpFunctionParameter/OpTypeFunction for struct types. This way it should be clear what is actually happening + it is still a pointer.

(will be cross-posting to https://github.com/KhronosGroup/SPIRV-Headers/issues since I think this is a spec bug that at the very least requires some explicit text that mentions the correct behavior)

Freshly compiled with clang-5.0, running checks via make check-all
Ubuntu 16.04

******************** TEST 'Clang :: CodeGenOpenCL/builtins-r600.cl' FAILED ********************
Script:
--
/home/robin/Development/SPIRVExperimental/llvm/build/./bin/clang -cc1 -internal-isystem /home/robin/Development/SPIRVExperimental/llvm/build/bin/../lib/clang/3.6.1/include -nostdsysteminc -triple r600-unknown-unknown -S -emit-llvm -o - /home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl | /home/robin/Development/SPIRVExperimental/llvm/build/./bin/FileCheck /home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl
/home/robin/Development/SPIRVExperimental/llvm/build/./bin/clang -cc1 -internal-isystem /home/robin/Development/SPIRVExperimental/llvm/build/bin/../lib/clang/3.6.1/include -nostdsysteminc -triple amdgcn-unknown-unknown -S -emit-llvm -o - /home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl | /home/robin/Development/SPIRVExperimental/llvm/build/./bin/FileCheck /home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl
--
Exit Code: 1

Command Output (stderr):
--
/home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl:8:11: error: expected string not found in input
// CHECK: call { double, i1 } @llvm.AMDGPU.div.scale.f64(double %a, double %b, i1 true)
          ^
<stdin>:6:32: note: scanning from here
define void @test_div_scale_f64(double addrspace(1)* %out, i32 addrspace(1)* %flagout, double %a, double %b) #0 {
                               ^
<stdin>:19:7: note: possible intended match here
 %2 = call { double, i1 } @llvm.AMDGPU.div.scale.f64(double %0, double %1, i1 true)
      ^
/home/robin/Development/SPIRVExperimental/llvm/tools/clang/test/CodeGenOpenCL/builtins-r600.cl:21:11: error: expected string not found in input
// CHECK: call { float, i1 } @llvm.AMDGPU.div.scale.f32(float %a, float %b, i1 true)
          ^
<stdin>:38:32: note: scanning from here
define void @test_div_scale_f32(float addrspace(1)* %out, i32 addrspace(1)* %flagout, float %a, float %b) #0 {
                               ^
<stdin>:51:7: note: possible intended match here
 %2 = call { float, i1 } @llvm.AMDGPU.div.scale.f32(float %0, float %1, i1 true)
      ^

--

********************

According to SPIR-V specification, Scope and memory semantics "Must be an of a 32-bit integer scalar that contains a mask."
Translator encodes them as literals instead of IDs.

Is GLSL supported? If do GLSL -> .spv using glslangValidator then .spv -> LLVM IR using llvm-spirv -r, I get InvalidBuiltinSetName: Expects OpenCL.std. Actual is GLSL.std.450.

During validation process of #121 pull request I found a few other issues:

• some arguments of OpControlBarrier are dropped during conversion from SPIR-V to OpenCL 1.2/2.0
• some arguments of OpMemoryBarrier are dropped during conversion from SPIR-V to OpenCL 1.2/2.0
• OpMemoryBarrier produces always mem_fence, but it should produce atomic_work_item_fence for OpenCL 2.0

I see this assertion when I use -emit-spirv with these files:
https://github.com/KhronosGroup/SPIR/blob/spirv-1.1/test/OpenCL/OpenCL22/attributes/TestBasicAttributes_kernel.cl
https://github.com/KhronosGroup/libclcxx/blob/master/test/libclcxx/address_spaces/TestAddrSpaceTypeConstantPtr.cl
https://github.com/KhronosGroup/libclcxx/blob/master/test/libclcxx/address_spaces/TestAddrSpaceTypeGlobalPtr.cl
https://github.com/KhronosGroup/libclcxx/blob/master/test/libclcxx/address_spaces/TestAddrSpaceTypeLocalPtr.cl
https://github.com/KhronosGroup/libclcxx/blob/master/test/libclcxx/address_spaces/TestAddrSpaceTypePrivatePtr.cl

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=c++ assert-memset.cl
clang: .../llvm/lib/SPIRV/SPIRVWriter.cpp:1296: SPIRV::SPIRVValue* SPIRV::LLVMToSPIRV::transIntrinsicInst(llvm::IntrinsicInst*, SPIRV::SPIRVBasicBlock*):
Assertion `!"Can't translate llvm.memset with non-zero value argument"' failed.

Here my test case:

struct S1
{
    int x;
    int y;
};

S1 foo11()
{
    return S1();
}

S1 foo12()
{
    return S1{};
}

S1 foo13()
{
    S1 s;
    return s;
}

class S2
{
    int x;
    int y;
};

S2 foo21()
{
    return S2();
}

S2 foo22()
{
    return S2{};
}

S2 foo23()
{
    S2 s;
    return s;
}

foo11 - fail
foo12 - ok
foo13 - ok

foo21 - fail
foo22 - fail
foo23 - ok

Full stack trace:
assert-memset.txt
LLVM:
assert-memset.ll.txt

(GCC 6.3.1 on linux and VS2015 on windows in Debug)

Current implementation assumes that last argument of the enqueue_kernel call is bitcast from llvm block type to void()*.
This assumption is not always true;

To reproduce run llvm-as < extExecution.ll.txt | llvm-spir

extExecution.ll.txt
extExecution.cl.txt

llvm-spirv tool crashes during translation of attached kernel.

test_image.zip

Note that tool doesn't crash if sampFun from the example has __kernel qualifier.

https://www.khronos.org/registry/spir-v/specs/1.1/SPIRV.html#OpGetKernelMaxNumSubgroups
https://www.khronos.org/registry/spir-v/specs/1.1/SPIRV.html#OpGetKernelLocalSizeForSubgroupCount

Due to recent changes, a new regression was created.

Please see the attached .cl file
test.zip

Compiled using clang (into .bc file) in command line:
clang -cc1 -internal-isystem C:\llvmspirv\build\Debug\lib\clang\3.6.1\include -nostdsysteminc -x cl -cl-std=CL1.2 -include C:\llvmspirv\build\Debug\lib\clang\3.6.1\include\opencl-12.h -emit-llvm-bc -triple spir64-unknown-unknown -O0

Ran using llvm-spirv in command line:
llvm-spirv -spirv-debug -spirv-text test.bc

Trying to convert following code to SPIR-V throws an exception.

OpenCL-C code:

struct SomeStruct
{
float16 f;
int i;
};

__kernel void test_struct(const __global struct SomeStruct* in, __global struct SomeStruct* out)
{
struct SomeStruct tmp = *in;
*out = tmp;
}

OpenCL-C to LLVM-IR compiled with SPIRV-LLVM/SPIR (clang version 3.6.1 (https://github.com/KhronosGroup/SPIR d7e44c3b27581e54ca0e522987d1ade2bd29b70d) (https://github.com/KhronosGroup/SPIRV-LLVM.git fffc52b7fb3552c045c398d68e22216c8d770c00))
LLVM-IR code (excerpt):

%struct.SomeStruct = type { <16 x float>, i32, [60 x i8] }
define spir_kernel void @test_struct(%struct.SomeStruct addrspace(1)* nocapture readonly %in, %struct.SomeStruct addrspace(1)* nocapture %out) #0 {
  %1 = bitcast %struct.SomeStruct addrspace(1)* %in to i8 addrspace(1)*
  %2 = bitcast %struct.SomeStruct addrspace(1)* %out to i8 addrspace(1)*
  call void @llvm.memmove.p1i8.p1i8.i32(i8 addrspace(1)* %2, i8 addrspace(1)* %1, i32 128, i32 64, i1 false)
  ret void
}

declare void @llvm.memmove.p1i8.p1i8.i32(i8 addrspace(1)* nocapture, i8 addrspace(1)* nocapture readonly, i32, i32, i1)

As you can see, this should be a simple mem-copy, but llvm-spirvthrows:

InvalidFunctionCall: Unexpected llvm intrinsic:  [Src: /opt/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1314  ]
Fails to save LLVM as SPIRV: InvalidFunctionCall: Unexpected llvm intrinsic:  [Src: /opt/SPIRV-LLVM/lib/SPIRV/SPIRVWriter.cpp:1314  ]

pull request #17 caused regression for the following kernel:

__kernel void test_fn( __global char2 *src, __global short *dst )
{
int tid = get_global_id( 0 );
short tmp = as_short( src[ tid ] );
dst[ tid ] = tmp;
}

clang generates the following code:

%call = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2
%sext = shl i64 %call, 32
%idxprom = ashr exact i64 %sext, 32
%arrayidx = getelementptr inbounds <2 x i8> addrspace(1)* %src, i64 %idxprom
%0 = load <2 x i8> addrspace(1)* %arrayidx, align 2, !tbaa !9
%astype = bitcast <2 x i8> %0 to i16
%arrayidx2 = getelementptr inbounds i16 addrspace(1)* %dst, i64 %idxprom
store i16 %astype, i16 addrspace(1)* %arrayidx2, align 2, !tbaa !12
ret void

SPIRV-V reader translates bitcast <2 x i8> %0 to i16 incorrectly to

call spir_func i16 @_Z13convert_shortDv2_c(<2 x i8> %1)

Another example:
__kernel void math_kernel8( __global int8* out, __global float8* in1, __global float8* in2 )
{
int i = get_global_id(0);
out[i] = isequal( in1[i], in2[i] );
}

The LLVM/SPIR-V translator needs to generate conversion instructions for i1 x 8 -> i32 x 8 for representing isequal() as instruction in SPIR-V. SPIR-V reader should not translate it to OCL builtin function.

There is a small bug in the SPIRVEnum.h file (https://github.com/KhronosGroup/SPIRV-LLVM/blob/khronos/spirv-3.6.1/lib/SPIRV/libSPIRV/SPIRVEnum.h#L442) which causes that "Shader" and "Matrix" capabilities are added to the OpenCL kernels when StorageClassFunction is used.
The storage class function is common for shaders and kernels.

BTW, StorageClassPushConstant is missing in the mapping.

Currently we have copy of spirv.hpp file that defines opcodes of SPIR-V instructions. Khronos now provides the latest official version of this header file in the separate repository.

We should integrate Khronos headers into our build process, so we can use the latest version in llvm and clang projects, since the version we are using at the moment is already outdated.

https://github.com/KhronosGroup/SPIRV-Headers

In case of complex constant initialization of automatic variables, SPIRV contains global variable declaration with Storage Class: Function.
ArrInit.txt

CL code example:
kernel void test() { uint arr[] = { 123, 456 }; }

Expected:

         ...
          %9 = OpVariable %8 UniformConstant %7
         %12 = OpFunction %10 None %11
         ...
               OpFunctionEnd

Actual:

         ...
          %9 = OpVariable %8 Function %7
         %12 = OpFunction %10 None %11
         ...
               OpFunctionEnd

I am not sure whether or not this is a real issue, but target spirv of llvm-config returns spirvinfo, spirvdesc and spirvcodegen but not spirvlib. Is it per design or was it just forgotten?

OpDecorationGroup doesn't support more than one decoration at a time. For instance, if multiple instructions shares 2+ decorations (constant, alignment), 2+ decoration groups are created by SPIR-V generator instead of one group that will apply both decorations.

Opened by Alexey Bader.

When generating a SPIR-V binary from the sames files as in #3 (LLVM code and OpenCL kernel), no OpCapability Kernel is generated (required by OpEntryPoint Kernel, OpMemoryModel OpenCL, OpDecorate FuncParamAttr, OpDecorate Constant), nor OpCapability Linkage, which is required by OpDecorate LinkageAttributes

Arguments order of OpBuildNDRange instruction is wrong.
SPIRV Spec
CL code:

kernel void test() {
  ndrange_t ndrange = ndrange_1D(123);
}

Expected:
OpBuildNDRange <Type> <id> GlobalWorkSize LocalWorkSize GlobalWorkOffset

         ...
         %13 = OpConstant %6 123
         %14 = OpConstant %6 0
          %4 = OpFunction %2 None %3
          %5 = OpLabel
         %11 = OpVariable %10 Function
         %12 = OpBuildNDRange %9 %13 %14 %14
               OpStore %11 %12
               OpReturn
               OpFunctionEnd

Actual:
OpBuildNDRange <Type> <id> GlobalWorkOffset GlobalWorkSize LocalWorkSize

         ...
         %13 = OpConstant %6 0
         %14 = OpConstant %6 123
         %15 = OpConstant %6 1
          %4 = OpFunction %2 None %3
          %5 = OpLabel
         %11 = OpVariable %10 Function
         %12 = OpBuildNDRange %9 %13 %14 %15
               OpStore %11 %12
               OpReturn
               OpFunctionEnd

In this SPIR-V binary (produced from this .bc, out of this OpenCL file), OpTypeImage ends up being encoded as 1900 0900 0700 0000 0600 0000 0100 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000. Either the last word is bogus, or it is the AccessQualifier and OpTypeImage is missing it from its operands count.

Also, the same OpTypeImage is defined twice, with the exact same parameters (well, except for the result id):

• 1900 0900 0700 0000 0600 0000 0100 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000
• 1900 0900 1500 0000 0600 0000 0100 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000

Translating this LLVM code using llvm-spirv results in this SPIR-V binary, where the OpDecorate LinkageAttributes only has the linkage type but no name. (The corresponding binary stream is 4700 0400 0500 0000 2900 0000 0000 0000.)

The code assumes the linkage type is the first attribute, which seems wrong when looking at the specification. Furthermore, there is no way to set both attributes as only the SPIRVDecorate* constructor modifies the Literals vector storing the attributes, and it only takes one word.

Note: If needed, the initial OpenCL kernel can be found here.

SPRI-V converter hits the following assert during conversion from the attached LLVM IR to SPIR-V
Assertion `HasVariWC && WC >= WordCount && "Invalid word count"' failed.
To reproduce:
$ llvm-as enqueue.ll.txt -o enqueue.bc && llvm-spirv enqueue.bc
Call stack enqueue.bt.txt

The problematic LLVM IL enqueue.ll.txt is compiled from the following OpenCL C 2.0 source code
void __kernel one(__global float * inout) {
_inout = cos(_inout);
}

void __kernel two(__global float * inout) {
enqueue_kernel(get_default_queue(), CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange_1D(1), ^{one(inout);});
}

My understanding is that OpName is suppose to keep information about source level names, but it looks like it currently refers to LLVM IR variables names, which makes it useless.

Opened by Alexey Bader.

There are 2 issues seen during translation of atomic_compare_exchange_strong and atomic_compare_exchange_weak calls to SPIR-V:

1. 5th operand of Inst is value, and its type should be scalar value not the pointer
2. AtomicCompareExchange must return the same type as Value but it returns bool
3. Operands are incorrectly ordered
   Fix for the issues mentioned above were implemented in pull request #26.
   Translation of OpAtomicCompareExchange[Weak] instruction to atomic_compare_exchange_[strong|weak] call has to be fixed. Argument types, argument order and different semantics of the operation in SPIR-V and OpenCL 2.0 have to be taken into account.

SPIR-V reader produces incorrectly mangled enqueue_kernel built-in. I suppose what some other built-ins are also affected. Compare the reproducer output and the mangled name produced by the SPIR 2.0 generator (I expect SPIR-V generator committed by Alexey Bader uses the same mangling).

repro.zip.
$ ./llvm-spirv repro.bc && ./llvm-spirv -r repro.spv -o out.bc && ./llvm-dis < out.bc | grep "declare.*enqueue_kernel"
$ _Z14enqueue_kernel9ocl_queue22kernel_enqueue_flags_tP9ndrange_tjPU3AS412ocl_clkeventS2_U13block_pointerFvvE

SPIR 2.0 name:
_Z14enqueue_kernel9ocl_queuei9ndrange_tjPKU3AS413ocl_clk_eventPU3AS413ocl_clk_eventU13block_pointerFvvE
taken from here: https://github.com/KhronosGroup/SPIR-Tools/wiki/SPIR-2.0-built-in-functions#enqueuing-kernels

I see here few discrepancies:

1. flags are mangled as int by SPIR 2.0 -> "i" VS enum mangling by the reader -> "22kernel_enqueue_flags_t" (which is rather bug in SPIR 2.0 generator)
2. ndrange is passed by value so should be mangled as 9ndrange_t but it is mangled as a pointer to ndrange_t (P9ndrange_t)
3. event_wait_list should be mangled as a pointer to constant (i.e. PKU3AS413ocl_clk_event)
4. event_ret is mangled as S2_ by the reader
5. 12ocl_clkevent instead of 13ocl_clk_event

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)

Not sure if Vulkan is currently supported or not, but it seems only OpenCL is?

In this SPIR-V binary (produced from this .bc, out of this OpenCL file), OpTypeImage is defined twice, with the exact same parameters (well, except for the result id):

• 1900 0900 0700 0000 0600 0000 0100 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000
• 1900 0900 1500 0000 0600 0000 0100 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000 0000

For example there is undefined behavior in SPIRVMap with using static maps if converter works in parallel.

While compiling LLVM, I got these errors:

[12:06:51][Step 1/2] Generating Mangler.cpp dependencies ...
[12:06:51][Step 1/2] ../../../Mangler/Mangler.cpp:17:27: fatal error: SPIRVInternal.h: No such file or directory
[12:06:51][Step 1/2] #include "SPIRVInternal.h"
[12:06:51][Step 1/2] ^
[12:06:51][Step 1/2] compilation terminated.

Also:

[12:08:15][Step 1/2] Compiling SPIRVUtil.cpp ...
[12:08:18][Step 1/2] ../../../SPIRVUtil.cpp: In function 'std::__1::string SPIRV::getSPIRVImageSampledTypeName(SPIRV::SPIRVType*)':
[12:08:18][Step 1/2] ../../../SPIRVUtil.cpp:1250:8: error: suggest explicit braces to avoid ambiguous 'else' [-Werror=parentheses]
[12:08:18][Step 1/2] if (Ty->getIntegerBitWidth() == 32)
[12:08:18][Step 1/2] ^
[12:08:20][Step 1/2] cc1plus: all warnings being treated as errors

After compiling the attached OpenCL kernel using clang -cc1 -emit-llvm-bc -triple spir-unknown-unknown -cl-spir-compile-options "" -include /usr/local/include/opencl_spir.h -o hello_word.bc ~/softwares/compute_test/hello_word.cl, I ran llvm-spirv on the generated llvm code, which produced the attached hello_world.spv file. Feeding it to spirv-dis from the SPIRV-Tools repository resulted in the error "error: 15: Error: Result Id is 0".
The result id of OpExtInstImport is indeed 0 in the SPIR-V binary, which is invalid by the SPIR-V spec, as IDs should be the interval 0 < id < max_bound.

While trying to compile the new spirv-target branch, I get the following error:

In file included from ../lib/Target/SPIRV/SPIRVISelDAGToDAG.cpp:44:
../lib/Target/SPIRV/SPIRVTargetMachine.h:41:10: fatal error: 'SPIRVSubTarget.h' file not found
#include "SPIRVSubTarget.h"

@yxsamliu Did you forget to push this file? :-)

when a kernel uses the llvm.fmuladd intrinsic we're not generating correct code.

The generated SPIR-V is this:

     OpName %8 "llvm.fmuladd.f32"
     OpDecorate %8 LinkageAttributes "llvm.fmuladd.f32" Import
%8 = OpFunction %6 Pure %7

%46 = OpFunctionCall %6 %8 %65 %47 %20

I think llvm-spir needs to emit a opfmul + opfadd.

See attached for an example.

FancyJuliaSet.txt

code to reproduce:

float negate(float in)
{
    return -in;
}

Compilation commands:

CLANG=$WORK_DIR/local/bin/clang++
LLVM_SPIRV=$WORK_DIR/local/bin/llvm-spirv

$CLANG -cc1 -emit-llvm-bc -triple spir64-unknown-unknown -include $SPIR_INC_FILE -o $FILE_NAME.bc $FILE_NAME.cl
$LLVM_SPIRV $FILE_NAME.bc -o $FILE_NAME.spv64

Error caught running it through the spirv-val validator

Trips assert assert(Ch == '\0' && "Invalid string in SPIRV") in const SPIRVDecoder& operator>>(const SPIRVDecoder&I, std::string& Str) in SPIRVStream.cpp

This is because the std::stream I.IS does not have std::noskipws set

It is perfectly legal to have spaces in the names of llvm type names.

kernels_ocl120_64-nocrash.spv.txt
kernels_ocl120_64-crash.spv.txt

The original OpenCL-C code:

struct SomeStruct
{
float16 f;
int i;
};

__kernel void test_struct(const __global struct SomeStruct* in, __global struct SomeStruct* out)
{
struct SomeStruct tmp = *in;
tmp.i = 42;
*out = tmp;
}

The LLVM-IR code (excerpt), compiled with clang version 3.6.1 (https://github.com/KhronosGroup/SPIR d7e44c3b27581e54ca0e522987d1ade2bd29b70d) (https://github.com/KhronosGroup/SPIRV-LLVM.git fffc52b7fb3552c045c398d68e22216c8d770c00):

%struct.SomeStruct = type { <16 x float>, i32, [60 x i8] }

define spir_kernel void @test_struct(%struct.SomeStruct addrspace(1)* nocapture readonly %in, %struct.SomeStruct addrspace(1)* nocapture %out) #0 {
  %tmp.sroa.5 = alloca [60 x i8], align 4
  %1 = getelementptr inbounds [60 x i8]* %tmp.sroa.5, i32 0, i32 0
  call void @llvm.lifetime.start(i64 60, i8* %1)
  %2 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %in, i32 0, i32 0
  %3 = load <16 x float> addrspace(1)* %2, align 64
  %4 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %in, i32 0, i32 2, i32 0
  call void @llvm.memcpy.p0i8.p1i8.i32(i8* %1, i8 addrspace(1)* %4, i32 60, i32 4, i1 false)
  %5 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 0
  store <16 x float> %3, <16 x float> addrspace(1)* %5, align 64
  %6 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 1
  store i32 42, i32 addrspace(1)* %6, align 64
  %7 = getelementptr inbounds %struct.SomeStruct addrspace(1)* %out, i32 0, i32 2, i32 0
  call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* %7, i8* %1, i32 60, i32 4, i1 false)
  call void @llvm.lifetime.end(i64 60, i8* %1)
  ret void
}
declare void @llvm.lifetime.start(i64, i8* nocapture)
declare void @llvm.memcpy.p0i8.p1i8.i32(i8* nocapture, i8 addrspace(1)* nocapture readonly, i32, i32, i1)
declare void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)* nocapture, i8* nocapture readonly, i32, i32, i1)
declare void @llvm.lifetime.end(i64, i8* nocapture)

llvm-spirv crashes with following stack-trace:

Stack dump:
0.	Program arguments: /opt/SPIRV-LLVM/build/bin/llvm-spirv -o /tmp/out.spv 
0  llvm-spirv      0x00000000005f5972
1  llvm-spirv      0x00000000005f4241
2  libpthread.so.0 0x00007f09af1ab5c0
3  llvm-spirv      0x00000000004b0d70
4  llvm-spirv      0x00000000004b1683
5  llvm-spirv      0x00000000004acefd
6  llvm-spirv      0x00000000004ad256
7  llvm-spirv      0x00000000004b1f08
8  llvm-spirv      0x00000000004b2240
9  llvm-spirv      0x00000000004b226b
10 llvm-spirv      0x000000000058bcac
11 llvm-spirv      0x00000000004b22ce
12 llvm-spirv      0x0000000000407a73
13 libc.so.6       0x00007f09ae54c401 __libc_start_main + 241
14 llvm-spirv      0x000000000040abea
Segmentation fault (core dumped)

Analysis with valgrind prints this:

Process terminating with default action of signal 11 (SIGSEGV)
Access not within mapped region at address 0x8
  at 0x4CED70: SPIRV::SPIRVEntry::setModule(SPIRV::SPIRVModule*) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x45F298: SPIRV::SPIRVDecoder::getEntry() (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x44793F: SPIRV::operator>>(std::istream&, SPIRV::SPIRVModule&) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x458F13: SPIRV::ConvertSPIRV(std::istream&, llvm::raw_ostream&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >&, bool, bool) (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x40AFBB: convertSPIRV()::{lambda(llvm::raw_ostream&)#1}::operator()(llvm::raw_ostream&) const (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x40B2DB: convertSPIRV() (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)
  by 0x40798B: main (in /opt/SPIRV-LLVM/build/bin/llvm-spirv)

... which looks like access to a nullptr

When generating a SPIR-V binary from the sames files as in #3 (LLVM code and OpenCL kernel), get_local_id(X) is marked as an exported function, which seems wrong as the kernel doesn't define but rather import it.

LIT test fails on RH6.5 on assertion in llvm::Value::setName (on debug build).

To reproduce:
llvm-lit -v path-to-llvm/test/SPIRV/image-unoptimized.ll
image-unoptimized.bt.txt

eraseUselessFunctions erases all function declarations with no use. This is too generic and may not be user's intention. Should be replaced with a more specific one which only erase useless functions generated by converter.

Reproducer:
set PATH_TO_GEN to spirv-1.0 generator install directory
And run (in bash)
${PATH_TO_GEN}/bin/clang -cc1 -x cl -cl-std=CL2.0 -triple spir64-unknonw-unknown -emit-spirv -include ${PATH_TO_GEN}/lib/clang/3.6.1/include/opencl-20.h repro.cl -o /tmp/out.spv && ./llvm-spirv -to-text /tmp/out.spv -o - | grep -i default

Expected:
GetDefaultQueue operation
Actual:
7 Name 4 "get_default_queue"
9 Decorate 4 LinkageAttributes "get_default_queue" Import

Hi, I just want to compiler a simple vector add example(adapted from NVIDIA's example) with dynamic parallelism in OpenCL 2.x but end up getting crashed.
Here is the CL kernel code:

__kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements)
{
    int tid = get_global_id(0);
    device_queue devQ = get_default_queue();
    ndrange ndrange1(32);

    auto myblock = [=](void)->void{
      int iGID = get_global_id(0) + tid * 32;
      // add the vector elements
      c[iGID] = a[iGID] + b[iGID];
    };

    auto err_ret = devQ.enqueue_kernel(enqueue_policy::wait_kernel,
                                      ndrange1,
                                      myblock);

}

The error I got is:

clang: /path/to/llvm/lib/SPIRV/SPIRVWriter.cpp:464: SPIRV::SPIRVType* SPIRV::LLVMToSPIRV::transType(llvm::Type*): Assertion `!ET->isFunctionTy() && "Function pointer type is not allowed"' failed.

After I turn on the SPIRV debugging message, I got more detailed information about the location of the error:

[mapValue]   call spir_func void @"_ZNU3AS42cl9__details16__enqueue_helperIZ9VectorAddE3$_0NS0_8__paramsIJEEENS0_5__seqImJEEES6_S4_S4_E32__get_enqueue_kernel_static_dataEOU3AS4S2_OU3AS4NS_5tupleIJEEE"(%class.anon.0 addrspace(4)* sret %1, %class.anon addrspace(4)* dereferenceable(32) %0, %"class.cl::tuple<>" addrspace(4)* dereferenceable(1) %args) => 99�����
[transValue]   %call2 = call spir_func void (i8 addrspace(4)*)* ()* @"_ZNU3AS42cl9__details16__enqueue_helperIZ9VectorAddE3$_0NS0_8__paramsIJEEENS0_5__seqImJEEES6_S4_S4_E28__get_enqueue_kernel_wrapperEv"()
[transType] void (i8 addrspace(4)*)* ()
[transType] void (i8 addrspace(4)*)*

It seems that the problem is raised in libclcxx. When the SPIRV backend use transType(llvm::Type*) to transform function parameters, transType(llvm::Type*) seem to assert that if the input Type is pointer type, then it must not be function pointer type. But it's a little bit strange since there are many chances in OpenCL 2.x that function parameters are lambda functions or function pointers.

Quote from SPIR-V specification section 2.4:
"9. All type declarations (OpTypeXXX instructions), all constant instructions, and all global variable declarations (all OpVariable instructions whose Storage Class is not Function). All operands in all these instructions must be declared before being used. Otherwise, they can be in any order. This section is also the first section to allow use of OpLine debug information."

LLVM/SPIR-V translator place type declarations first, constant declarations next and global variables after that. That order violates requirement "All operands in all these instructions must be declared before being used." for the following OpenCL code:

constant float f[2] = {0.f, 0.f};

SPIR-V code produced by translator:

     %2 = OpTypeInt 32 0
     %4 = OpTypeFloat 32
     %5 = OpTypeArray %4 %3 ; <<< here we %3 is used but defined below.
     %7 = OpTypePointer UniformConstant %5
     %3 = OpConstant %2 2
     %6 = OpConstantNull %5
     %8 = OpVariable %7 UniformConstant %6

Attached is an example of the code where LLVM/SPIR-V converter translates sampler initializer to OpConstant instead of expected OpConstantSampler.

test.zip

The documentation (https://github.com/KhronosGroup/SPIRV-LLVM/blob/khronos/spirv-3.6.1/docs/SPIRVRepresentationInLLVM.rst#id9) says that the built-in variables should have names:
__spirv_BuiltIn{Name}

In the SPIR-V converter (https://github.com/KhronosGroup/SPIRV-LLVM/blob/khronos/spirv-3.6.1/lib/SPIRV/libSPIRV/SPIRVEnum.h#L229) these built-ins don't have "BuiltIn" prefix.
Is this generator bug or documentation?

Personally I'd prefer the names with "BuiltIn" prefix.