[SYCL] Add generation of device image with specialization constants replaced by default values

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

@@ -0,0 +1,26 @@
+; Test checks handling of bool specialization constant.
+
+; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: FileCheck %s -input-file %t_1.ll --implicit-check-not="SpecConst"
+
+; CHECK: %bool1 = trunc i8 1 to i1
+; CHECK: %frombool = zext i1 %bool1 to i8
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::specialization_id" = type { i8 }
+%struct.A = type { i8 }
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { i8 1 }, align 4
+
+; Function Attrs: convergent nounwind
+declare dso_local spir_func noundef zeroext i1 @_Z37__sycl_getScalar2020SpecConstantValueIbET_PKcPKvS4_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
+
+define spir_kernel void @kernel() {
+entry:
+  %bool = call spir_func noundef zeroext i1 @_Z37__sycl_getScalar2020SpecConstantValueIbET_PKcPKvS4_(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)* getelementptr inbounds (%"class.sycl::_V1::specialization_id", %"class.sycl::_V1::specialization_id" addrspace(1)* @_ZL1c, i64 0, i32 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef null) #4
+  %frombool = zext i1 %bool to i8
+  ret void
+}

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

@@ -0,0 +1,90 @@
+; Test checks the content of simple generated device image.
+; It checks scalar, sret and "return by value" versions of SpecConstant functions.
+; Also test checks generated symbols.
+
+; RUN: sycl-post-link -split=auto -spec-const=native -symbols -S -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: FileCheck %s -input-file %t.table -check-prefix=CHECK-TABLE
+; RUN: FileCheck %s -input-file %t_0.prop -check-prefix=CHECK-PROP0
+; RUN: FileCheck %s -input-file %t_1.prop -check-prefix=CHECK-PROP1
+; RUN: FileCheck %s -input-file %t_0.ll -check-prefix=CHECK-IR0
+; RUN: FileCheck %s -input-file %t_1.ll -check-prefix=CHECK-IR1 --implicit-check-not "SpecConstant"
+; RUN: FileCheck %s -input-file %t_0.sym -check-prefix=CHECK-SYM0
+; RUN: FileCheck %s -input-file %t_1.sym -check-prefix=CHECK-SYM1
+
+; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop|{{.*}}_0.sym
+; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop|{{.*}}_1.sym
+
+; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP1: specConstsReplacedWithDefault=1|1
+
+; CHECK-IR0: call i32 @_Z20__spirv_SpecConstantii
+; CHECK-IR0: call %struct.B @_Z29__spirv_SpecConstantCompositeiii_Rstruct.B
+; CHECK-IR0: call %struct.A @_Z29__spirv_SpecConstantCompositeistruct.B_Rstruct.A
+
+; CHECK-IR1: store %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } }, %struct.A addrspace(4)* %a.ascast.i, align 4
+
+; Check that %scalar value has been replaced by global value.
+; CHECK-IR1-NOT: %scalar = call
+; CHECK-IR1: %scalar2 = add i32 123, 1
+
+; Check that %returned_spec_const value has been replaced by global value.
+; CHECK-IR1-NOT: %returned_spec_const = call
+; CHECK-IR1: %sc.e = extractvalue %struct.C { i32 1 }, 0
+
+; CHECK-SYM0: kernel
+; CHECK-SYM0-EMPTY:
+
+; CHECK-SYM1: kernel
+; CHECK-SYM1-EMPTY:
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::specialization_id" = type { %struct.A }
+%"class.sycl::_V1::specialization_id.2" = type { i32 }
+%"class.sycl::_V1::specialization_id.3" = type { %struct.C }
+%struct.A = type { i32, %struct.B }
+%struct.B = type { i32, i32, i32 }
+%struct.C = type { i32 }
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@__usid_str.1 = private unnamed_addr constant [33 x i8] c"uidcac21ed8fab7d507____ZL6valueS\00", align 1
+@__usid_str.2 = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8f____ZL1b\00", align 1
+@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
+@_ZL6valueS = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.2" { i32 123 }, align 4
+@_ZL1b = internal addrspace(1) constant %"class.sycl::_V1::specialization_id.3" { %struct.C { i32 1 } }, align 4
+
+; This constant checks `zeroinitializer` field.
+@_ZL1d = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B zeroinitializer } }, align 4
+
+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)
+
+declare spir_func %struct.C @_Z40__sycl_getComposite2020SpecConstantValueI1CET_PKcPKvS5_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
+
+declare dso_local spir_func noundef i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* noundef, i8 addrspace(4)* noundef, i8 addrspace(4)* noundef)
+
+; Function for testing symbol generation
+define spir_func void @func() {
+  ret void
+}
+
+define spir_kernel void @kernel() {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  %scalar = call spir_func i32 @_Z37__sycl_getScalar2020SpecConstantValueIiET_PKcPKvS4_(i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([33 x i8], [33 x i8]* @__usid_str.1, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id.2" addrspace(1)* @_ZL6valueS to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
+  %scalar2 = add i32 %scalar, 1
+
+  %returned_spec_const = call spir_func %struct.C @_Z40__sycl_getComposite2020SpecConstantValueI1CET_PKcPKvS5_(i8 addrspace(4)* noundef addrspacecast (i8* getelementptr inbounds ([28 x i8], [28 x i8]* @__usid_str.2, i64 0, i64 0) to i8 addrspace(4)*), i8 addrspace(4)* noundef addrspacecast (i8 addrspace(1)* bitcast (%"class.sycl::_V1::specialization_id.3" addrspace(1)* @_ZL1b to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
+  %sc.e = extractvalue %struct.C %returned_spec_const, 0
+  %scalar3 = add i32 %sc.e, 1
+
+  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)* @_ZL1d to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
+
+
+  call void @func()
+  ret void
+}

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

@@ -0,0 +1,55 @@
+; Test checks generation of device image of esimd kernel.
+
+; RUN: sycl-post-link -split=auto -split-esimd -lower-esimd -O2 -spec-const=native -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: FileCheck %s -input-file=%t.table -check-prefix=CHECK-TABLE
+; RUN: FileCheck %s -input-file=%t_1.prop -check-prefix=CHECK-PROP
+; RUN: FileCheck %s -input-file=%t_esimd_1.prop -check-prefix=CHECK-ESIMD-PROP
+
+; CHECK-TABLE: {{.*}}_esimd_0.bc|{{.*}}_esimd_0.prop
+; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
+; CHECK-TABLE: {{.*}}_esimd_1.bc|{{.*}}_esimd_1.prop
+; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
+
+; CHECK-PROP: specConstsReplacedWithDefault=1|1
+
+; CHECK-ESIMD-PROP: isEsimdImage=1|1
+; CHECK-ESIMD-PROP: specConstsReplacedWithDefault=1|1
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::specialization_id" = type { %struct.A }
+%struct.A = type { i32, %struct.B }
+%struct.B = type { i32, i32, i32 }
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
+
+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)
+
+define spir_kernel void @func1() !kernel_arg_buffer_location !7 !sycl_kernel_omit_args !8 {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}
+
+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 {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}
+
+!1 = !{}
+!2 = !{i32 1}
+!3 = !{!"none"}
+!4 = !{!"char*"}
+!5 = !{!""}
+!6 = !{i1 false}
+!7 = !{i32 -1}
+!8 = !{i1 true}

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

@@ -0,0 +1,51 @@
+; Test checks generation of device images for splitted kernels.
+
+; RUN: sycl-post-link -split=kernel -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
+; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
+; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
+; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
+; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
+
+; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
+; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
+; CHECK-TABLE: {{.*}}_2.bc|{{.*}}_2.prop
+; CHECK-TABLE: {{.*}}_3.bc|{{.*}}_3.prop
+
+; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP1: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP2-NOT: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP3: specConstsReplacedWithDefault=1|1
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::specialization_id" = type { %struct.A }
+%struct.A = type { i32, %struct.B }
+%struct.B = type { i32, i32, i32 }
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
+
+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)
+
+define spir_kernel void @kernel1() {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}
+
+define spir_kernel void @kernel2() {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}

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

@@ -0,0 +1,54 @@
+; Test checks generation of device images for splitted kernels by source.
+
+; RUN: sycl-post-link -split=source -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
+; RUN: cat %t_0.prop | FileCheck %s -check-prefix=CHECK-PROP0
+; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
+; RUN: cat %t_2.prop | FileCheck %s -check-prefix=CHECK-PROP2
+; RUN: cat %t_3.prop | FileCheck %s -check-prefix=CHECK-PROP3
+
+; CHECK-TABLE: {{.*}}_0.bc|{{.*}}_0.prop
+; CHECK-TABLE: {{.*}}_1.bc|{{.*}}_1.prop
+; CHECK-TABLE: {{.*}}_2.bc|{{.*}}_2.prop
+; CHECK-TABLE: {{.*}}_3.bc|{{.*}}_3.prop
+
+; CHECK-PROP0-NOT: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP1: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP2-NOT: specConstsReplacedWithDefault=1|1
+
+; CHECK-PROP3: specConstsReplacedWithDefault=1|1
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%"class.sycl::_V1::specialization_id" = type { %struct.A }
+%struct.A = type { i32, %struct.B }
+%struct.B = type { i32, i32, i32 }
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@_ZL1c = internal addrspace(1) constant %"class.sycl::_V1::specialization_id" { %struct.A { i32 3, %struct.B { i32 3, i32 2, i32 1 } } }, align 4
+
+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)
+
+define spir_kernel void @kernel1() #0 {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}
+
+define spir_kernel void @kernel2() #1 {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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)
+  ret void
+}
+
+attributes #0 = { "sycl-module-id"="TU1.cpp" }
+attributes #1 = { "sycl-module-id"="TU2.cpp" }

