Skip to content

Commit c4c9bdb

Browse files
maksimsabmaksimsab
committed
[SYCL] Add generation of device image with specialization constants replaced by default values
1 parent 5d92897 commit c4c9bdb

File tree

9 files changed

+383
-0
lines changed

9 files changed

+383
-0
lines changed

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/device-image.ll

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,44 @@
1+
; Test checks the content of simple generated device image.
2+
3+
; RUN: sycl-post-link -split=auto -spec-const=rt -symbols -S -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE -DPATH=%t
5+
; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
6+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
7+
; RUN: cat %t_0.ll | FileCheck %s -check-prefix=CHECK-IR0
8+
; RUN: cat %t_1.ll | FileCheck %s -check-prefix=CHECK-IR1
9+
10+
; CHECK-TABLE: [[PATH]]_0.ll|[[PATH]]_0.prop|[[PATH]]_0.sym
11+
; CHECK-TABLE: [[PATH]]_1.ll|[[PATH]]_1.prop|[[PATH]]_1.sym
12+
13+
; CHECK-PROP0-NOT: defaultSpecConstants=1|1
14+
; CHECK-PROP0-NOT: originalImage
15+
16+
; CHECK-PROP1: defaultSpecConstants=1|1
17+
18+
; CHECK-IR0: call i32 @_Z20__spirv_SpecConstantii
19+
; CHECK-IR0: call %struct.B @_Z29__spirv_SpecConstantCompositeiii_Rstruct.B
20+
; CHECK-IR0: call %struct.A @_Z29__spirv_SpecConstantCompositeistruct.B_Rstruct.A
21+
22+
; CHECK-IR1-NOT: SpecConstant
23+
; CHECK-IR1: call void @llvm.memcpy.p4i8.p4i8.i64(i8 addrspace(4)* align 8 %1, i8 addrspace(4)* align 8 addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i64 16, i1 false)
24+
25+
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"
26+
target triple = "spir64-unknown-unknown"
27+
28+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
29+
%struct.A = type { i32, %struct.B }
30+
%struct.B = type { i32, i32, i32 }
31+
32+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
33+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
34+
35+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
36+
37+
define spir_kernel void @func1() {
38+
entry:
39+
%a.i = alloca %struct.A, align 4
40+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
41+
%0 = bitcast %struct.A* %a.i to i8*
42+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
43+
ret void
44+
}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/esimd.ll

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
; Test checks generation of device image of esimd kernel.
2+
3+
; RUN: sycl-post-link -split=auto -emit-param-info -symbols -emit-exported-symbols -split-esimd -lower-esimd -O2 -spec-const=rt -device-globals -o %t.table %s -generate-device-image-default-spec-consts && \
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE -DPATH=%t && \
5+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP && \
6+
; RUN: cat %t_esimd_1.prop | FileCheck %s -check-prefix=CHECK-ESIMD-PROP
7+
8+
; CHECK-TABLE: [[PATH]]_esimd_0.bc|[[PATH]]_esimd_0.prop|[[PATH]]_esimd_0.sym
9+
; CHECK-TABLE: [[PATH]]_0.bc|[[PATH]]_0.prop|[[PATH]]_0.sym
10+
; CHECK-TABLE: [[PATH]]_esimd_1.bc|[[PATH]]_esimd_1.prop|[[PATH]]_esimd_1.sym
11+
; CHECK-TABLE: [[PATH]]_1.bc|[[PATH]]_1.prop|[[PATH]]_1.sym
12+
13+
; CHECK-PROP: defaultSpecConstants=1|1
14+
15+
; CHECK-ESIMD-PROP: isEsimdImage=1|1
16+
; CHECK-ESIMD-PROP: defaultSpecConstants=1|1
17+
18+
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"
19+
target triple = "spir64-unknown-unknown"
20+
21+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
22+
%struct.A = type { i32, %struct.B }
23+
%struct.B = type { i32, i32, i32 }
24+
25+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
26+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
27+
28+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
29+
30+
define spir_kernel void @func1() !kernel_arg_buffer_location !7 !sycl_kernel_omit_args !8 {
31+
entry:
32+
%a.i = alloca %struct.A, align 4
33+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
34+
%0 = bitcast %struct.A* %a.i to i8*
35+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
36+
ret void
37+
}
38+
39+
define spir_kernel void @func2(i8 addrspace(1)* noundef align 1 %_arg__specialization_constants_buffer) !sycl_explicit_simd !1 !kernel_arg_addr_space !2 !kernel_arg_access_qual !3 !kernel_arg_type !4 !kernel_arg_base_type !4 !kernel_arg_type_qual !5 !kernel_arg_accessor_ptr !6 {
40+
entry:
41+
%a.i = alloca %struct.A, align 4
42+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
43+
%0 = bitcast %struct.A* %a.i to i8*
44+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
45+
ret void
46+
}
47+
48+
!1 = !{}
49+
!2 = !{i32 1}
50+
!3 = !{!"none"}
51+
!4 = !{!"char*"}
52+
!5 = !{!""}
53+
!6 = !{i1 false}
54+
!7 = !{i32 -1}
55+
!8 = !{i1 true}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-kernel.ll

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,53 @@
1+
; Test checks generation of device images for splitted kernels.
2+
3+
; RUN: sycl-post-link -split=kernel -symbols -o %t.table %s -generate-device-image-default-spec-consts
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE -DPATH=%t
5+
; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
6+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
7+
; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
8+
; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
9+
10+
; CHECK-TABLE: [[PATH]]_0.bc|[[PATH]]_0.prop|[[PATH]]_0.sym
11+
; CHECK-TABLE: [[PATH]]_1.bc|[[PATH]]_1.prop|[[PATH]]_1.sym
12+
; CHECK-TABLE: [[PATH]]_2.bc|[[PATH]]_2.prop|[[PATH]]_2.sym
13+
; CHECK-TABLE: [[PATH]]_3.bc|[[PATH]]_3.prop|[[PATH]]_3.sym
14+
15+
; CHECK-PROP0-NOT: defaultSpecConstants=1|1
16+
; CHECK-PROP0-NOT: originalImage
17+
18+
; CHECK-PROP1: defaultSpecConstants=1|1
19+
20+
; CHECK-PROP2-NOT: defaultSpecConstants=1|1
21+
; CHECK-PROP2-NOT: originalImage
22+
23+
; CHECK-PROP3: defaultSpecConstants=1|1
24+
25+
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"
26+
target triple = "spir64-unknown-unknown"
27+
28+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
29+
%struct.A = type { i32, %struct.B }
30+
%struct.B = type { i32, i32, i32 }
31+
32+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
33+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
34+
35+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
36+
37+
define spir_kernel void @kernel1() {
38+
entry:
39+
%a.i = alloca %struct.A, align 4
40+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
41+
%0 = bitcast %struct.A* %a.i to i8*
42+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
43+
ret void
44+
}
45+
46+
define spir_kernel void @kernel2() {
47+
entry:
48+
%a.i = alloca %struct.A, align 4
49+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
50+
%0 = bitcast %struct.A* %a.i to i8*
51+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
52+
ret void
53+
}

