|
1 |
| -// RUN: %clangxx -fsycl-device-only -S -Xclang -emit-llvm %s -o - | FileCheck %s |
| 1 | +// RUN: %clangxx -fsycl-device-only -fsycl-targets=%sycl_triple -S %s -o - | FileCheck %s |
2 | 2 |
|
3 |
| -#include <sycl/sycl.hpp> |
4 | 3 | #include <sycl/ext/intel/experimental/bfloat16.hpp>
|
| 4 | +#include <sycl/sycl.hpp> |
5 | 5 |
|
6 | 6 | using sycl::ext::intel::experimental::bfloat16;
|
7 | 7 |
|
8 | 8 | SYCL_EXTERNAL uint16_t some_bf16_intrinsic(uint16_t x, uint16_t y);
|
9 | 9 |
|
10 |
| -__attribute__((noinline)) |
11 |
| -float op(float a, float b) { |
12 |
| - bfloat16 A {a}; |
13 |
| -// CHECK: [[A:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float %a) |
14 |
| -// CHECK-NOT: fptoui |
| 10 | +__attribute__((noinline)) float op(float a, float b) { |
| 11 | + // CHECK: define {{.*}} spir_func float @_Z2opff(float [[a:%.*]], float [[b:%.*]]) |
| 12 | + bfloat16 A{a}; |
| 13 | + // CHECK: [[A:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float [[a]]) |
| 14 | + // CHECK-NOT: fptoui |
15 | 15 |
|
16 |
| - bfloat16 B {b}; |
17 |
| -// CHECK: [[B:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float %b) |
18 |
| -// CHECK-NOT: fptoui |
| 16 | + bfloat16 B{b}; |
| 17 | + // CHECK: [[B:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float [[b]]) |
| 18 | + // CHECK-NOT: fptoui |
19 | 19 |
|
20 | 20 | bfloat16 C = A + B;
|
21 |
| -// CHECK: [[A_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[A]]) |
22 |
| -// CHECK: [[B_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[B]]) |
23 |
| -// CHECK: [[Add:%.*]] = fadd float [[A_float]], [[B_float]] |
24 |
| -// CHECK: [[C:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float [[Add]]) |
25 |
| -// CHECK-NOT: uitofp |
26 |
| -// CHECK-NOT: fptoui |
| 21 | + // CHECK: [[A_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[A]]) |
| 22 | + // CHECK: [[B_float:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[B]]) |
| 23 | + // CHECK: [[Add:%.*]] = fadd float [[A_float]], [[B_float]] |
| 24 | + // CHECK: [[C:%.*]] = tail call spir_func zeroext i16 @_Z27__spirv_ConvertFToBF16INTELf(float [[Add]]) |
| 25 | + // CHECK-NOT: uitofp |
| 26 | + // CHECK-NOT: fptoui |
27 | 27 |
|
28 | 28 | bfloat16 D = some_bf16_intrinsic(A, C);
|
29 |
| -// CHECK: [[D:%.*]] = tail call spir_func zeroext i16 @_Z19some_bf16_intrinsictt(i16 zeroext [[A]], i16 zeroext [[C]]) |
30 |
| -// CHECK-NOT: uitofp |
31 |
| -// CHECK-NOT: fptoui |
| 29 | + // CHECK: [[D:%.*]] = tail call spir_func zeroext i16 @_Z19some_bf16_intrinsictt(i16 zeroext [[A]], i16 zeroext [[C]]) |
| 30 | + // CHECK-NOT: uitofp |
| 31 | + // CHECK-NOT: fptoui |
32 | 32 |
|
33 | 33 | return D;
|
34 |
| -// CHECK: [[RetVal:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[D]]) |
35 |
| -// CHECK: ret float [[RetVal]] |
36 |
| -// CHECK-NOT: uitofp |
37 |
| -// CHECK-NOT: fptoui |
| 34 | + // CHECK: [[RetVal:%.*]] = tail call spir_func float @_Z27__spirv_ConvertBF16ToFINTELt(i16 zeroext [[D]]) |
| 35 | + // CHECK: ret float [[RetVal]] |
| 36 | + // CHECK-NOT: uitofp |
| 37 | + // CHECK-NOT: fptoui |
38 | 38 | }
|
39 | 39 |
|
40 | 40 | int main(int argc, char *argv[]) {
|
|
0 commit comments