llvm/test/tools/sycl-post-link/spec-constants/default-value/struct-with-padding.ll

@@ -0,0 +1,32 @@
+; Test checks that struct with padding is handled correctly.
+
+; RUN: sycl-post-link -split=auto -spec-const=native -S -o %t.table %s -generate-device-image-default-spec-consts
+; RUN: cat %t.table | FileCheck %s -check-prefix=CHECK-TABLE
+; RUN: cat %t_1.prop | FileCheck %s -check-prefix=CHECK-PROP1
+; RUN: cat %t_1.ll | FileCheck %s -check-prefix=CHECK-IR1 --implicit-check-not SpecConstant
+
+; CHECK-TABLE: {{.*}}_0.ll|{{.*}}_0.prop
+; CHECK-TABLE: {{.*}}_1.ll|{{.*}}_1.prop
+
+; CHECK-PROP1: specConstsReplacedWithDefault=1|1
+
+; CHECK-IR1: store { float, i32, i8 } { float 0x40091EB860000000, i32 42, i8 8 }, { float, i32, i8 } addrspace(4)* %1, align 4
+
+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"
+target triple = "spir64-unknown-unknown"
+
+%struct.A = type <{ float, i32, i8, [3 x i8] }>
+
+@__usid_str = private unnamed_addr constant [28 x i8] c"uida046125e6e1c1f8d____ZL1c\00", align 1
+@_ZL1c = internal addrspace(1)  constant { { float, i32, i8 } } { { float, i32, i8 } { float 0x40091EB860000000, i32 42, i8 8 } }, align 4
+
+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)
+
+define spir_kernel void @func1() {
+entry:
+  %a.i = alloca %struct.A, align 4
+  %a.ascast.i = addrspacecast %struct.A* %a.i to %struct.A addrspace(4)*
+  %0 = bitcast %struct.A* %a.i to i8*
+  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 ({ { float, i32, i8 } } addrspace(1)* @_ZL1c to i8 addrspace(1)*) to i8 addrspace(4)*), i8 addrspace(4)* noundef null)
+  ret void
+}

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