‎llvm/test/tools/sycl-post-link/spec-constants/default-value/split-by-source.ll

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
; Test checks generation of device images for splitted kernels by source.
2+
3+
; RUN: sycl-post-link -split=source -symbols -o %t.table %s -generate-device-image-default-spec-consts && \
4+
; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE -DPATH=%t
5+
; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
6+
; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
7+
; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
8+
; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
9+
10+
; CHECK-TABLE: [[PATH]]_0.bc|[[PATH]]_0.prop|[[PATH]]_0.sym
11+
; CHECK-TABLE: [[PATH]]_1.bc|[[PATH]]_1.prop|[[PATH]]_1.sym
12+
; CHECK-TABLE: [[PATH]]_2.bc|[[PATH]]_2.prop|[[PATH]]_2.sym
13+
; CHECK-TABLE: [[PATH]]_3.bc|[[PATH]]_3.prop|[[PATH]]_3.sym
14+
15+
; CHECK-PROP0-NOT: defaultSpecConstants=1|1
16+
; CHECK-PROP0-NOT: originalImage
17+
18+
; CHECK-PROP1: defaultSpecConstants=1|1
19+
20+
; CHECK-PROP2-NOT: defaultSpecConstants=1|1
21+
; CHECK-PROP2-NOT: originalImage
22+
23+
; CHECK-PROP3: defaultSpecConstants=1|1
24+
25+
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"
26+
target triple = "spir64-unknown-unknown"
27+
28+
%"class.sycl::_V1::specialization_id" = type { %struct.A }
29+
%struct.A = type { i32, %struct.B }
30+
%struct.B = type { i32, i32, i32 }
31+
32+
@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
33+
@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
34+
35+
declare spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
36+
37+
define spir_kernel void @kernel1() #0 {
38+
entry:
39+
%a.i = alloca %struct.A, align 4
40+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
41+
%0 = bitcast %struct.A* %a.i to i8*
42+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
43+
ret void
44+
}
45+
46+
define spir_kernel void @kernel2() #1 {
47+
entry:
48+
%a.i = alloca %struct.A, align 4
49+
%a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
50+
%0 = bitcast %struct.A* %a.i to i8*
51+
call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI1AET_PKcPKvS5_(%struct.A addrspace(4)* sret(%struct.A) align 4 %a.ascast.i, i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
52+
ret void
53+
}
54+
55+
attributes #0 = { "sycl-module-id"="TU1.cpp" }
56+
attributes #1 = { "sycl-module-id"="TU2.cpp" }

