@@ -13,17 +13,17 @@ SYCL_EXTERNAL void foo(long x, sycl::half y);
13
13
__attribute__ ((noinline)) float op(float a, float b) {
14
14
// CHECK: define {{.*}} spir_func float @_Z2opff(float [[a:%.*]], float [[b:%.*]])
15
15
bfloat16 A{a};
16
- // CHECK: [[A:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}} [[a]].addr.ascast )
16
+ // CHECK: [[A:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}})
17
17
// CHECK-NOT: fptoui
18
18
19
19
bfloat16 B{b};
20
- // CHECK: [[B:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}} [[b]].addr.ascast )
20
+ // CHECK: [[B:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}})
21
21
// CHECK-NOT: fptoui
22
22
23
23
bfloat16 C = A + B;
24
- // CHECK: [[RTCASTI:%ref.tmp.ascast.i ]] = addrspacecast float* [[RT:%ref.tmp.i ]] to float addrspace(4)*
25
- // CHECK: [[A_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}} %1 )
26
- // CHECK: [[B_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}} %4 )
24
+ // CHECK: [[RTCASTI:%.* ]] = addrspacecast float* [[RT:%.* ]] to float addrspace(4)*
25
+ // CHECK: [[A_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}})
26
+ // CHECK: [[B_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}})
27
27
// CHECK: [[Add:%.*]] = fadd float [[A_float]], [[B_float]]
28
28
// CHECK: store float [[Add]], float* [[RT]], align 4
29
29
// CHECK: [[C:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}}) [[RTCASTI]])
@@ -32,22 +32,22 @@ __attribute__((noinline)) float op(float a, float b) {
32
32
// CHECK-NOT: fptoui
33
33
34
34
long L = bfloat16 (3 .14f );
35
- // CHECK: [[L:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}} %ref.tmp1.ascast )
35
+ // CHECK: [[L:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}})
36
36
// CHECK: [[P8:%.*]] = addrspacecast i16* [[VI9:%.*]] to i16 addrspace(4)*
37
37
// CHECK: store i16 [[L]], i16* [[VI9]]
38
38
// CHECK: [[L_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}} [[P8]])
39
39
// CHECK: [[L:%.*]] = fptosi float [[L_float]] to i{{32|64}}
40
40
41
41
sycl::half H = bfloat16 (2 .71f );
42
- // CHECK: [[H:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}} %ref.tmp3.ascast )
42
+ // CHECK: [[H:%.*]] = call spir_func zeroext i16 @__devicelib_ConvertFToBF16INTEL(float {{.*}})
43
43
// CHECK: [[P11:%.*]] = addrspacecast i16* [[VI13:%.*]] to i16 addrspace(4)*
44
44
// CHECK: store i16 [[H]], i16* [[VI13]], align 2
45
45
// CHECK: [[H_float:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}} [[P11]])
46
46
// CHECK: [[H:%.*]] = fptrunc float [[H_float]] to half
47
47
foo (L, H);
48
48
49
49
return A;
50
- // CHECK: [[RetVal:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}} %2 )
50
+ // CHECK: [[RetVal:%.*]] = call spir_func float @__devicelib_ConvertBF16ToFINTEL(i16 {{.*}})
51
51
// CHECK: ret float [[RetVal]]
52
52
// CHECK-NOT: uitofp
53
53
// CHECK-NOT: fptoui
0 commit comments