@@ -562,6 +562,21 @@ void ModuleDesc::cleanup() {
   MPM.run(*M, MAM);
 }
 
+bool ModuleDesc::isSpecConstantDefault() const {
+  return Props.IsSpecConstantDefault;
+}
+
+void ModuleDesc::setSpecConstantDefault(bool Value) {
+  Props.IsSpecConstantDefault = Value;
+}
+
+ModuleDesc ModuleDesc::clone() const {
+  std::unique_ptr<Module> NewModule = CloneModule(getModule());
+  ModuleDesc NewMD(std::move(NewModule));
+  NewMD.EntryPoints.Props = EntryPoints.Props;
+  return NewMD;
+}
+
 #ifndef NDEBUG
 void ModuleDesc::verifyESIMDProperty() const {
   if (EntryPoints.Props.HasESIMD == SyclEsimdSplitStatus::SYCL_AND_ESIMD) {
@@ -629,6 +644,12 @@ void EntryPointGroup::rebuildFromNames(const std::vector<std::string> &Names,
   });
 }
 
+void EntryPointGroup::rebuild(const Module &M) {
+  for (const Function &F : M.functions())
+    if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
+      Functions.insert(const_cast<Function *>(&F));
+}
+
 namespace {
 // This is a helper class, which allows to group/categorize function based on
 // provided rules. It is intended to be used in device code split