‎llvm/tools/sycl-post-link/ModuleSplitter.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -562,6 +562,28 @@ void ModuleDesc::cleanup() {
562562
MPM.run(*M, MAM);
563563
}
564564

565+
bool ModuleDesc::isSpecConstantDefault() const {
566+
return Props.IsSpecConstantDefault;
567+
}
568+
569+
void ModuleDesc::setSpecConstantDefault(bool Value) {
570+
Props.IsSpecConstantDefault = Value;
571+
}
572+
573+
void ModuleDesc::setOriginalImageName(const std::string &Name) {
574+
OriginalImageName = Name;
575+
}
576+
577+
const std::optional<std::string> &ModuleDesc::tryGetOriginalImageName() const {
578+
return OriginalImageName;
579+
}
580+
581+
ModuleDesc CreateModuleDescWithDefaultSpecConstants(std::unique_ptr<Module> M) {
582+
ModuleDesc NewMD(std::move(M));
583+
NewMD.setSpecConstantDefault(true);
584+
return NewMD;
585+
}
586+
565587
#ifndef NDEBUG
566588
void ModuleDesc::verifyESIMDProperty() const {
567589
if (EntryPoints.Props.HasESIMD == SyclEsimdSplitStatus::SYCL_AND_ESIMD) {
@@ -629,6 +651,16 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
629651
});
630652
}
631653

