Skip to content

Commit fa89ada

Browse files
bb-syclbb-sycl
committed
[SPIR-V] Fix opaque pointers support in intel/llvm/llvm-spirv
It cherry-picks several missing patches from KHR translator
2 parents 1a8941c + 52d07d6 commit fa89ada

File tree

53 files changed

+100
-138
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

53 files changed

+100
-138
lines changed

‎llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -4097,11 +4097,7 @@ SPIRVValue *LLVMToSPIRVBase::transIntrinsicInst(IntrinsicInst *II,
40974097
return DbgTran->createDebugValuePlaceholder(cast<DbgValueInst>(II), BB);
40984098
case Intrinsic::annotation: {
40994099
SPIRVType *Ty = transScavengedType(II);
4100-
4101-
GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(II->getArgOperand(1));
4102-
if (!GEP)
4103-
return nullptr;
4104-
Constant *C = cast<Constant>(GEP->getOperand(0));
4100+
Constant *C = cast<Constant>(II->getArgOperand(1)->stripPointerCasts());
41054101
StringRef AnnotationString;
41064102
if (!getConstantStringInfo(C, AnnotationString))
41074103
return nullptr;

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_blocking_pipes/PipeBlocking.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,10 +3,10 @@
33
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
44
; FIXME: add more negative test cases
55
; RUN: llvm-spirv %t.spt -to-binary -o %t.spv
6-
; RUN: llvm-spirv -r %t.spv -o %t.bc
6+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc
77
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

9-
; RUN: llvm-spirv -r %t.spv -o %t.bc --spirv-target-env=SPV-IR
9+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc --spirv-target-env=SPV-IR
1010
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
1111

1212
; ModuleID = 'test/CodeGenOpenCL/pipe_builtin.cl'

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_device_side_avc_motion_esimation/subgroup_avc_intel_generic.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,9 +3,9 @@
33
// RUN: llvm-spirv %t.spv --to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44
// There is no validation for SPV_INTEL_device_side_avc_motion_estimation implemented in
55
// SPIRV-Tools. TODO: spirv-val %t.spv
6-
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
6+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
77
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefixes=CHECK-LLVM-COMMON,CHECK-LLVM
8-
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
8+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
99
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefixes=CHECK-LLVM-COMMON,CHECK-LLVM-SPIRV
1010
// RUN: llvm-spirv %t.rev.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_device_side_avc_motion_estimation -o %t.spv
1111
// RUN: llvm-spirv %t.spv --to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_device_side_avc_motion_esimation/subgroup_avc_intel_types.spt

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -93,9 +93,9 @@
9393
1 FunctionEnd
9494

9595
; RUN: llvm-spirv %s -to-binary -o %t.spv
96-
; RUN: llvm-spirv -r %t.spv -o %t.bc
96+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc
9797
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefixes=CHECK-COMMON,CHECK-LLVM
98-
; RUN: llvm-spirv -r %t.spv -o %t.bc --spirv-target-env=SPV-IR
98+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc --spirv-target-env=SPV-IR
9999
; RUN: llvm-dis -opaque-pointers=0 < %t.bc | FileCheck %s --check-prefixes=CHECK-COMMON,CHECK-LLVM-SPIRV
100100

101101
; CHECK-LLVM: %opencl.intel_sub_group_avc_mce_payload_t = type opaque

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_device_side_avc_motion_esimation/subgroup_avc_intel_vme_image.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,9 +3,9 @@
33
// RUN: llvm-spirv %t.spv --to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44
// There is no validation for SPV_INTEL_device_side_avc_motion_estimation implemented in
55
// SPIRV-Tools. TODO: spirv-val %t.spv
6-
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
6+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
77
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefixes=CHECK-LLVM-COMMON,CHECK-LLVM
8-
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
8+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
99
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefixes=CHECK-LLVM-COMMON,CHECK-LLVM-SPIRV
1010
// RUN: llvm-spirv %t.rev.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_device_side_avc_motion_estimation -o %t.rev.spv
1111
// RUN: llvm-spirv %t.rev.spv --to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_loop_controls/FPGAIVDepLoopAttrOnClosure.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -41,7 +41,7 @@
4141
; RUN: llvm-spirv -to-text %t.spv -o %t.spt
4242
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
4343

44-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
44+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
4545
; RUN: llvm-dis -opaque-pointers=0 %t.rev.bc -o %t.rev.ll
4646

4747
; CHECK-LLVM is the base prefix, which includes simple checks for

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/IntelFPGAMemoryAccesses.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -41,8 +41,8 @@
4141
; return 0;
4242
; }
4343

44-
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
45-
; RUN: llvm-spirv %t.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_fpga_memory_accesses -o %t.spv
44+
; RUN: llvm-as %s -o %t.bc
45+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_memory_accesses -o %t.spv
4646
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
4747

4848
; RUN: llvm-spirv -r -emit-opaque-pointers %t.spv -o %t.rev.bc

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_accesses/intel_fpga_lsu_optimized.ll

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -49,12 +49,12 @@
4949
; }
5050

