Skip to content

Disabling mem2reg by default #349

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand Down Expand Up @@ -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");
Expand Down
29 changes: 17 additions & 12 deletions llvm-spirv/test/AtomicCompareExchange_cl20.ll
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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]]

Expand Down
2 changes: 1 addition & 1 deletion llvm-spirv/test/DebugInfo/BuiltinCallLocation.cl
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
4 changes: 2 additions & 2 deletions llvm-spirv/test/DebugInfo/DebugDeclareUnused.cl
Original file line number Diff line number Diff line change
@@ -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

Expand Down
2 changes: 1 addition & 1 deletion llvm-spirv/test/OpSwitch32.ll
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
2 changes: 1 addition & 1 deletion llvm-spirv/test/OpSwitch64.ll
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
47 changes: 47 additions & 0 deletions llvm-spirv/test/image-unoptimized.cl
Original file line number Diff line number Diff line change
@@ -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});
}
77 changes: 0 additions & 77 deletions llvm-spirv/test/image-unoptimized.ll

This file was deleted.