Skip to content

Commit 9bf9af4

Browse files
author
Yihan Wang
authored
[SYCLomatic] Address compiler breaking change of sycl math APIs (#2259)
Signed-off-by: Wang, Yihan <[email protected]>
1 parent 2398edb commit 9bf9af4

File tree

9 files changed

+36
-43
lines changed

9 files changed

+36
-43
lines changed

‎clang/lib/DPCT/AsmMigration.cpp‎

Lines changed: 4 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1491,12 +1491,7 @@ class SYCLGen : public SYCLGenBase {
14911491
if (tryEmitStmt(Op, Inst->getInputOperand(0)))
14921492
return SYCLGenError();
14931493

1494-
std::string TypeString;
1495-
if (tryEmitType(TypeString, Inst->getType(0)))
1496-
return SYCLGenError();
1497-
1498-
std::string ReplaceString =
1499-
MapNames::getClNamespace() + MathFn.str() + "<" + TypeString + ">(";
1494+
std::string ReplaceString = MapNames::getClNamespace() + MathFn.str() + '(';
15001495
if (Inst->getOpcode() == asmtok::op_ex2)
15011496
ReplaceString += "2, ";
15021497
ReplaceString += Op + ")";
@@ -1541,15 +1536,13 @@ class SYCLGen : public SYCLGenBase {
15411536
if (emitStmt(Inst->getOutputOperand()))
15421537
return SYCLGenError();
15431538
OS() << " = ";
1544-
std::string Op[3], TypeString;
1539+
std::string Op[3];
15451540
for (int i = 0; i < 3; ++i)
15461541
if (tryEmitStmt(Op[i], Inst->getInputOperand(i)))
15471542
return SYCLGenError();
1548-
if (tryEmitType(TypeString, Inst->getType(0)))
1549-
return SYCLGenError();
15501543

1551-
OS() << MapNames::getClNamespace() << "abs_diff<" << TypeString << ">("
1552-
<< Op[0] << ", " << Op[1] << ") + " << Op[2];
1544+
OS() << MapNames::getClNamespace() << "abs_diff(" << Op[0] << ", " << Op[1]
1545+
<< ") + " << Op[2];
15531546
endstmt();
15541547
return SYCLGenSuccess();
15551548
}

‎clang/test/dpct/asm/cos.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,10 @@
1111
__global__ void cos() {
1212
float f32;
1313

14-
// CHECK: f32 = sycl::cos<float>(1.0f);
14+
// CHECK: f32 = sycl::cos(1.0f);
1515
asm("cos.approx.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1616

17-
// CHECK: f32 = sycl::cos<float>(1.0f);
17+
// CHECK: f32 = sycl::cos(1.0f);
1818
asm("cos.approx.ftz.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1919
}
2020

‎clang/test/dpct/asm/ex2.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,10 +10,10 @@
1010

1111
__global__ void ex2() {
1212
float f32;
13-
// CHECK: f32 = sycl::pow<float>(2, 3.4f);
13+
// CHECK: f32 = sycl::pow(2, 3.4f);
1414
asm("ex2.approx.f32 %0, %1;" : "=f"(f32) : "f"(3.4f));
1515

16-
// CHECK: f32 = sycl::pow<float>(2, 3.4f);
16+
// CHECK: f32 = sycl::pow(2, 3.4f);
1717
asm("ex2.approx.ftz.f32 %0, %1;" : "=f"(f32) : "f"(3.4f));
1818
}
1919

‎clang/test/dpct/asm/lg2.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,10 +10,10 @@
1010

1111
__global__ void lg2() {
1212
float f32;
13-
// CHECK: f32 = sycl::log2<float>(1.0f);
13+
// CHECK: f32 = sycl::log2(1.0f);
1414
asm("lg2.approx.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1515

16-
// CHECK: f32 = sycl::log2<float>(1.0f);
16+
// CHECK: f32 = sycl::log2(1.0f);
1717
asm("lg2.approx.ftz.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1818
}
1919

‎clang/test/dpct/asm/rsqrt.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -12,13 +12,13 @@ __global__ void rsqrt() {
1212
float f32;
1313
double f64;
1414

15-
// CHECK: f32 = sycl::rsqrt<float>(1.0f);
15+
// CHECK: f32 = sycl::rsqrt(1.0f);
1616
asm("rsqrt.approx.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1717

18-
// CHECK: f32 = sycl::rsqrt<float>(1.0f);
18+
// CHECK: f32 = sycl::rsqrt(1.0f);
1919
asm("rsqrt.approx.ftz.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
2020

21-
// CHECK: f64 = sycl::rsqrt<double>(1.0);
21+
// CHECK: f64 = sycl::rsqrt(1.0);
2222
asm("rsqrt.approx.f64 %0, %1;" : "=d"(f64) : "d"(1.0));
2323
}
2424
// clang-format on

‎clang/test/dpct/asm/sad.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -16,22 +16,22 @@ __global__ void sad() {
1616
int64_t i64;
1717
uint64_t u64;
1818

19-
// CHECK: i16 = sycl::abs_diff<int16_t>(i16, i16) + i16;
19+
// CHECK: i16 = sycl::abs_diff(i16, i16) + i16;
2020
asm("sad.s16 %0, %1, %2, %3;" : "=h"(i16) : "h"(i16), "h"(i16), "h"(i16));
2121

22-
// CHECK: u16 = sycl::abs_diff<uint16_t>(u16, u16) + u16;
22+
// CHECK: u16 = sycl::abs_diff(u16, u16) + u16;
2323
asm("sad.u16 %0, %1, %2, %3;" : "=h"(u16) : "h"(u16), "h"(u16), "h"(u16));
2424

25-
// CHECK: i32 = sycl::abs_diff<int32_t>(i32, i32) + i32;
25+
// CHECK: i32 = sycl::abs_diff(i32, i32) + i32;
2626
asm("sad.s32 %0, %1, %2, %3;" : "=r"(i32) : "r"(i32), "r"(i32), "r"(i32));
2727

28-
// CHECK: u32 = sycl::abs_diff<uint32_t>(u32, u32) + u32;
28+
// CHECK: u32 = sycl::abs_diff(u32, u32) + u32;
2929
asm("sad.u32 %0, %1, %2, %3;" : "=r"(u32) : "r"(u32), "r"(u32), "r"(u32));
3030

31-
// CHECK: i64 = sycl::abs_diff<int64_t>(i64, i64) + i64;
31+
// CHECK: i64 = sycl::abs_diff(i64, i64) + i64;
3232
asm("sad.s64 %0, %1, %2, %3;" : "=l"(i64) : "l"(i64), "l"(i64), "l"(i64));
3333

34-
// CHECK: u64 = sycl::abs_diff<uint64_t>(u64, u64) + u64;
34+
// CHECK: u64 = sycl::abs_diff(u64, u64) + u64;
3535
asm("sad.u64 %0, %1, %2, %3;" : "=l"(u64) : "l"(u64), "l"(u64), "l"(u64));
3636
}
3737

‎clang/test/dpct/asm/sin.cu‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,10 +11,10 @@
1111
__global__ void sin() {
1212
float f32;
1313

14-
// CHECK: f32 = sycl::sin<float>(1.0f);
14+
// CHECK: f32 = sycl::sin(1.0f);
1515
asm("sin.approx.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1616

17-
// CHECK: f32 = sycl::sin<float>(1.0f);
17+
// CHECK: f32 = sycl::sin(1.0f);
1818
asm("sin.approx.ftz.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1919
}
2020

‎clang/test/dpct/asm/sqrt.cu‎

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -11,56 +11,56 @@
1111
// CHECK: void sqrt() {
1212
// CHECK-NEXT: float f32;
1313
// CHECK-NEXT: double f64;
14-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
15-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
14+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
15+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
1616
// CHECK-NEXT: /*
1717
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
1818
// CHECK-NEXT: */
19-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
19+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
2020
// CHECK-NEXT: /*
2121
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
2222
// CHECK-NEXT: */
23-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
23+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
2424
// CHECK-NEXT: /*
2525
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
2626
// CHECK-NEXT: */
27-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
27+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
2828
// CHECK-NEXT: /*
2929
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
3030
// CHECK-NEXT: */
31-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
31+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
3232
// CHECK-NEXT: /*
3333
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
3434
// CHECK-NEXT: */
35-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
35+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
3636
// CHECK-NEXT: /*
3737
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
3838
// CHECK-NEXT: */
39-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
39+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
4040
// CHECK-NEXT: /*
4141
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
4242
// CHECK-NEXT: */
43-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
43+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
4444
// CHECK-NEXT: /*
4545
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
4646
// CHECK-NEXT: */
47-
// CHECK-NEXT: f32 = sycl::sqrt<float>(1.0f);
47+
// CHECK-NEXT: f32 = sycl::sqrt(1.0f);
4848
// CHECK-NEXT: /*
4949
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
5050
// CHECK-NEXT: */
51-
// CHECK-NEXT: f64 = sycl::sqrt<double>(1.0);
51+
// CHECK-NEXT: f64 = sycl::sqrt(1.0);
5252
// CHECK-NEXT: /*
5353
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
5454
// CHECK-NEXT: */
55-
// CHECK-NEXT: f64 = sycl::sqrt<double>(1.0);
55+
// CHECK-NEXT: f64 = sycl::sqrt(1.0);
5656
// CHECK-NEXT: /*
5757
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
5858
// CHECK-NEXT: */
59-
// CHECK-NEXT: f64 = sycl::sqrt<double>(1.0);
59+
// CHECK-NEXT: f64 = sycl::sqrt(1.0);
6060
// CHECK-NEXT: /*
6161
// CHECK-NEXT: DPCT1013:{{.*}}: The rounding mode could not be specified and the generated code may have different accuracy than the original code. Verify the correctness. SYCL math built-in function rounding mode is aligned with OpenCL C 1.2 standard.
6262
// CHECK-NEXT: */
63-
// CHECK-NEXT: f64 = sycl::sqrt<double>(1.0);
63+
// CHECK-NEXT: f64 = sycl::sqrt(1.0);
6464
// CHECK-NEXT: }
6565
__global__ void sqrt() {
6666
float f32;

‎clang/test/dpct/asm/tanh.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,7 @@
1111
__global__ void tanh() {
1212
float f32;
1313

14-
// CHECK: f32 = sycl::tanh<float>(1.0f);
14+
// CHECK: f32 = sycl::tanh(1.0f);
1515
asm("tanh.approx.f32 %0, %1;" : "=f"(f32) : "f"(1.0f));
1616
}
1717

0 commit comments

Comments
 (0)