5151
; Check that translation of optimized IR doesn't crash:
52-
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
53-
; RUN: llvm-spirv %t.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_fpga_memory_accesses -o %t.spv
52+
; RUN: llvm-as %s -o %t.bc
53+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_memory_accesses -o %t.spv
5454

5555
; Check that reverse translation restore ptr.annotations correctly:
56-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
57-
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
56+
; RUN: llvm-spirv -r -emit-opaque-pointers %t.spv -o %t.rev.bc
57+
; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
5858

5959
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
6060
target triple = "spir64-unknown-unknown"
@@ -87,33 +87,33 @@ entry:
8787
%4 = addrspacecast i32 addrspace(1)* %add.ptr.i27 to i32 addrspace(4)*
8888
%5 = tail call dereferenceable(4) i32 addrspace(4)* @llvm.ptr.annotation.p4i32(i32 addrspace(4)* %4, i8* getelementptr inbounds ([26 x i8], [26 x i8]* @.str, i64 0, i64 0), i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str.1, i64 0, i64 0), i32 0, i8* null) #2
8989
%6 = load i32, i32 addrspace(4)* %5, align 4, !tbaa !9
90-
; CHECK-LLVM: [[PTR_i27:[%0-9a-z.]+]] = getelementptr inbounds i32, i32 addrspace(1)* {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
91-
; CHECK-LLVM: [[PTR_i:[%0-9a-z.]+]] = getelementptr inbounds i32, i32 addrspace(1)* {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
92-
; CHECK-LLVM: [[PTR_i27_AS_CAST:[%0-9a-z.]+]] = addrspacecast i32 addrspace(1)* [[PTR_i27]] to i32 addrspace(4)*
93-
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call i32 addrspace(4)* @llvm.ptr.annotation.p4i32.p0i8(i32 addrspace(4)* [[PTR_i27_AS_CAST]], i8* getelementptr inbounds ({{.*}} [[PTR_i27_ANNOT_STR]]
94-
; CHECK-LLVM: [[PTR_RESULT_LOAD:[%0-9a-z.]+]] = load i32, i32 addrspace(4)* [[PTR_ANNOT_CALL]]
90+
; CHECK-LLVM: [[PTR_i27:[%0-9a-z.]+]] = getelementptr inbounds i32, ptr addrspace(1) {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
91+
; CHECK-LLVM: [[PTR_i:[%0-9a-z.]+]] = getelementptr inbounds i32, ptr addrspace(1) {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
92+
; CHECK-LLVM: [[PTR_i27_AS_CAST:[%0-9a-z.]+]] = addrspacecast ptr addrspace(1) [[PTR_i27]] to ptr addrspace(4)
93+
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p0(ptr addrspace(4) [[PTR_i27_AS_CAST]], ptr [[PTR_i27_ANNOT_STR]]
94+
; CHECK-LLVM: [[PTR_RESULT_LOAD:[%0-9a-z.]+]] = load i32, ptr addrspace(4) [[PTR_ANNOT_CALL]]
9595
%add.ptr.i15.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i27, i64 1
9696
%7 = addrspacecast i32 addrspace(1)* %add.ptr.i15.i to i32 addrspace(4)*
9797
%8 = tail call dereferenceable(4) i32 addrspace(4)* @llvm.ptr.annotation.p4i32(i32 addrspace(4)* %7, i8* getelementptr inbounds ([28 x i8], [28 x i8]* @.str.2, i64 0, i64 0), i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str.1, i64 0, i64 0), i32 0, i8* null) #2
9898
%9 = load i32, i32 addrspace(4)* %8, align 4, !tbaa !9
99-
; CHECK-LLVM: [[PTR_i15_i:[%0-9a-z.]+]] = getelementptr inbounds i32, i32 addrspace(1)* {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
100-
; CHECK-LLVM: [[PTR_i15_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast i32 addrspace(1)* [[PTR_i15_i]] to i32 addrspace(4)*
101-
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call i32 addrspace(4)* @llvm.ptr.annotation.p4i32.p0i8(i32 addrspace(4)* [[PTR_i15_i_AS_CAST]], i8* getelementptr inbounds ({{.*}} [[PTR_i15_i_ANNOT_STR]]
102-
; CHECK-LLVM: [[PTR_RESULT_LOAD_1:[%0-9a-z.]+]] = load i32, i32 addrspace(4)* [[PTR_ANNOT_CALL]]
99+
; CHECK-LLVM: [[PTR_i15_i:[%0-9a-z.]+]] = getelementptr inbounds i32, ptr addrspace(1) {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
100+
; CHECK-LLVM: [[PTR_i15_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast ptr addrspace(1) [[PTR_i15_i]] to ptr addrspace(4)
101+
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p0(ptr addrspace(4) [[PTR_i15_i_AS_CAST]], ptr [[PTR_i15_i_ANNOT_STR]]
102+
; CHECK-LLVM: [[PTR_RESULT_LOAD_1:[%0-9a-z.]+]] = load i32, ptr addrspace(4) [[PTR_ANNOT_CALL]]
103103
%10 = addrspacecast i32 addrspace(1)* %add.ptr.i to i32 addrspace(4)*
104104
%11 = tail call i32 addrspace(4)* @llvm.ptr.annotation.p4i32(i32 addrspace(4)* %10, i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str.3, i64 0, i64 0), i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str.1, i64 0, i64 0), i32 0, i8* null) #2
105105
store i32 %6, i32 addrspace(4)* %11, align 4, !tbaa !9
106-
; CHECK-LLVM: [[PTR_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast i32 addrspace(1)* [[PTR_i]] to i32 addrspace(4)*
107-
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call i32 addrspace(4)* @llvm.ptr.annotation.p4i32.p0i8(i32 addrspace(4)* [[PTR_i_AS_CAST]], i8* getelementptr inbounds ({{.*}} [[PTR_i_ANNOT_STR]]
108-
; CHECK-LLVM: store i32 [[PTR_RESULT_LOAD]], i32 addrspace(4)* [[PTR_ANNOT_CALL]]
106+
; CHECK-LLVM: [[PTR_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast ptr addrspace(1) [[PTR_i]] to ptr addrspace(4)
107+
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p0(ptr addrspace(4) [[PTR_i_AS_CAST]], ptr [[PTR_i_ANNOT_STR]]
108+
; CHECK-LLVM: store i32 [[PTR_RESULT_LOAD]], ptr addrspace(4) [[PTR_ANNOT_CALL]]
109109
%add.ptr.i.i = getelementptr inbounds i32, i32 addrspace(1)* %add.ptr.i, i64 1
110110
%12 = addrspacecast i32 addrspace(1)* %add.ptr.i.i to i32 addrspace(4)*
111111
%13 = tail call i32 addrspace(4)* @llvm.ptr.annotation.p4i32(i32 addrspace(4)* %12, i8* getelementptr inbounds ([25 x i8], [25 x i8]* @.str.4, i64 0, i64 0), i8* getelementptr inbounds ([14 x i8], [14 x i8]* @.str.1, i64 0, i64 0), i32 0, i8* null) #2
112112
store i32 %9, i32 addrspace(4)* %13, align 4, !tbaa !9
113-
; CHECK-LLVM: [[PTR_i_i:[%0-9a-z.]+]] = getelementptr inbounds i32, i32 addrspace(1)* {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
114-
; CHECK-LLVM: [[PTR_i_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast i32 addrspace(1)* [[PTR_i_i]] to i32 addrspace(4)*
115-
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call i32 addrspace(4)* @llvm.ptr.annotation.p4i32.p0i8(i32 addrspace(4)* [[PTR_i_i_AS_CAST]], i8* getelementptr inbounds ({{.*}} [[PTR_i_i_ANNOT_STR]]
116-
; CHECK-LLVM: store i32 [[PTR_RESULT_LOAD_1]], i32 addrspace(4)* [[PTR_ANNOT_CALL]]
113+
; CHECK-LLVM: [[PTR_i_i:[%0-9a-z.]+]] = getelementptr inbounds i32, ptr addrspace(1) {{[%0-9a-z._]+}}, i64 {{[%0-9a-z.]+}}
114+
; CHECK-LLVM: [[PTR_i_i_AS_CAST:[%0-9a-z.]+]] = addrspacecast ptr addrspace(1) [[PTR_i_i]] to ptr addrspace(4)
115+
; CHECK-LLVM: [[PTR_ANNOT_CALL:[%0-9a-z.]+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p0(ptr addrspace(4) [[PTR_i_i_AS_CAST]], ptr [[PTR_i_i_ANNOT_STR]]
116+
; CHECK-LLVM: store i32 [[PTR_RESULT_LOAD_1]], ptr addrspace(4) [[PTR_ANNOT_CALL]]
117117
ret void
118118
}
119119

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributes.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -205,8 +205,8 @@
205205
; LLVM IR compilation command:
206206
; clang -cc1 -triple spir -disable-llvm-passes -fsycl-is-device -emit-llvm intel-fpga-local-var.cpp
207207

208-
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
209-
; RUN: llvm-spirv %t.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_fpga_memory_attributes -o %t.spv
208+
; RUN: llvm-as %s -o %t.bc
209+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_memory_attributes -o %t.spv
210210
; RUN: llvm-spirv %t.spv --spirv-ext=+SPV_INTEL_fpga_memory_attributes -to-text -o %t.spt
211211
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
212212

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_memory_attributes/IntelFPGAMemoryAttributesForStruct.ll

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -181,15 +181,15 @@
181181
; LLVM IR compilation command:
182182
; clang -cc1 -triple spir -disable-llvm-passes -fsycl-is-device -emit-llvm intel-fpga-local-var.cpp
183183

184-
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
185-
; RUN: llvm-spirv %t.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_fpga_memory_attributes -o %t.spv
184+
; RUN: llvm-as %s -o %t.bc
185+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_memory_attributes -o %t.spv
186186
; RUN: llvm-spirv %t.spv --spirv-ext=+SPV_INTEL_fpga_memory_attributes -to-text -o %t.spt
187187
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
188188

189-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
189+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
190190
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
191191

192-
; RUN: llvm-spirv -spirv-text -r %t.spt -o %t.rev.bc
192+
; RUN: llvm-spirv -spirv-text -r -emit-opaque-pointers=0 %t.spt -o %t.rev.bc
193193
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
194194

195195
; TODO: add a bunch of different tests for --spirv-ext option

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_fpga_reg/IntelFPGAReg.ll

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -54,13 +54,13 @@
5454
; A ca(213);
5555
; A cb = __builtin_intel_fpga_reg(ca);
5656

57-
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
57+
; RUN: llvm-as %s -o %t.bc
5858
; FIXME: add more negative test cases
59-
; RUN: llvm-spirv %t.bc -opaque-pointers=0 --spirv-ext=+SPV_INTEL_fpga_reg -o %t.spv
59+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_reg -o %t.spv
6060
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
6161
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
6262

63-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
63+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
6464
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
6565

6666
; CHECK-SPIRV: Capability FPGARegINTEL

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_io_pipes/PipeStorageIOINTEL.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3,10 +3,10 @@
33
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
44
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
55

6-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
6+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
77
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

9-
; RUN: llvm-spirv -spirv-text -r %t.spt -o %t.rev.bc
9+
; RUN: llvm-spirv -spirv-text -r -emit-opaque-pointers=0 %t.spt -o %t.rev.bc
1010
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
1111

1212
; CHECK-LLVM-DAG: %spirv.ConstantPipeStorage = type { i32, i32, i32 }

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
; RUN: llvm-spirv %t.bc -opaque-pointers=0 -spirv-ext=+SPV_INTEL_joint_matrix -o %t.spv
33
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44

5-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
5+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
66
; RUN: llvm-dis -opaque-pointers=0 %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
77

88
; CHECK-SPIRV: Capability JointMatrixINTEL

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_bfloat16.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
99
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
1010

11-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
11+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
1212
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
1313

1414
; CHECK-REGULARIZED: %[[Alloca:.*]] = alloca %"class.cl::sycl::ext::intel::experimental::bfloat16", align 2

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_half.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
; RUN: llvm-spirv %t.bc -opaque-pointers=0 -spirv-ext=+SPV_INTEL_joint_matrix -o %t.spv
33
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44

5-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
5+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
66
; RUN: llvm-dis -opaque-pointers=0 %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
77

88
; CHECK-SPIRV-DAG: TypeFloat [[#FloatTy:]] 32

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/joint_matrix_tf32.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
55
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
66

7-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
7+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
88
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
99

1010
; CHECK-SPIRV-DAG: Capability TensorFloat32RoundingINTEL

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/opaque_joint_matrix.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
; RUN: llvm-spirv %t.spv -to-text -o %t.spt
44
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
55

6-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc -opaque-pointers=0
6+
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc -emit-opaque-pointers=0
77
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

99
; CHECK-SPIRV-DAG: Capability JointMatrixINTEL

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_masked_gather_scatter/intel-basic-vector-pointers.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
; RUN: llvm-spirv %t.spv --to-text -o %t.spt
44
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
55

6-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
6+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
77
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

99
; RUN: not llvm-spirv %t.bc -opaque-pointers=0 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_masked_gather_scatter/intel-gather-scatter.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,7 @@
33
; RUN: llvm-spirv %t.spv --to-text -o %t.spt
44
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
55

6-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
6+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
77
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
88

99
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_masked_gather_scatter -o %t.spv -spirv-allow-unknown-intrinsics

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_media_block_io/SPV_INTEL_media_block_io.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,9 +2,9 @@
22
// RUN: llvm-spirv --spirv-ext=+SPV_INTEL_media_block_io %t.bc -o %t.spv
33
// RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
44
// RUN: spirv-val %t.spv
5-
// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
5+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.rev.bc
66
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
7-
// RUN: llvm-spirv -r --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
7+
// RUN: llvm-spirv -r -emit-opaque-pointers=0 --spirv-target-env=SPV-IR %t.spv -o %t.rev.bc
88
// RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-SPV-IR
99

1010
uchar __attribute__((overloadable)) intel_sub_group_media_block_read_uc(int2 src_offset, int width, int height, read_only image2d_t image);

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_subgroups/cl_intel_sub_groups.ll

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -36,9 +36,9 @@
3636
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
3737
; RUN: llvm-spirv %t.bc -o - -spirv-text --spirv-ext=+SPV_INTEL_subgroups | FileCheck %s --check-prefix=CHECK-SPIRV
3838
; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_INTEL_subgroups
39-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
39+
; RUN: llvm-spirv -r %t.spv -emit-opaque-pointers=0 -o %t.rev.bc
4040
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
41-
; RUN: llvm-spirv -r %t.spv -o %t.rev.bc --spirv-target-env=SPV-IR
41+
; RUN: llvm-spirv -r %t.spv -emit-opaque-pointers=0 -o %t.rev.bc --spirv-target-env=SPV-IR
4242
; RUN: llvm-dis -opaque-pointers=0 < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-SPIRV
4343
; RUN: llvm-spirv %t.rev.bc -o - -spirv-text --spirv-ext=+SPV_INTEL_subgroups | FileCheck %s --check-prefix=CHECK-SPIRV
4444

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_vector_compute/buffer_surface_intel.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
22
; RUN: llvm-spirv %t.bc -opaque-pointers=0 -o %t.spv --spirv-ext=+SPV_INTEL_vector_compute --spirv-allow-unknown-intrinsics
33
; RUN: llvm-spirv %t.spv -o %t.spt --to-text
4-
; RUN: llvm-spirv -r %t.spv -o %t.bc
4+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc
55
; RUN: llvm-dis -opaque-pointers=0 %t.bc -o %t.ll
66
; RUN: FileCheck %s --input-file %t.spt -check-prefix=SPV
77
; RUN: FileCheck %s --input-file %t.ll -check-prefix=LLVM

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_vector_compute/decoration_media_block_io.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
; RUN: llvm-as -opaque-pointers=0 %s -o %t.bc
22
; RUN: llvm-spirv %t.bc -opaque-pointers=0 -o %t.spv --spirv-ext=+SPV_INTEL_vector_compute --spirv-allow-unknown-intrinsics=llvm.genx
33
; RUN: llvm-spirv %t.spv -o %t.spt --to-text
4-
; RUN: llvm-spirv -r %t.spv -o %t.bc
4+
; RUN: llvm-spirv -r -emit-opaque-pointers=0 %t.spv -o %t.bc
55
; RUN: llvm-dis -opaque-pointers=0 %t.bc -o %t.ll
66
; RUN: FileCheck %s --input-file %t.spt -check-prefix=SPV
77
; RUN: FileCheck %s --input-file %t.ll -check-prefix=LLVM

0 commit comments

Comments
 (0)