@@ -19,6 +19,7 @@ $_ZTSZ4mainEUlN2cl4sycl14kernel_handlerEE_ = comdat any
19
19
20
20
@__usid_str = private unnamed_addr constant [32 x i8 ] c "ef880fa09cf7a9d7____ZL8coeff_id\00 " , align 1
21
21
@_ZL8coeff_id = internal addrspace (1 ) constant %"class.cl::sycl::specialization_id" { %struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer , i64 0 , [8 x i8 ] undef } }, align 32
22
+ @__usid_str.0 = private unnamed_addr constant [33 x i8 ] c "df991fa0adf9bad8____ZL8coeff_id2\00 " , align 1
22
23
@_ZL8coeff_id2 = internal addrspace (1 ) constant %"class.cl::sycl::specialization_id.1" { %struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer , i64 0 , [7 x i8 ] undef , i8 undef } }, align 32
23
24
24
25
; Function Attrs: convergent norecurse
@@ -38,7 +39,7 @@ define weak_odr dso_local spir_kernel void @_ZTSZ4mainEUlN2cl4sycl14kernel_handl
38
39
%4 = alloca %struct.coeff2_str_aligned_t , align 32
39
40
%5 = addrspacecast %struct.coeff2_str_aligned_t* %4 to %struct.coeff2_str_aligned_t addrspace (4 )*
40
41
%6 = bitcast %struct.coeff2_str_aligned_t* %4 to i8*
41
- call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff2_str_aligned_tET_PKcPKvS5_ (%struct.coeff2_str_aligned_t addrspace (4 )* sret (%struct.coeff2_str_aligned_t ) align 32 %5 , i8 addrspace (4 )* noundef addrspacecast (i8* getelementptr inbounds ([32 x i8 ], [32 x i8 ]* @__usid_str , i64 0 , i64 0 ) to i8 addrspace (4 )*), i8 addrspace (4 )* noundef addrspacecast (i8 addrspace (1 )* bitcast (%"class.cl::sycl::specialization_id.1" addrspace (1 )* @_ZL8coeff_id2 to i8 addrspace (1 )*) to i8 addrspace (4 )*), i8 addrspace (4 )* noundef null ) #4
42
+ call spir_func void @_Z40__sycl_getComposite2020SpecConstantValueI19coeff2_str_aligned_tET_PKcPKvS5_ (%struct.coeff2_str_aligned_t addrspace (4 )* sret (%struct.coeff2_str_aligned_t ) align 32 %5 , i8 addrspace (4 )* noundef addrspacecast (i8* getelementptr inbounds ([33 x i8 ], [33 x i8 ]* @__usid_str.0 , i64 0 , i64 0 ) to i8 addrspace (4 )*), i8 addrspace (4 )* noundef addrspacecast (i8 addrspace (1 )* bitcast (%"class.cl::sycl::specialization_id.1" addrspace (1 )* @_ZL8coeff_id2 to i8 addrspace (1 )*) to i8 addrspace (4 )*), i8 addrspace (4 )* noundef null ) #4
42
43
; CHECK-IR: %[[#NS7:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID4:]], float 0.000000e+00)
43
44
; CHECK-IR: %[[#NS8:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID5:]], float 0.000000e+00)
44
45
; CHECK-IR: %[[#NS9:]] = call float @_Z20__spirv_SpecConstantif(i32 [[#SCID6:]], float 0.000000e+00)
@@ -65,6 +66,8 @@ attributes #4 = { convergent }
65
66
!opencl.spir.version = !{!3 }
66
67
!spirv.Source = !{!4 }
67
68
!llvm.ident = !{!5 }
69
+ ; CHECK-IR: !sycl.specialization-constants = !{![[#MN0:]], ![[#MN1:]]}
70
+ ; CHECK-IR: !sycl.specialization-constants-default-values = !{![[#MN2:]], ![[#MN3:]]}
68
71
69
72
!0 = !{!"libcpmt" }
70
73
!1 = !{i32 1 , !"wchar_size" , i32 2 }
@@ -74,6 +77,10 @@ attributes #4 = { convergent }
74
77
!5 = !{!"clang version 14.0.0" }
75
78
!6 = !{i32 -1 }
76
79
!7 = !{i1 true }
80
+ ; CHECK-IR: ![[#MN0]] = !{!"ef880fa09cf7a9d7____ZL8coeff_id", i32 0, i32 0, i32 4, i32 1, i32 4, i32 4, i32 2, i32 8, i32 4, i32 3, i32 16, i32 8, i32 -1, i32 24, i32 8}
81
+ ; CHECK-IR: ![[#MN1]] = !{!"df991fa0adf9bad8____ZL8coeff_id2", i32 5, i32 0, i32 4, i32 6, i32 4, i32 4, i32 7, i32 8, i32 4, i32 8, i32 16, i32 8, i32 -1, i32 31, i32 1}
82
+ ; CHECK-IR: ![[#MN2]] = !{%struct.coeff_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [8 x i8] undef }}
83
+ ; CHECK-IR: ![[#MN3]] = !{%struct.coeff2_str_aligned_t { %"class.std::array" zeroinitializer, i64 0, [7 x i8] undef, i8 undef }}
77
84
78
85
; CHECK-PROP: [SYCL/specialization constants]
79
86
; CHECK-PROP-NEXT: ef880fa09cf7a9d7____ZL8coeff_id=2|
0 commit comments