654+
void EntryPointGroup::rebuild(const Module &M) {
655+
std::vector<std::string> Names;
656+
for (const Function &F : M.functions()) {
657+
if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
658+
continue;
659+
660+
Functions.insert(const_cast<Function *>(&F));
661+
}
662+
}
663+
632664
namespace {
633665
// This is a helper class, which allows to group/categorize function based on
634666
// provided rules. It is intended to be used in device code split

‎llvm/tools/sycl-post-link/ModuleSplitter.h

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -107,8 +107,12 @@ class ModuleDesc {
107107
EntryPointGroup EntryPoints;
108108
bool IsTopLevel = false;
109109

110+
// Name of the original device image from which this MD was generated from.
111+
std::optional<std::string> OriginalImageName;
112+
110113
public:
111114
struct Properties {
115+
bool IsSpecConstantDefault = false;
112116
bool SpecConstsMet = false;
113117
};
114118
std::string Name = "";
@@ -183,12 +187,20 @@ class ModuleDesc {
183187
// Cleans up module IR - removes dead globals, debug info etc.
184188
void cleanup();
185189

190+
bool isSpecConstantDefault() const;
191+
void setSpecConstantDefault(bool Value);
192+
193+
void setOriginalImageName(const std::string &Name);
194+
const std::optional<std::string> &tryGetOriginalImageName() const;
195+
186196
#ifndef NDEBUG
187197
void verifyESIMDProperty() const;
188198
void dump() const;
189199
#endif // NDEBUG
190200
};
191201

202+
ModuleDesc CreateModuleDescWithDefaultSpecConstants(std::unique_ptr<Module> M);
203+
192204
// Module split support interface.
193205
// It gets a module (in a form of module descriptor, to get additional info) and
194206
// a collection of entry points groups. Each group specifies subset entry points

‎llvm/tools/sycl-post-link/SpecConstants.cpp

Lines changed: 66 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,10 +14,15 @@
1414
#include "llvm/ADT/APInt.h"
1515
#include "llvm/ADT/StringMap.h"
1616
#include "llvm/ADT/StringRef.h"
17+
#include "llvm/IR/IRBuilder.h"
1718
#include "llvm/IR/InstIterator.h"
1819
#include "llvm/IR/Instruction.h"
1920
#include "llvm/IR/Instructions.h"
2021
#include "llvm/IR/Operator.h"
22+
#include "llvm/Transforms/Utils/Cloning.h"
23+
24+
#include <memory>
25+
#include <vector>
2126

2227
using namespace llvm;
2328

@@ -914,3 +919,64 @@ bool SpecConstantsPass::collectSpecConstantDefaultValuesMetadata(
914919

915920
return true;
916921
}
922+
923+
void replaceSpecConstsWithDefaultValue(const Module &M,
924+
const std::vector<CallInst *> &CIs) {
925+
const DataLayout &DL = M.getDataLayout();
926+
for (CallInst *CI : CIs) {
927+
Value *Dst = CI->getArgOperand(0);
928+
Value *Src = CI->getArgOperand(2);
929+
Align DstAlign = DL.getABITypeAlign(Dst->getType());
930+
Align SrcAlign = DL.getABITypeAlign(Src->getType());
931+
Type *StructureType =
932+
Src->stripPointerCasts()->getType()->getContainedType(0);
933+
uint64_t SpecConstSize = DL.getTypeAllocSize(StructureType);
934+
935+
IRBuilder B(CI);
936+
B.CreateMemCpy(Dst, DstAlign, Src, SrcAlign, SpecConstSize);
937+
CI->removeFromParent();
938+
CI->deleteValue();
939+
}
940+
}
941+
942+
static bool checkModuleContainsSpecConsts(const Module &M) {
943+
for (const Function &F : M.functions()) {
944+
if (F.getName().startswith(SYCL_GET_SCALAR_2020_SPEC_CONST_VAL) ||
945+
F.getName().startswith(SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL))
946+
return true;
947+
}
948+
949+
return false;
950+
}
951+
952+
/// Function generates a copy of Module and replaces all spec constant uses
953+
/// memcpy intrinsic.
954+
/// If the module doesn't contain spec constants then nullptr is returned.
955+
std::unique_ptr<Module>
956+
llvm::generateDeviceImageWithDefaultSpecConstants(const Module &M) {
957+
if (!checkModuleContainsSpecConsts(M))
958+
return nullptr;
959+
960+
std::unique_ptr<Module> M2 = CloneModule(M);
961+
std::vector<Function *> SpecConstDeclarations;
962+
std::vector<CallInst *> CIs;
963+
for (Function &F : M2->functions()) {
964+
if (!F.isDeclaration())
965+
continue;
966+
967+
if (!F.getName().startswith(SYCL_GET_SCALAR_2020_SPEC_CONST_VAL) &&
968+
!F.getName().startswith(SYCL_GET_COMPOSITE_2020_SPEC_CONST_VAL))
969+
continue;
970+
971+
SpecConstDeclarations.push_back(&F);
972+
for (User *U : F.users())
973+
if (auto *CI = dyn_cast<CallInst>(U))
974+
CIs.push_back(CI);
975+
}
976+
977+
replaceSpecConstsWithDefaultValue(*M2, CIs);
978+
for (Function *F : SpecConstDeclarations)
979+
F->removeFromParent();
980+
981+
return std::move(M2);
982+
}

‎llvm/tools/sycl-post-link/SpecConstants.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,7 @@
1919
#include "llvm/IR/Module.h"
2020
#include "llvm/IR/PassManager.h"
2121

22+
#include <memory>
2223
#include <vector>
2324

2425
namespace llvm {
@@ -70,4 +71,7 @@ class SpecConstantsPass : public PassInfoMixin<SpecConstantsPass> {
7071
bool SetValAtRT;
7172
};
7273

74+
std::unique_ptr<Module>
75+
generateDeviceImageWithDefaultSpecConstants(const Module &M);
76+
7377
} // namespace llvm

0 commit comments

Comments
 (0)