Skip to content

Commit 2bc58be

Browse files
AlexeySachkovvmaksimo
AlexeySachkov
authored and
vmaksimo
committed
Add possibility to lower BuiltIn-s into SPIR-V friendly IR
1 parent 40efc99 commit 2bc58be

File tree

4 files changed

+59
-35
lines changed

4 files changed

+59
-35
lines changed

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

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -290,7 +290,14 @@ Value *SPIRVToLLVM::mapFunction(SPIRVFunction *BF, Function *F) {
290290
// %d = extractelement <3 x i64> %5, i32 idx
291291
bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV,
292292
SPIRVBuiltinVariableKind Kind) {
293-
std::string FuncName = SPIRSPIRVBuiltinVariableMap::rmap(Kind);
293+
std::string FuncName;
294+
// TODO: we should always produce SPIR-V friendly IR and apply lowering later
295+
// if needed
296+
if (BM->getDesiredBIsRepresentation() != BIsRepresentation::SPIRVFriendlyIR) {
297+
FuncName = SPIRSPIRVBuiltinVariableMap::rmap(Kind);
298+
} else {
299+
FuncName = std::string(GV->getName());
300+
}
294301
Type *ReturnTy = GV->getType()->getPointerElementType();
295302
// Some SPIR-V builtin variables are translated to a function with an index
296303
// argument.

‎llvm-spirv/test/transcoding/builtin_vars.ll

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,9 @@
22
; RUN: llvm-spirv %t.bc -o %t.spv
33
; RUN: spirv-val %t.spv
44
; RUN: llvm-spirv -r %t.spv -o %t.out.bc
5-
; RUN: llvm-dis %t.out.bc -o - | FileCheck %s
5+
; RUN: llvm-dis %t.out.bc -o - | FileCheck %s --check-prefix=CHECK-OCL
6+
; RUN: llvm-spirv -r %t.spv --spirv-target-env=SPV-IR -o %t.out.bc
7+
; RUN: llvm-dis %t.out.bc -o - | FileCheck %s --check-prefix=CHECK-SPV
68

79
target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
810
target triple = "spir-unknown-unknown"
@@ -13,7 +15,8 @@ target triple = "spir-unknown-unknown"
1315
define spir_kernel void @f() #0 !kernel_arg_addr_space !0 !kernel_arg_access_qual !0 !kernel_arg_type !0 !kernel_arg_base_type !0 !kernel_arg_type_qual !0 {
1416
entry:
1517
%0 = load i32, i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @__spirv_BuiltInGlobalLinearId to i32 addrspace(4)*), align 4
16-
; CHECK: %0 = call spir_func i32 @_Z20get_global_linear_idv() #1
18+
; CHECK-OCL: %0 = call spir_func i32 @_Z20get_global_linear_idv() #1
19+
; CHECK-SPV: %0 = call spir_func i32 @_Z29__spirv_BuiltInGlobalLinearIdv() #1
1720
ret void
1821
}
1922

‎llvm-spirv/test/transcoding/builtin_vars_arithmetics.ll

Lines changed: 32 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,9 @@
33
; RUN: spirv-val %t.spv
44
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
55
; RUN: llvm-spirv %t.spv -r -o %t.rev.bc
6-
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
6+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM,CHECK-LLVM-OCL
7+
; RUN: llvm-spirv %t.spv -r --spirv-target-env=SPV-IR -o %t.rev.bc
8+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM,CHECK-LLVM-SPV
79

