forked from intel/llvm
-
Notifications
You must be signed in to change notification settings - Fork 0
Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[SYCL] Test corrections after moving bfloat16 support out of experime…
…ntal status. (intel/llvm-test-suite#1129) Tests changes for intel#6524 Signed-off-by: Rajiv Deodhar <rajiv.deodhar@intel.com> Co-authored-by: JackAKirk <jack.kirk@codeplay.com>
- Loading branch information
Showing
24 changed files
with
239 additions
and
82 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,70 @@ | ||
| // UNSUPPORTED: hip | ||
| // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-device-code-split=per_kernel %s -o %t.out | ||
| // Currently the feature is supported only on CPU and GPU, natively or by | ||
| // software emulation. | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
| // RUNx: %ACC_RUN_PLACEHOLDER %t.out | ||
|
|
||
| //==---------- bfloat16_conversions.cpp - SYCL bfloat16 type test ---------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===---------------------------------------------------------------------===// | ||
|
|
||
| #include <iostream> | ||
| #include <sycl/sycl.hpp> | ||
|
|
||
| using namespace sycl; | ||
|
|
||
| template <typename T> T calculate(T a, T b) { | ||
| sycl::ext::oneapi::bfloat16 x = -a; | ||
| sycl::ext::oneapi::bfloat16 y = b; | ||
| sycl::ext::oneapi::bfloat16 z = x + y; | ||
| T result = z; | ||
| return result; | ||
| } | ||
|
|
||
| template <typename T> int test_device(queue Q) { | ||
| T data[3] = {-7.0f, 8.1f, 0.0f}; | ||
|
|
||
| buffer<T, 1> buf{data, 3}; | ||
| Q.submit([&](handler &cgh) { | ||
| accessor numbers{buf, cgh, read_write}; | ||
| cgh.single_task([=]() { numbers[2] = calculate(numbers[0], numbers[1]); }); | ||
| }); | ||
|
|
||
| host_accessor hostOutAcc{buf, read_only}; | ||
| std::cout << "Device Result = " << hostOutAcc[2] << std::endl; | ||
| if (hostOutAcc[2] == 15.125f) | ||
| return 0; | ||
| return 1; | ||
| } | ||
|
|
||
| template <typename T> int test_host() { | ||
| T a{-5.6f}; | ||
| T b{-1.1f}; | ||
| T result = calculate(a, b); | ||
| std::cout << "Host Result = " << result << std::endl; | ||
| if (result == 4.5f) | ||
| return 0; | ||
| return 1; | ||
| } | ||
|
|
||
| int main() { | ||
| queue Q; | ||
| int result; | ||
| result = test_host<sycl::half>(); | ||
| result |= test_host<float>(); | ||
| if (Q.get_device().has(aspect::fp16)) | ||
| result |= test_device<sycl::half>(Q); | ||
| result |= test_device<float>(Q); | ||
| if (result) | ||
| std::cout << "FAIL\n"; | ||
| else | ||
| std::cout << "PASS\n"; | ||
|
|
||
| return result; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,83 @@ | ||
| /// | ||
| /// Check if bfloat16 example works using fallback libraries | ||
| /// | ||
|
|
||
| // REQUIRES: opencl-aot, ocloc, cpu, gpu-intel-gen9 | ||
| // UNSUPPORTED: cuda | ||
| // CUDA is not compatible with SPIR. | ||
|
|
||
| // RUN: %clangxx -fsycl %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device gen9" %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device *" %s -o %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen "-device gen9" %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
| // RUN: %GPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64,spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64,spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" %s -o %t.out | ||
| // RUN: %CPU_RUN_PLACEHOLDER %t.out | ||
|
|
||
| #include <sycl/sycl.hpp> | ||
|
|
||
| using namespace sycl; | ||
| using sycl::ext::oneapi::bfloat16; | ||
|
|
||
| float foo(float a, float b) { | ||
| // Convert from float to bfloat16. | ||
| bfloat16 A{a}; | ||
| bfloat16 B{b}; | ||
|
|
||
| // Convert A and B from bfloat16 to float, do addition on floating-point | ||
| // numbers, then convert the result to bfloat16 and store it in C. | ||
| bfloat16 C = A + B; | ||
|
|
||
| // Return the result converted from bfloat16 to float. | ||
| return C; | ||
| } | ||
|
|
||
| int main(int argc, char *argv[]) { | ||
| float data[3] = {7.0f, 8.1f, 0.0f}; | ||
|
|
||
| float result_host = foo(7.0f, 8.1f); | ||
| std::cout << "CPU Result = " << result_host << std::endl; | ||
| if (std::abs(15.1f - result_host) > 0.1f) { | ||
| std::cout << "Test failed. Expected CPU Result ~= 15.1" << std::endl; | ||
| return 1; | ||
| } | ||
|
|
||
| queue deviceQueue; | ||
| buffer<float, 1> buf{data, 3}; | ||
|
|
||
| deviceQueue.submit([&](handler &cgh) { | ||
| accessor numbers{buf, cgh, read_write}; | ||
| cgh.single_task([=]() { numbers[2] = foo(numbers[0], numbers[1]); }); | ||
| }); | ||
|
|
||
| host_accessor hostOutAcc{buf, read_only}; | ||
| float result_device = hostOutAcc[2]; | ||
| std::cout << "GPU Result = " << result_device << std::endl; | ||
| if (std::abs(result_host - result_device) > 0.1f) { | ||
| std::cout << "Test failed. CPU Result !~= GPU result" << std::endl; | ||
| return 1; | ||
| } | ||
|
|
||
| return 0; | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Oops, something went wrong.