intel · AlexeySotkin · Jul 16, 2019
@@ -87,7 +87,7 @@ using namespace OCLUtil;
 
 namespace SPIRV {
 
-cl::opt<bool> SPIRVMemToReg("spirv-mem2reg", cl::init(true),
+cl::opt<bool> SPIRVMemToReg("spirv-mem2reg", cl::init(false),
                             cl::desc("LLVM/SPIR-V translation enable mem2reg"));
 
 cl::opt<bool> SPIRVNoDerefAttr(
@@ -258,7 +258,7 @@ SPIRVType *LLVMToSPIRV::transType(Type *T) {
     return mapType(T, BM->addFloatType(T->getPrimitiveSizeInBits()));
 
   // A pointer to image or pipe type in LLVM is translated to a SPIRV
-  // sampler or pipe type.
+  // (non-pointer) image or pipe type.
   if (T->isPointerTy()) {
     auto ET = T->getPointerElementType();
     assert(!ET->isFunctionTy() && "Function pointer type is not allowed");

@@ -11,9 +11,6 @@ target triple = "spir-unknown-unknown"
 ; Int64Atomics capability must be declared only if atomic builtins have 64-bit integers arguments.
 ; CHECK-NOT: Capability Int64Atomics
 
-; CHECK: Name [[Pointer:[0-9]+]] "object"
-; CHECK: Name [[ComparatorPtr:[0-9]+]] "expected"
-; CHECK: Name [[Value:[0-9]+]] "desired"
 ; CHECK: 4 TypeInt [[int:[0-9]+]] 32 0
 ; CHECK: Constant [[int]] [[DeviceScope:[0-9]+]] 1
 ; CHECK: Constant [[int]] [[SequentiallyConsistent_MS:[0-9]+]] 16
@@ -22,9 +19,9 @@ target triple = "spir-unknown-unknown"
 
 ; Function Attrs: nounwind
 define spir_func void @test(i32 addrspace(4)* %object, i32 addrspace(4)* %expected, i32 %desired) #0 {
-; CHECK: FunctionParameter [[int_ptr]] [[Pointer]]
-; CHECK: FunctionParameter [[int_ptr]] [[ComparatorPtr]]
-; CHECK: FunctionParameter [[int]] [[Value]]
+; CHECK: FunctionParameter [[int_ptr]] [[object:[0-9]+]]
+; CHECK: FunctionParameter [[int_ptr]] [[expected:[0-9]+]]
+; CHECK: FunctionParameter [[int]] [[desired:[0-9]+]]
 
 entry:
   %object.addr = alloca i32 addrspace(4)*, align 4
@@ -39,11 +36,16 @@ entry:
   %0 = load i32 addrspace(4)*, i32 addrspace(4)** %object.addr, align 4
   %1 = load i32 addrspace(4)*, i32 addrspace(4)** %expected.addr, align 4
   %2 = load i32, i32* %desired.addr, align 4
-
-  %call = call spir_func zeroext i1 @_Z30atomic_compare_exchange_strongPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %0, i32 addrspace(4)* %1, i32 %2)
-; CHECK: Load [[int]] [[Comparator:[0-9]+]] [[ComparatorPtr]]
+; CHECK: Store [[object_addr:[0-9]+]] [[object]]
+; CHECK: Store [[expected_addr:[0-9]+]] [[expected]]
+; CHECK: Store [[desired_addr:[0-9]+]] [[desired]]
+; CHECK: Load [[int_ptr]] [[Pointer:[0-9]+]] [[object_addr]]
+; CHECK: Load [[int_ptr]] [[exp:[0-9]+]] [[expected_addr]]
+; CHECK: Load [[int]] [[Value:[0-9]+]] [[desired_addr]]
+; CHECK: Load [[int]] [[Comparator:[0-9]+]] [[exp]]
 ; CHECK-NEXT: 9 AtomicCompareExchange [[int]] [[Result:[0-9]+]] [[Pointer]] [[DeviceScope]] [[SequentiallyConsistent_MS]] [[SequentiallyConsistent_MS]] [[Value]] [[Comparator]]
-; CHECK-NEXT: Store [[ComparatorPtr]] [[Result]]
+  %call = call spir_func zeroext i1 @_Z30atomic_compare_exchange_strongPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %0, i32 addrspace(4)* %1, i32 %2)
+; CHECK-NEXT: Store [[exp]] [[Result]]
 ; CHECK-NEXT: IEqual [[bool]] [[CallRes:[0-9]+]] [[Result]] [[Comparator]]
 ; CHECK-NOT: [[Result]]
   %frombool = zext i1 %call to i8
@@ -57,10 +59,13 @@ entry:
   %5 = load i32 addrspace(4)*, i32 addrspace(4)** %expected.addr, align 4
   %6 = load i32, i32* %desired.addr, align 4
 
+; CHECK: Load [[int_ptr]] [[Pointer:[0-9]+]] [[object_addr]]
+; CHECK: Load [[int_ptr]] [[exp:[0-9]+]] [[expected_addr]]
+; CHECK: Load [[int]] [[Value:[0-9]+]] [[desired_addr]]
+; CHECK: Load [[int]] [[ComparatorWeak:[0-9]+]] [[exp]]
   %call2 = call spir_func zeroext i1 @_Z28atomic_compare_exchange_weakPVU3AS4U7_AtomiciPU3AS4ii(i32 addrspace(4)* %4, i32 addrspace(4)* %5, i32 %6)
-; CHECK: Load [[int]] [[ComparatorWeak:[0-9]+]] [[ComparatorPtr]]
 ; CHECK-NEXT: 9 AtomicCompareExchangeWeak [[int]] [[Result:[0-9]+]] [[Pointer]] [[DeviceScope]] [[SequentiallyConsistent_MS]] [[SequentiallyConsistent_MS]] [[Value]] [[ComparatorWeak]]
-; CHECK-NEXT: Store [[ComparatorPtr]] [[Result]]
+; CHECK-NEXT: Store [[exp]] [[Result]]
 ; CHECK-NEXT: IEqual [[bool]] [[CallRes:[0-9]+]] [[Result]] [[ComparatorWeak]]
 ; CHECK-NOT: [[Result]]
 

@@ -8,7 +8,7 @@
 // CHECK-SPIRV: Label
 // CHECK-SPIRV: ExtInst {{.*}} DebugScope
 // CHECK-SPIRV: ExtInst {{.*}} sin
-// CHECK-LLVM: call spir_func float @_Z3sinf(float %x) {{.*}} !dbg ![[loc:[0-9]+]]
+// CHECK-LLVM: call spir_func float @_Z3sinf(float %{{.*}}) {{.*}} !dbg ![[loc:[0-9]+]]
 // CHECK-LLVM: ![[loc]] = !DILocation(line: 14, column: 10, scope: !{{.*}})
 float f(float x) {
   return sin(x);

@@ -1,7 +1,7 @@
 // Check that we can translate llvm.dbg.declare for a local variable which was
-// deleted by mem2reg pass(enabled by default in llvm-spirv)
+// deleted by mem2reg pass(disabled by default in llvm-spirv)
 
-// RUN: %clang_cc1 %s -triple spir -disable-llvm-passes -debug-info-kind=standalone -emit-llvm-bc -o - | llvm-spirv -o %t.spv
+// RUN: %clang_cc1 %s -triple spir -disable-llvm-passes -debug-info-kind=standalone -emit-llvm-bc -o - | llvm-spirv -spirv-mem2reg -o %t.spv
 // RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 // RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM
 

@@ -30,7 +30,7 @@ target triple = "spir64-unknown-unknown"
 
 ;CHECK-LLVM: test_32
 ;CHECK-LLVM: entry
-;CHECK-LLVM: switch i32 %conv, label %sw.epilog
+;CHECK-LLVM: switch i32 %0, label %sw.epilog
 ;CHECK-LLVM: i32 0, label %sw.bb
 ;CHECK-LLVM: i32 1, label %sw.bb1
 

@@ -33,7 +33,7 @@ target triple = "spir64-unknown-unknown"
 
 ;CHECK-LLVM: test_64
 ;CHECK-LLVM: entry
-;CHECK-LLVM: switch i64 %call, label %sw.epilog [
+;CHECK-LLVM: switch i64 %0, label %sw.epilog [
 ;CHECK-LLVM: i64 0, label %sw.bb
 ;CHECK-LLVM: i64 1, label %sw.bb1
 ;CHECK-LLVM: i64 21474836481, label %sw.bb3

@@ -0,0 +1,47 @@
+// RUN: %clang_cc1 %s -triple spir -O0 -emit-llvm-bc -o %t.bc
+// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s
+// RUN: spirv-val %t.spv
+
+// CHECK: TypeImage [[TypeImage:[0-9]+]]
+// CHECK: TypeSampler [[TypeSampler:[0-9]+]]
+// CHECK: TypePointer [[TypeImagePtr:[0-9]+]] {{[0-9]+}} [[TypeImage]]
+// CHECK: TypePointer [[TypeSamplerPtr:[0-9]+]] {{[0-9]+}} [[TypeSampler]]
+
+// CHECK: FunctionParameter [[TypeImage]] [[srcimg:[0-9]+]]
+// CHECK: FunctionParameter [[TypeSampler]] [[sampler:[0-9]+]]
+
+// CHECK: Variable [[TypeImagePtr]] [[srcimg_addr:[0-9]+]]
+// CHECK: Variable [[TypeSamplerPtr]] [[sampler_addr:[0-9]+]]
+
+// CHECK: Store [[srcimg_addr]] [[srcimg]]
+// CHECK: Store [[sampler_addr]] [[sampler]]
+
+// CHECK: Load {{[0-9]+}} [[srcimg_val:[0-9]+]] [[srcimg_addr]]
+// CHECK: Load {{[0-9]+}} [[sampler_val:[0-9]+]] [[sampler_addr]]
+
+// CHECK: SampledImage {{[0-9]+}} {{[0-9]+}} [[srcimg_val]] [[sampler_val]]
+// CHECK-NEXT: ImageSampleExplicitLod
+
+// CHECK: Load {{[0-9]+}} [[srcimg_val:[0-9]+]] [[srcimg_addr]]
+// CHECK: ImageQuerySizeLod {{[0-9]+}} {{[0-9]+}} [[srcimg_val]]
+
+// Excerpt from opencl-c-base.h
+typedef float float4 __attribute__((ext_vector_type(4)));
+typedef int int2 __attribute__((ext_vector_type(2)));
+typedef __SIZE_TYPE__ size_t;
+
+// Excerpt from opencl-c.h to speed up compilation.
+#define __ovld __attribute__((overloadable))
+#define __purefn __attribute__((pure))
+#define __cnfn __attribute__((const))
+size_t __ovld __cnfn get_global_id(unsigned int dimindx);
+int __ovld __cnfn get_image_width(read_only image2d_t image);
+float4 __purefn __ovld read_imagef(read_only image2d_t image, sampler_t sampler, int2 coord);
+
+
+__kernel void test_fn(image2d_t srcimg, sampler_t sampler, global float4 *results) {
+  int tid_x = get_global_id(0);
+  int tid_y = get_global_id(1);
+  results[tid_x + tid_y * get_image_width(srcimg)] = read_imagef(srcimg, sampler, (int2){tid_x, tid_y});
+}