[SYCL][New offload model] Add SYCL E2E tests for --offload-new-driver option and fix failing tests #14730
There are no files selected for viewing
This file contains hidden or 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 hidden or 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 hidden or 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 hidden or 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 hidden or 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 hidden or 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,71 @@ | ||
| //==--- aot.cpp - Simple vector addition (AOT compilation example) --------==// | ||
| // | ||
| // 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 <sycl/detail/core.hpp> | ||
|
|
||
| #include <array> | ||
| #include <iostream> | ||
|
|
||
| constexpr sycl::access::mode sycl_read = sycl::access::mode::read; | ||
| constexpr sycl::access::mode sycl_write = sycl::access::mode::write; | ||
|
|
||
| template <typename T> class Vadd; | ||
|
|
||
| template <typename T, size_t N> | ||
| void vadd(const std::array<T, N> &A, const std::array<T, N> &B, | ||
| std::array<T, N> &C) { | ||
| sycl::queue Queue([](sycl::exception_list ExceptionList) { | ||
| for (std::exception_ptr ExceptionPtr : ExceptionList) { | ||
| try { | ||
| std::rethrow_exception(ExceptionPtr); | ||
| } catch (sycl::exception &E) { | ||
| std::cerr << E.what(); | ||
| } catch (...) { | ||
| std::cerr << "Unknown async exception was caught." << std::endl; | ||
| } | ||
| } | ||
| }); | ||
|
|
||
| sycl::range<1> numOfItems{N}; | ||
| sycl::buffer bufA(A.data(), numOfItems); | ||
| sycl::buffer bufB(B.data(), numOfItems); | ||
| sycl::buffer bufC(C.data(), numOfItems); | ||
|
|
||
| Queue.submit([&](sycl::handler &cgh) { | ||
| sycl::accessor accA{bufA, cgh, sycl::read_only}; | ||
| sycl::accessor accB{bufB, cgh, sycl::read_only}; | ||
| sycl::accessor accC{bufC, cgh, sycl::write_only}; | ||
|
|
||
| cgh.parallel_for<Vadd<T>>(numOfItems, [=](sycl::id<1> wiID) { | ||
| accC[wiID] = accA[wiID] + accB[wiID]; | ||
| }); | ||
| }); | ||
|
|
||
| Queue.wait_and_throw(); | ||
| } | ||
|
|
||
| int main() { | ||
| const size_t array_size = 4; | ||
| std::array<int, array_size> A = {{1, 2, 3, 4}}, B = {{1, 2, 3, 4}}, C; | ||
| std::array<float, array_size> D = {{1.f, 2.f, 3.f, 4.f}}, | ||
| E = {{1.f, 2.f, 3.f, 4.f}}, F; | ||
| vadd(A, B, C); | ||
| vadd(D, E, F); | ||
| for (unsigned int i = 0; i < array_size; i++) { | ||
| if (C[i] != A[i] + B[i]) { | ||
| std::cout << "Incorrect result (element " << i << " is " << C[i] << "!\n"; | ||
| return 1; | ||
| } | ||
| if (F[i] != D[i] + E[i]) { | ||
| std::cout << "Incorrect result (element " << i << " is " << F[i] << "!\n"; | ||
| return 1; | ||
| } | ||
| } | ||
| std::cout << "Correct result!\n"; | ||
| return 0; | ||
| } |
This file contains hidden or 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,23 @@ | ||
| #include "split-per-source.h" | ||
|
There was a problem hiding this comment. do we need a copyright header? There was a problem hiding this comment. This file was lifted from here: ./sycl/test-e2e/DeviceCodeSplit/Inputs/split-per-source-second-file.cpp where I do not find a copyright header. |
||
|
|
||
| void runKernelsFromFile2() { | ||
| sycl::queue Q; | ||
| int Data = 0; | ||
| { | ||
| sycl::buffer<int, 1> Buf(&Data, sycl::range<1>(1)); | ||
| auto KernelID1 = sycl::get_kernel_id<File2Kern1>(); | ||
| auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>( | ||
| Q.get_context(), {KernelID1}); | ||
| auto Krn = KB.get_kernel(KernelID1); | ||
|
|
||
| std::vector<sycl::kernel_id> KernelIDStorage = KB.get_kernel_ids(); | ||
| assert(KernelIDStorage.size() == 1); | ||
| assert(KernelIDStorage[0] == KernelID1); | ||
|
|
||
| Q.submit([&](sycl::handler &Cgh) { | ||
| auto Acc = Buf.get_access<sycl::access::mode::read_write>(Cgh); | ||
| Cgh.single_task<File2Kern1>(Krn, [=]() { Acc[0] = 3; }); | ||
| }); | ||
| } | ||
| assert(Data == 3); | ||
| } | ||
This file contains hidden or 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,7 @@ | ||
| #include <sycl/detail/core.hpp> | ||
|
|
||
| class File1Kern1; | ||
| class File1Kern2; | ||
| class File2Kern1; | ||
|
|
||
| void runKernelsFromFile2(); |
This file contains hidden or 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,6 @@ | ||
| // REQUIRES: opencl-aot, cpu | ||
|
There was a problem hiding this comment. we also need There was a problem hiding this comment. I am not aware of this setting. But I can add that here. Again this test was lifted from ./sycl/test-e2e/DeviceCodeSplit/aot-cpu.cpp and I added --offload-new-driver |
||
|
|
||
| // Test with `--offload-new-driver` | ||
| // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source -fsycl-targets=spir64_x86_64 -I %S/Inputs -o %t.out %S/split-per-source-main.cpp %S/Inputs/split-per-source-second-file.cpp \ | ||
| // RUN: -fsycl-dead-args-optimization --offload-new-driver | ||
|
There was a problem hiding this comment.
There was a problem hiding this comment. I have lifted this file again from ./sycl/test-e2e/DeviceCodeSplit/aot-cpu.cpp. May be we can have a cleanup PR later on? |
||
| // RUN: %{run} %t.out | ||
This file contains hidden or 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,13 @@ | ||
| // REQUIRES: ocloc, gpu | ||
| // UNSUPPORTED: cuda || hip | ||
| // CUDA does neither support device code splitting nor SPIR. | ||
| // Test with `--offload-new-driver` | ||
| // | ||
| // RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ | ||
| // RUN: -fsycl-targets=spir64_gen \ | ||
| // RUN: -Xsycl-target-backend=spir64_gen \ | ||
| // RUN: "-device tgllp" -I %S/Inputs -o %t.out \ | ||
| // RUN: %S/split-per-source-main.cpp \ | ||
| // RUN: %S/Inputs/split-per-source-second-file.cpp \ | ||
| // RUN: -fsycl-dead-args-optimization --offload-new-driver | ||
| // RUN: %{run} %t.out |
This file contains hidden or 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,20 @@ | ||
| //==--- cpu.cpp - AOT compilation for cpu devices using opencl-aot --------==// | ||
| // | ||
| // 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 | ||
| // | ||
| //===---------------------------------------------------------------------===// | ||
|
|
||
| // REQUIRES: opencl-aot, cpu | ||
|
There was a problem hiding this comment. need There was a problem hiding this comment. Same as above comments. Thanks |
||
|
|
||
| // Test with `--offload-new-driver` | ||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %S/Inputs/aot.cpp -o %t.out | ||
| // RUN: %{run} %t.out | ||
|
|
||
| // Test that opencl-aot can handle multiple build options. | ||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %S/Inputs/aot.cpp -Xsycl-target-backend "--bo=-g" -Xsycl-target-backend "--bo=-cl-opt-disable" -o %t2.out | ||
|
|
||
| // Test that opencl-aot can handle march option. | ||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %S/Inputs/aot.cpp -Xsycl-target-backend "--march=avx512" | ||
| // RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 --offload-new-driver %S/Inputs/aot.cpp -Xsycl-target-backend "--march=wsm" | ||
Oops, something went wrong.
Oops, something went wrong.
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
nit: if we are returning, do we need to
std::move?There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I see that std::move has been used wherever we are returning Err. I followed the trend.