810
; The IR was generated from the following source:
911
; #include <CL/sycl.hpp>
@@ -39,6 +41,35 @@
3941
; CHECK-SPIRV: Decorate [[GlobalOffset]] LinkageAttributes "__spirv_BuiltInGlobalOffset" Import
4042
; CHECK-SPIRV: Decorate [[GlobalSize]] LinkageAttributes "__spirv_BuiltInGlobalSize" Import
4143
; CHECK-SPIRV: Decorate [[GlobalInvocationId]] LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
44+
;
45+
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x 64> addrspace(4)*
46+
; CHECK-LLVM-NOT: load <3 x i64>
47+
; CHECK-LLVM-OCL: %[[Id0:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 0) #1
48+
; CHECK-LLVM-SPV: %[[Id0:[0-9]+]] = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 0) #1
49+
; CHECK-LLVM: %[[FirstVec:[0-9]+]] = insertelement <3 x i64> undef, i64 %[[Id0]], i32 0
50+
; CHECK-LLVM-OCL: %[[Id1:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 1) #1
51+
; CHECK-LLVM-SPV: %[[Id1:[0-9]+]] = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 1) #1
52+
; CHECK-LLVM: %[[SecondVec:[0-9]+]] = insertelement <3 x i64> %[[FirstVec]], i64 %[[Id1]], i32 1
53+
; CHECK-LLVM-OCL: %[[Id2:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 2) #1
54+
; CHECK-LLVM-SPV: %[[Id2:[0-9]+]] = call spir_func i64 @_Z33__spirv_BuiltInGlobalInvocationIdi(i32 2) #1
55+
; CHECK-LLVM: %[[GlobIdVec:[0-9]+]] = insertelement <3 x i64> %[[SecondVec]], i64 %[[Id2]], i32 2
56+
; CHECK-LLVM: %{{[0-9]+}} = extractelement <3 x i64> %[[GlobIdVec]], i32 1
57+
; CHECK-LLVM: %{{[0-9]+}} = extractelement <3 x i64> %[[GlobIdVec]], i32 0
58+
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalSize to <3 x 64> addrspace(4)*
59+
; CHECK-LLVM-NOT: load <3 x i64>
60+
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalOffset to <3 x 64> addrspace(4)*
61+
; CHECK-LLVM-NOT: load <3 x i64>
62+
; CHECK-LLVM-OCL: %[[GOffset0:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 0) #1
63+
; CHECK-LLVM-SPV: %[[GOffset0:[0-9]+]] = call spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32 0) #1
64+
; CHECK-LLVM: %[[FirstVec2:[0-9]+]] = insertelement <3 x i64> undef, i64 %[[GOffset0]], i32 0
65+
; CHECK-LLVM-OCL: %[[GOffset1:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 1) #1
66+
; CHECK-LLVM-SPV: %[[GOffset1:[0-9]+]] = call spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32 1) #1
67+
; CHECK-LLVM: %[[SecondVec2:[0-9]+]] = insertelement <3 x i64> %[[FirstVec2]], i64 %[[GOffset1]], i32 1
68+
; CHECK-LLVM-OCL: %[[GOffset2:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 2) #1
69+
; CHECK-LLVM-SPV: %[[GOffset2:[0-9]+]] = call spir_func i64 @_Z25__spirv_BuiltInGlobalSizei(i32 2) #1
70+
; CHECK-LLVM: %[[GOffsetVec:[0-9]+]] = insertelement <3 x i64> %[[SecondVec2]], i64 %[[GOffset2]], i32 2
71+
; CHECK-LLVM %20 = sub <3 x i64> %[[GlobSizeVec]], %[[GOffsetVec]]
72+
; CHECK-LLVM %21 = sub <3 x i64> %[[GlobSizeVec]], %[[GOffSetVec]]
4273

4374
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"
4475
target triple = "spir64-unknown-linux-sycldevice"
@@ -62,33 +93,11 @@ entry:
6293
%agg.tmp5.sroa.0.sroa.0.0.copyload = load i64, i64* %agg.tmp5.sroa.0.sroa.0.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx, align 8
6394
%agg.tmp5.sroa.0.sroa.2.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx69 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi2EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 1
6495
%agg.tmp5.sroa.0.sroa.2.0.copyload = load i64, i64* %agg.tmp5.sroa.0.sroa.2.0.agg.tmp5.sroa.0.0..sroa_cast.sroa_idx69, align 8
65-
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x 64> addrspace(4)*
66-
; CHECK-LLVM-NOT: load <3 x i64>
67-
; CHECK-LLVM: %[[Id0:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 0) #1
68-
; CHECK-LLVM: %[[FirstVec:[0-9]+]] = insertelement <3 x i64> undef, i64 %[[Id0]], i32 0
69-
; CHECK-LLVM: %[[Id1:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 1) #1
70-
; CHECK-LLVM: %[[SecondVec:[0-9]+]] = insertelement <3 x i64> %[[FirstVec]], i64 %[[Id1]], i32 1
71-
; CHECK-LLVM: %[[Id2:[0-9]+]] = call spir_func i64 @_Z13get_global_idj(i32 2) #1
72-
; CHECK-LLVM: %[[GlobIdVec:[0-9]+]] = insertelement <3 x i64> %[[SecondVec]], i64 %[[Id2]], i32 2
7396
%0 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
74-
; CHECK-LLVM: %{{[0-9]+}} = extractelement <3 x i64> %[[GlobIdVec]], i32 1
75-
; CHECK-LLVM: %{{[0-9]+}} = extractelement <3 x i64> %[[GlobIdVec]], i32 0
7697
%1 = extractelement <3 x i64> %0, i64 1
7798
%2 = extractelement <3 x i64> %0, i64 0
78-
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalSize to <3 x 64> addrspace(4)*
79-
; CHECK-LLVM-NOT: load <3 x i64>
8099
%3 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalSize to <3 x i64> addrspace(4)*), align 32
81-
; CHECK-LLVM-NOT: addrspacecast <3 x 64> addrspace(1)* @__spirv_BuiltInGlobalOffset to <3 x 64> addrspace(4)*
82-
; CHECK-LLVM-NOT: load <3 x i64>
83-
; CHECK-LLVM: %[[GOffset0:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 0) #1
84-
; CHECK-LLVM: %[[FirstVec2:[0-9]+]] = insertelement <3 x i64> undef, i64 %[[GOffset0]], i32 0
85-
; CHECK-LLVM: %[[GOffset1:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 1) #1
86-
; CHECK-LLVM: %[[SecondVec2:[0-9]+]] = insertelement <3 x i64> %[[FirstVec2]], i64 %[[GOffset1]], i32 1
87-
; CHECK-LLVM: %[[GOffset2:[0-9]+]] = call spir_func i64 @_Z17get_global_offsetj(i32 2) #1
88-
; CHECK-LLVM: %[[GOffsetVec:[0-9]+]] = insertelement <3 x i64> %[[SecondVec2]], i64 %[[GOffset2]], i32 2
89100
%4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalOffset to <3 x i64> addrspace(4)*), align 32
90-
; CHECK-LLVM %20 = sub <3 x i64> %[[GlobSizeVec]], %[[GOffsetVec]]
91-
; CHECK-LLVM %21 = sub <3 x i64> %[[GlobSizeVec]], %[[GOffSetVec]]
92101
%5 = sub <3 x i64> %0, %4
93102
%6 = sub <3 x i64> %0, %4
94103
%7 = extractelement <3 x i64> %6, i64 0

‎llvm-spirv/test/transcoding/builtin_vars_opt.ll

Lines changed: 14 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,9 @@
33
; RUN: spirv-val %t.spv
44
; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
55
; RUN: llvm-spirv %t.spv -r -o %t.rev.bc
6-
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM
6+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM,CHECK-LLVM-OCL
7+
; RUN: llvm-spirv %t.spv -r --spirv-target-env=SPV-IR -o %t.rev.bc
8+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM,CHECK-LLVM-SPV
79

810
; The IR was generated from the following source:
911
; #include <CL/sycl.hpp>
@@ -37,6 +39,17 @@
3739
; CHECK-SPIRV: Decorate [[#SG_MaxSize_BI:]] BuiltIn 37
3840
; CHECK-SPIRV: Decorate [[#SG_MaxSize_BI:]] Constant
3941
; CHECK-SPIRV: Decorate [[#SG_MaxSize_BI:]] LinkageAttributes "__spirv_BuiltInSubgroupMaxSize" Import
42+
;
43+
; CHECK-LLVM-OCL-NOT: @__spirv_BuiltInSubgroupMaxSize
44+
; CHECK-LLVM-NOT: addrspacecast i32 addrspace(1)* @__spirv_BuiltInSubgroupMaxSize to i32 addrspace(4)*
45+
; CHECK-LLVM-LABEL: if.then.i
46+
; CHECK-LLVM-NOT: load
47+
; CHECK-LLVM-OCL: call spir_func i32 @_Z22get_max_sub_group_sizev()
48+
; CHECK-LLVM-SPV: call spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
49+
; CHECK-LLVM-LABEL: cond.false.i:
50+
; CHECK-LLVM-NOT: load
51+
; CHECK-LLVM-OCL: call spir_func i32 @_Z22get_max_sub_group_sizev()
52+
; CHECK-LLVM-SPV: call spir_func i32 @_Z30__spirv_BuiltInSubgroupMaxSizev()
4053

4154
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"
4255
target triple = "spir64-unknown-linux-sycldevice"
@@ -47,7 +60,6 @@ target triple = "spir64-unknown-linux-sycldevice"
4760

4861
$_ZTS10sycl_subgrIiLi0EE = comdat any
4962

50-
; CHECK-LLVM-NOT: @__spirv_BuiltInSubgroupMaxSize
5163
@__spirv_BuiltInSubgroupMaxSize = external dso_local local_unnamed_addr addrspace(1) constant i32, align 4
5264

5365

@@ -59,14 +71,10 @@ entry:
5971
%add.ptr.i = getelementptr inbounds i32, i32 addrspace(1)* %_arg_1, i64 %1
6072
%2 = and i32 %_arg_, 1
6173
%tobool.not.i = icmp eq i32 %2, 0
62-
; CHECK-LLVM-NOT: addrspacecast i32 addrspace(1)* @__spirv_BuiltInSubgroupMaxSize to i32 addrspace(4)*
6374
%3 = addrspacecast i32 addrspace(1)* @__spirv_BuiltInSubgroupMaxSize to i32 addrspace(4)*
6475
br i1 %tobool.not.i, label %if.end.i, label %if.then.i
6576

6677
if.then.i: ; preds = %entry
67-
; CHECK-LLVM: if.then.i
68-
; CHECK-LLVM-NOT: load
69-
; CHECK-LLVM: call spir_func i32 @_Z22get_max_sub_group_sizev()
7078
%4 = load i32, i32 addrspace(4)* %3, align 4, !noalias !8
7179
%ptridx.ascast.i14.i = addrspacecast i32 addrspace(1)* %add.ptr.i to i32 addrspace(4)*
7280
store i32 %4, i32 addrspace(4)* %ptridx.ascast.i14.i, align 4
@@ -78,9 +86,6 @@ if.end.i: ; preds = %if.then.i, %entry
7886
br i1 %tobool4.not.i, label %cond.false.i, label %"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_7nd_itemILi1EEEE_clES5_.exit"
7987

8088
cond.false.i: ; preds = %if.end.i
81-
; CHECK-LLVM: cond.false.i:
82-
; CHECK-LLVM-NOT: load
83-
; CHECK-LLVM: call spir_func i32 @_Z22get_max_sub_group_sizev()
8489
%5 = load i32, i32 addrspace(4)* %3, align 4, !noalias !11
8590
br label %"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_7nd_itemILi1EEEE_clES5_.exit"
8691

0 commit comments

Comments
 (0)