From a7472f8cef1de098f2133d9bc6784982bd076bb1 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Fri, 2 Sep 2022 04:09:27 -0700 Subject: [PATCH] [SYCL] Remove host run and dependencies from SYCL/DeprecatedFeatures tests This commit removes the host run and any assumptions and operations related to the host device from the tests in SYCL/DeprecatedFeatures. Co-authored-by: Sachkov, Alexey Signed-off-by: Larsen, Steffen --- .../ESIMD/Inputs/esimd_test_utils.hpp | 7 +- .../FunctionPointers/fp-as-kernel-arg.cpp | 21 +- .../pass-fp-through-buffer.cpp | 21 +- .../SpecConsts1.2.1/composite-in-functor.cpp | 1 - .../SpecConsts1.2.1/composite-type.cpp | 1 - .../multiple-usages-of-composite.cpp | 1 - .../SpecConsts1.2.1/spec_const_hw.cpp | 1 - .../SpecConsts1.2.1/spec_const_neg.cpp | 1 - .../SpecConsts1.2.1/spec_const_redefine.cpp | 1 - .../specialization_constants.cpp | 1 - .../specialization_constants_negative.cpp | 1 - .../specialization_constants_override.cpp | 1 - .../unpacked-composite-type.cpp | 1 - SYCL/DeprecatedFeatures/buffer_interop.cpp | 135 ++++----- SYCL/DeprecatedFeatures/deprecated.cpp | 1 - SYCL/DeprecatedFeatures/get-options.cpp | 13 +- SYCL/DeprecatedFeatures/get_backend.cpp | 53 ++-- .../kernel-and-program-interop.cpp | 281 +++++++++--------- .../DeprecatedFeatures/kernel-and-program.cpp | 80 +++-- SYCL/DeprecatedFeatures/kernel_interop.cpp | 3 - SYCL/DeprecatedFeatures/opencl-interop.cpp | 4 - SYCL/DeprecatedFeatures/platform.cpp | 21 +- SYCL/DeprecatedFeatures/program_link.cpp | 3 - SYCL/DeprecatedFeatures/queue_old_interop.cpp | 18 +- SYCL/DeprecatedFeatures/set_arg_interop.cpp | 190 ++++++------ 25 files changed, 393 insertions(+), 468 deletions(-) diff --git a/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp b/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp index 97be177788..d4e0d492d6 100644 --- a/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp +++ b/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp @@ -32,17 +32,14 @@ namespace esimd_test { // was returned for all devices, then the selection process will cause an // exception. class ESIMDSelector : public device_selector { - // Require GPU device unless HOST is requested in SYCL_DEVICE_FILTER env + // Require GPU device virtual int operator()(const device &device) const { if (const char *dev_filter = getenv("SYCL_DEVICE_FILTER")) { std::string filter_string(dev_filter); if (filter_string.find("gpu") != std::string::npos) return device.is_gpu() ? 1000 : -1; - if (filter_string.find("host") != std::string::npos) - return device.is_host() ? 1000 : -1; std::cerr - << "Supported 'SYCL_DEVICE_FILTER' env var values are 'gpu' and " - "'host', '" + << "Supported 'SYCL_DEVICE_FILTER' env var values is 'gpu' and '" << filter_string << "' does not contain such substrings.\n"; return -1; } diff --git a/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp b/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp index 63ebc7426e..c53a3f7b0b 100644 --- a/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp +++ b/SYCL/DeprecatedFeatures/FunctionPointers/fp-as-kernel-arg.cpp @@ -3,7 +3,6 @@ // CUDA does not support the function pointer as kernel argument extension. // RUN: %clangxx -Xclang -fsycl-allow-func-ptr -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // FIXME: This test should use runtime early exit once correct check for @@ -34,17 +33,15 @@ int main() { sycl::kernel KE = P.get_kernel(); auto FptrStorage = sycl::ext::oneapi::get_device_func_ptr(&add, "add", P, D); - if (!D.is_host()) { - // FIXME: update this check with query to supported extension - // For now, we don't have runtimes that report required OpenCL extension and - // it is hard to understand should this functionality be supported or not. - // So, let's skip this test if FptrStorage is 0, which means that by some - // reason we failed to obtain device function pointer. Just to avoid false - // alarms - if (0 == FptrStorage) { - std::cout << "Test PASSED. (it was actually skipped)" << std::endl; - return 0; - } + // FIXME: update this check with query to supported extension + // For now, we don't have runtimes that report required OpenCL extension and + // it is hard to understand should this functionality be supported or not. + // So, let's skip this test if FptrStorage is 0, which means that by some + // reason we failed to obtain device function pointer. Just to avoid false + // alarms + if (0 == FptrStorage) { + std::cout << "Test PASSED. (it was actually skipped)" << std::endl; + return 0; } sycl::buffer BufA(A.data(), sycl::range<1>(Size)); diff --git a/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp b/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp index f0cfcc3da2..fab00029c7 100644 --- a/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp +++ b/SYCL/DeprecatedFeatures/FunctionPointers/pass-fp-through-buffer.cpp @@ -3,7 +3,6 @@ // CUDA does not support the function pointer as kernel argument extension. // RUN: %clangxx -Xclang -fsycl-allow-func-ptr -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // FIXME: This test should use runtime early exit once correct check for @@ -40,17 +39,15 @@ int main() { auto DTAcc = DispatchTable.get_access(); DTAcc[0] = sycl::ext::oneapi::get_device_func_ptr(&add, "add", P, D); DTAcc[1] = sycl::ext::oneapi::get_device_func_ptr(&sub, "sub", P, D); - if (!D.is_host()) { - // FIXME: update this check with query to supported extension - // For now, we don't have runtimes that report required OpenCL extension - // and it is hard to understand should this functionality be supported or - // not. So, let's skip this test if DTAcc[i] is 0, which means that by - // some reason we failed to obtain device function pointer. Just to avoid - // false alarms - if (0 == DTAcc[0] || 0 == DTAcc[1]) { - std::cout << "Test PASSED. (it was actually skipped)" << std::endl; - return 0; - } + // FIXME: update this check with query to supported extension + // For now, we don't have runtimes that report required OpenCL extension + // and it is hard to understand should this functionality be supported or + // not. So, let's skip this test if DTAcc[i] is 0, which means that by + // some reason we failed to obtain device function pointer. Just to avoid + // false alarms + if (0 == DTAcc[0] || 0 == DTAcc[1]) { + std::cout << "Test PASSED. (it was actually skipped)" << std::endl; + return 0; } } diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp index 9ef1e05e50..62ab07faf8 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp @@ -1,7 +1,6 @@ // UNSUPPORTED: cuda || hip // // RUN: %clangxx -fsycl %s -D__SYCL_INTERNAL_API -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER // diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp index 61ed328be8..5795d0503c 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp @@ -1,7 +1,6 @@ // UNSUPPORTED: cuda || hip // // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp index 3259291f3e..815fc3e857 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp @@ -1,7 +1,6 @@ // UNSUPPORTED: cuda || hip // // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -v -// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER // diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp index 332879dda8..c7e1d6cd50 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp @@ -1,7 +1,6 @@ // UNSUPPORTED: cuda || hip // // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp index db018b03ba..4c363f9a49 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // Specialization constants are not supported on FPGA h/w and emulator. diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp index 59ac0dcdc9..0caf33700f 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp @@ -3,7 +3,6 @@ // FIXME Disable fallback assert so that it doesn't interferes with number of // program builds at run-time // RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out 2>&1 %CPU_CHECK_PLACEHOLDER // RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER // diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp index ae4e4e9f7c..391876159e 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // Specialization constants are not supported on FPGA h/w and emulator. diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp index 24308c3651..c7816b2ca0 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // Specialization constants are not supported on FPGA h/w and emulator. diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp index 16c1b34267..c9459aa869 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // Specialization constants are not supported on FPGA h/w and emulator. diff --git a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp index 2bdf1c4463..c68d9fd26a 100644 --- a/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp +++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp @@ -1,7 +1,6 @@ // UNSUPPORTED: cuda || hip // // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER // diff --git a/SYCL/DeprecatedFeatures/buffer_interop.cpp b/SYCL/DeprecatedFeatures/buffer_interop.cpp index 02b5abf372..183975be03 100644 --- a/SYCL/DeprecatedFeatures/buffer_interop.cpp +++ b/SYCL/DeprecatedFeatures/buffer_interop.cpp @@ -24,87 +24,82 @@ int main() { bool Failed = false; { queue Queue; - if (!Queue.is_host()) { - std::vector Data1(10, -1); - std::vector Data2(10, -2); - { - buffer BufferA(Data1.data(), range<1>(10)); - buffer BufferB(Data2); + std::vector Data1(10, -1); + std::vector Data2(10, -2); + { + buffer BufferA(Data1.data(), range<1>(10)); + buffer BufferB(Data2); - program Program(Queue.get_context()); - Program.build_with_source( - "kernel void override_source(global int* Acc) " - "{Acc[get_global_id(0)] = 0; }\n"); - sycl::kernel Kernel = Program.get_kernel("override_source"); - Queue.submit([&](handler &CGH) { - auto AccA = BufferA.get_access(CGH); - CGH.set_arg(0, AccA); - auto AccB = BufferB.get_access(CGH); - CGH.parallel_for(sycl::range<1>(10), Kernel); - }); - } // Data is copied back - for (int i = 0; i < 10; i++) { - if (Data2[i] != -2) { - std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " - << -2 << std::endl; - assert(false); - Failed = true; - } + program Program(Queue.get_context()); + Program.build_with_source("kernel void override_source(global int* Acc) " + "{Acc[get_global_id(0)] = 0; }\n"); + sycl::kernel Kernel = Program.get_kernel("override_source"); + Queue.submit([&](handler &CGH) { + auto AccA = BufferA.get_access(CGH); + CGH.set_arg(0, AccA); + auto AccB = BufferB.get_access(CGH); + CGH.parallel_for(sycl::range<1>(10), Kernel); + }); + } // Data is copied back + for (int i = 0; i < 10; i++) { + if (Data2[i] != -2) { + std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " << -2 + << std::endl; + assert(false); + Failed = true; } - for (int i = 0; i < 10; i++) { - if (Data1[i] != 0) { - std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " - << 0 << std::endl; - assert(false); - Failed = true; - } + } + for (int i = 0; i < 10; i++) { + if (Data1[i] != 0) { + std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " << 0 + << std::endl; + assert(false); + Failed = true; } } } { queue Queue; - if (!Queue.is_host()) { - std::vector Data1(10, -1); - std::vector Data2(10, -2); - { - buffer BufferA(Data1.data(), range<1>(10)); - buffer BufferB(Data2); - accessor - AccA(BufferA); - accessor - AccB(BufferB); + std::vector Data1(10, -1); + std::vector Data2(10, -2); + { + buffer BufferA(Data1.data(), range<1>(10)); + buffer BufferB(Data2); + accessor + AccA(BufferA); + accessor + AccB(BufferB); - program Program(Queue.get_context()); - Program.build_with_source( - "kernel void override_source_placeholder(global " - "int* Acc) {Acc[get_global_id(0)] = 0; }\n"); - sycl::kernel Kernel = Program.get_kernel("override_source_placeholder"); + program Program(Queue.get_context()); + Program.build_with_source( + "kernel void override_source_placeholder(global " + "int* Acc) {Acc[get_global_id(0)] = 0; }\n"); + sycl::kernel Kernel = Program.get_kernel("override_source_placeholder"); - Queue.submit([&](handler &CGH) { - CGH.require(AccA); - CGH.set_arg(0, AccA); - CGH.require(AccB); - CGH.parallel_for(sycl::range<1>(10), Kernel); - }); - } // Data is copied back - for (int i = 0; i < 10; i++) { - if (Data2[i] != -2) { - std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " - << -2 << std::endl; - assert(false); - Failed = true; - } + Queue.submit([&](handler &CGH) { + CGH.require(AccA); + CGH.set_arg(0, AccA); + CGH.require(AccB); + CGH.parallel_for(sycl::range<1>(10), Kernel); + }); + } // Data is copied back + for (int i = 0; i < 10; i++) { + if (Data2[i] != -2) { + std::cout << " Data2[" << i << "] is " << Data2[i] << " expected " << -2 + << std::endl; + assert(false); + Failed = true; } - for (int i = 0; i < 10; i++) { - if (Data1[i] != 0) { - std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " - << 0 << std::endl; - assert(false); - Failed = true; - } + } + for (int i = 0; i < 10; i++) { + if (Data1[i] != 0) { + std::cout << " Data1[" << i << "] is " << Data1[i] << " expected " << 0 + << std::endl; + assert(false); + Failed = true; } } } diff --git a/SYCL/DeprecatedFeatures/deprecated.cpp b/SYCL/DeprecatedFeatures/deprecated.cpp index 582dc50d25..434910ddf3 100644 --- a/SYCL/DeprecatedFeatures/deprecated.cpp +++ b/SYCL/DeprecatedFeatures/deprecated.cpp @@ -1,7 +1,6 @@ // RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out //==------------- deprecated.cpp - SYCL 2020 deprecation test --------------==// diff --git a/SYCL/DeprecatedFeatures/get-options.cpp b/SYCL/DeprecatedFeatures/get-options.cpp index c0ed222afe..22109cbb46 100644 --- a/SYCL/DeprecatedFeatures/get-options.cpp +++ b/SYCL/DeprecatedFeatures/get-options.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out @@ -34,16 +33,16 @@ int main() { PrgA.build_with_kernel_type(BuildOpts); assert(PrgA.get_compile_options().empty()); assert(PrgA.get_link_options().empty()); - assert(PrgA.get_build_options() == (PrgA.is_host() ? "" : BuildOpts)); + assert(PrgA.get_build_options() == BuildOpts); sycl::program PrgB{Ctx}; PrgB.compile_with_kernel_type(CompileOpts); - assert(PrgB.get_compile_options() == (PrgB.is_host() ? "" : CompileOpts)); + assert(PrgB.get_compile_options() == CompileOpts); assert(PrgB.get_link_options().empty()); - assert(PrgB.get_build_options() == (PrgB.is_host() ? "" : CompileOpts)); + assert(PrgB.get_build_options() == CompileOpts); PrgB.link(LinkOpts); - assert(PrgB.get_compile_options() == (PrgB.is_host() ? "" : CompileOpts)); - assert(PrgB.get_link_options() == (PrgB.is_host() ? "" : LinkOpts)); - assert(PrgB.get_build_options() == (PrgB.is_host() ? "" : LinkOpts)); + assert(PrgB.get_compile_options() == CompileOpts); + assert(PrgB.get_link_options() == LinkOpts); + assert(PrgB.get_build_options() == LinkOpts); } diff --git a/SYCL/DeprecatedFeatures/get_backend.cpp b/SYCL/DeprecatedFeatures/get_backend.cpp index 4ca91d044f..dfa60e1881 100644 --- a/SYCL/DeprecatedFeatures/get_backend.cpp +++ b/SYCL/DeprecatedFeatures/get_backend.cpp @@ -20,7 +20,6 @@ bool check(backend be) { case backend::ext_oneapi_level_zero: case backend::ext_oneapi_cuda: case backend::ext_oneapi_hip: - case backend::host: return true; default: return false; @@ -35,38 +34,36 @@ inline void return_fail() { int main() { for (const auto &plt : platform::get_platforms()) { - if (!plt.is_host()) { - if (check(plt.get_backend()) == false) { - return_fail(); - } + if (check(plt.get_backend()) == false) { + return_fail(); + } - context c(plt); - if (c.get_backend() != plt.get_backend()) { - return_fail(); - } + context c(plt); + if (c.get_backend() != plt.get_backend()) { + return_fail(); + } - program prog(c); - if (prog.get_backend() != plt.get_backend()) { - return_fail(); - } + program prog(c); + if (prog.get_backend() != plt.get_backend()) { + return_fail(); + } - default_selector sel; - queue q(c, sel); - if (q.get_backend() != plt.get_backend()) { - return_fail(); - } + default_selector sel; + queue q(c, sel); + if (q.get_backend() != plt.get_backend()) { + return_fail(); + } - auto device = q.get_device(); - if (device.get_backend() != plt.get_backend()) { - return_fail(); - } + auto device = q.get_device(); + if (device.get_backend() != plt.get_backend()) { + return_fail(); + } - unsigned char *HostAlloc = (unsigned char *)malloc_host(1, c); - auto e = q.memset(HostAlloc, 42, 1); - free(HostAlloc, c); - if (e.get_backend() != plt.get_backend()) { - return_fail(); - } + unsigned char *HostAlloc = (unsigned char *)malloc_host(1, c); + auto e = q.memset(HostAlloc, 42, 1); + free(HostAlloc, c); + if (e.get_backend() != plt.get_backend()) { + return_fail(); } } std::cout << "Passed" << std::endl; diff --git a/SYCL/DeprecatedFeatures/kernel-and-program-interop.cpp b/SYCL/DeprecatedFeatures/kernel-and-program-interop.cpp index 27fe4d303e..4d08c7f81d 100644 --- a/SYCL/DeprecatedFeatures/kernel-and-program-interop.cpp +++ b/SYCL/DeprecatedFeatures/kernel-and-program-interop.cpp @@ -1,7 +1,6 @@ // REQUIRES: opencl, opencl_icd // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -D__SYCL_INTERNAL_API -o %t.out %opencl_lib -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // No ACC device because Images are not supported there. @@ -28,74 +27,69 @@ int main() { int data = 0; // OpenCL interoperability kernel invocation - if (!q.is_host()) { + { + cl_int err; + sycl::context ctx = q.get_context(); + cl_context clCtx = ctx.get(); + cl_command_queue clQ = q.get(); + cl_mem clBuffer = + clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, sizeof(int), NULL, NULL); + err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, sizeof(int), &data, + 0, NULL, NULL); + assert(err == CL_SUCCESS); + clFinish(clQ); + sycl::program prog(ctx); + prog.build_with_source( + "kernel void SingleTask(global int* a) {*a+=1; }\n"); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(clBuffer); + cgh.single_task(prog.get_kernel("SingleTask")); + }); + q.wait(); + err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, sizeof(int), &data, + 0, NULL, NULL); + clReleaseCommandQueue(clQ); + clReleaseContext(clCtx); + assert(err == CL_SUCCESS); + assert(data == 1); + } + { + sycl::queue sycl_queue; + sycl::program prog(sycl_queue.get_context()); + prog.build_with_source("kernel void foo(global int* a, global int* b, " + "global int* c) {*a=*b+*c; }\n"); + int a = 13, b = 14, c = 15; { - cl_int err; - sycl::context ctx = q.get_context(); - cl_context clCtx = ctx.get(); - cl_command_queue clQ = q.get(); - cl_mem clBuffer = - clCreateBuffer(clCtx, CL_MEM_WRITE_ONLY, sizeof(int), NULL, NULL); - err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, sizeof(int), - &data, 0, NULL, NULL); - assert(err == CL_SUCCESS); - clFinish(clQ); - sycl::program prog(ctx); - prog.build_with_source( - "kernel void SingleTask(global int* a) {*a+=1; }\n"); - q.submit([&](sycl::handler &cgh) { - cgh.set_args(clBuffer); - cgh.single_task(prog.get_kernel("SingleTask")); + sycl::buffer bufa(&a, sycl::range<1>(1)); + sycl::buffer bufb(&b, sycl::range<1>(1)); + sycl::buffer bufc(&c, sycl::range<1>(1)); + sycl_queue.submit([&](sycl::handler &cgh) { + auto A = bufa.get_access(cgh); + auto B = bufb.get_access(cgh); + auto C = bufc.get_access(cgh); + cgh.set_args(A, B, C); + cgh.single_task(prog.get_kernel("foo")); }); - q.wait(); - err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, sizeof(int), &data, - 0, NULL, NULL); - clReleaseCommandQueue(clQ); - clReleaseContext(clCtx); - assert(err == CL_SUCCESS); - assert(data == 1); - } - { - sycl::queue sycl_queue; - sycl::program prog(sycl_queue.get_context()); - prog.build_with_source("kernel void foo(global int* a, global int* b, " - "global int* c) {*a=*b+*c; }\n"); - int a = 13, b = 14, c = 15; - { - sycl::buffer bufa(&a, sycl::range<1>(1)); - sycl::buffer bufb(&b, sycl::range<1>(1)); - sycl::buffer bufc(&c, sycl::range<1>(1)); - sycl_queue.submit([&](sycl::handler &cgh) { - auto A = bufa.get_access(cgh); - auto B = bufb.get_access(cgh); - auto C = bufc.get_access(cgh); - cgh.set_args(A, B, C); - cgh.single_task(prog.get_kernel("foo")); - }); - } - assert(a == b + c); } + assert(a == b + c); } { sycl::queue Queue; - if (!Queue.is_host()) { - sycl::sampler first(sycl::coordinate_normalization_mode::normalized, - sycl::addressing_mode::clamp, - sycl::filtering_mode::linear); - sycl::sampler second(sycl::coordinate_normalization_mode::unnormalized, - sycl::addressing_mode::clamp_to_edge, - sycl::filtering_mode::nearest); - sycl::program prog(Queue.get_context()); - prog.build_with_source( - "kernel void sampler_args(int a, sampler_t first, " - "int b, sampler_t second, int c) {}\n"); - sycl::kernel krn = prog.get_kernel("sampler_args"); - - Queue.submit([&](sycl::handler &cgh) { - cgh.set_args(0, first, 2, second, 3); - cgh.single_task(krn); - }); - } + sycl::sampler first(sycl::coordinate_normalization_mode::normalized, + sycl::addressing_mode::clamp, + sycl::filtering_mode::linear); + sycl::sampler second(sycl::coordinate_normalization_mode::unnormalized, + sycl::addressing_mode::clamp_to_edge, + sycl::filtering_mode::nearest); + sycl::program prog(Queue.get_context()); + prog.build_with_source("kernel void sampler_args(int a, sampler_t first, " + "int b, sampler_t second, int c) {}\n"); + sycl::kernel krn = prog.get_kernel("sampler_args"); + + Queue.submit([&](sycl::handler &cgh) { + cgh.set_args(0, first, 2, second, 3); + cgh.single_task(krn); + }); } } // Parallel for with range @@ -104,45 +98,42 @@ int main() { std::vector dataVec(10); std::iota(dataVec.begin(), dataVec.end(), 0); - if (!q.is_host()) { - cl_int err; - { - sycl::context ctx = q.get_context(); - cl_context clCtx = ctx.get(); - cl_command_queue clQ = q.get(); - cl_mem clBuffer = clCreateBuffer( - clCtx, CL_MEM_WRITE_ONLY, sizeof(int) * dataVec.size(), NULL, NULL); - err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, - sizeof(int) * dataVec.size(), dataVec.data(), - 0, NULL, NULL); - assert(err == CL_SUCCESS); - - sycl::program prog(ctx); - prog.build_with_source( - "kernel void ParallelFor(__global int* a, int v, __local int *l) " - "{ size_t index = get_global_id(0); l[index] = a[index];" - " l[index] += v; a[index] = l[index]; }\n"); - - q.submit([&](sycl::handler &cgh) { - const int value = 1; - auto local_acc = - sycl::accessor(sycl::range<1>(10), - cgh); - cgh.set_args(clBuffer, value, local_acc); - cgh.parallel_for(sycl::range<1>(10), prog.get_kernel("ParallelFor")); - }); - - q.wait(); - err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, - sizeof(int) * dataVec.size(), dataVec.data(), - 0, NULL, NULL); - clReleaseCommandQueue(clQ); - clReleaseContext(clCtx); - assert(err == CL_SUCCESS); - for (size_t i = 0; i < dataVec.size(); ++i) { - assert(dataVec[i] == i + 1); - } + cl_int err; + { + sycl::context ctx = q.get_context(); + cl_context clCtx = ctx.get(); + cl_command_queue clQ = q.get(); + cl_mem clBuffer = clCreateBuffer( + clCtx, CL_MEM_WRITE_ONLY, sizeof(int) * dataVec.size(), NULL, NULL); + err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, + sizeof(int) * dataVec.size(), dataVec.data(), + 0, NULL, NULL); + assert(err == CL_SUCCESS); + + sycl::program prog(ctx); + prog.build_with_source( + "kernel void ParallelFor(__global int* a, int v, __local int *l) " + "{ size_t index = get_global_id(0); l[index] = a[index];" + " l[index] += v; a[index] = l[index]; }\n"); + + q.submit([&](sycl::handler &cgh) { + const int value = 1; + auto local_acc = sycl::accessor( + sycl::range<1>(10), cgh); + cgh.set_args(clBuffer, value, local_acc); + cgh.parallel_for(sycl::range<1>(10), prog.get_kernel("ParallelFor")); + }); + + q.wait(); + err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, + sizeof(int) * dataVec.size(), dataVec.data(), 0, + NULL, NULL); + clReleaseCommandQueue(clQ); + clReleaseContext(clCtx); + assert(err == CL_SUCCESS); + for (size_t i = 0; i < dataVec.size(); ++i) { + assert(dataVec[i] == i + 1); } } } @@ -153,50 +144,48 @@ int main() { std::vector dataVec(10); std::iota(dataVec.begin(), dataVec.end(), 0); - if (!q.is_host()) { - cl_int err; - { - sycl::context ctx = q.get_context(); - cl_context clCtx = ctx.get(); - cl_command_queue clQ = q.get(); - cl_mem clBuffer = clCreateBuffer( - clCtx, CL_MEM_WRITE_ONLY, sizeof(int) * dataVec.size(), NULL, NULL); - err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, - sizeof(int) * dataVec.size(), dataVec.data(), - 0, NULL, NULL); - assert(err == CL_SUCCESS); - - sycl::program prog(ctx); - prog.build_with_source( - "kernel void ParallelForND( local int* l,global int* a)" - "{ size_t idx = get_global_id(0);" - " int pos = idx & 1;" - " int opp = pos ^ 1;" - " l[pos] = a[get_global_id(0)];" - " barrier(CLK_LOCAL_MEM_FENCE);" - " a[idx]=l[opp]; }"); - - // TODO is there no way to set local memory size via interoperability? - sycl::kernel krn = prog.get_kernel("ParallelForND"); - clSetKernelArg(krn.get(), 0, sizeof(int) * 2, NULL); - - q.submit([&](sycl::handler &cgh) { - cgh.set_arg(1, clBuffer); - cgh.parallel_for( - sycl::nd_range<1>(sycl::range<1>(10), sycl::range<1>(2)), krn); - }); - - q.wait(); - err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, - sizeof(int) * dataVec.size(), dataVec.data(), - 0, NULL, NULL); - clReleaseCommandQueue(clQ); - clReleaseContext(clCtx); - assert(err == CL_SUCCESS); - } - for (size_t i = 0; i < dataVec.size(); ++i) { - assert(dataVec[i] == (i ^ 1)); - } + cl_int err; + { + sycl::context ctx = q.get_context(); + cl_context clCtx = ctx.get(); + cl_command_queue clQ = q.get(); + cl_mem clBuffer = clCreateBuffer( + clCtx, CL_MEM_WRITE_ONLY, sizeof(int) * dataVec.size(), NULL, NULL); + err = clEnqueueWriteBuffer(clQ, clBuffer, CL_TRUE, 0, + sizeof(int) * dataVec.size(), dataVec.data(), + 0, NULL, NULL); + assert(err == CL_SUCCESS); + + sycl::program prog(ctx); + prog.build_with_source( + "kernel void ParallelForND( local int* l,global int* a)" + "{ size_t idx = get_global_id(0);" + " int pos = idx & 1;" + " int opp = pos ^ 1;" + " l[pos] = a[get_global_id(0)];" + " barrier(CLK_LOCAL_MEM_FENCE);" + " a[idx]=l[opp]; }"); + + // TODO is there no way to set local memory size via interoperability? + sycl::kernel krn = prog.get_kernel("ParallelForND"); + clSetKernelArg(krn.get(), 0, sizeof(int) * 2, NULL); + + q.submit([&](sycl::handler &cgh) { + cgh.set_arg(1, clBuffer); + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(10), sycl::range<1>(2)), krn); + }); + + q.wait(); + err = clEnqueueReadBuffer(clQ, clBuffer, CL_TRUE, 0, + sizeof(int) * dataVec.size(), dataVec.data(), 0, + NULL, NULL); + clReleaseCommandQueue(clQ); + clReleaseContext(clCtx); + assert(err == CL_SUCCESS); + } + for (size_t i = 0; i < dataVec.size(); ++i) { + assert(dataVec[i] == (i ^ 1)); } } } diff --git a/SYCL/DeprecatedFeatures/kernel-and-program.cpp b/SYCL/DeprecatedFeatures/kernel-and-program.cpp index d3c55404ab..59c6b6a1bc 100644 --- a/SYCL/DeprecatedFeatures/kernel-and-program.cpp +++ b/SYCL/DeprecatedFeatures/kernel-and-program.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %HOST_RUN_PLACEHOLDER %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out @@ -44,13 +43,11 @@ int main() { auto acc = buf.get_access(cgh); cgh.single_task(krn, [=]() { acc[0] = acc[0] + 1; }); }); - if (!q.is_host()) { - const std::string integrationHeaderKernelName = - sycl::detail::KernelInfo::getName(); - const std::string clKerneName = - krn.get_info(); - assert(integrationHeaderKernelName == clKerneName); - } + const std::string integrationHeaderKernelName = + sycl::detail::KernelInfo::getName(); + const std::string clKerneName = + krn.get_info(); + assert(integrationHeaderKernelName == clKerneName); } assert(data == 1); } @@ -94,45 +91,42 @@ int main() { std::iota(dataVec.begin(), dataVec.end(), 0); // Precompiled kernel invocation - // TODO run on host as well once local barrier is supported - if (!q.is_host()) { - { - sycl::range<1> numOfItems(dataVec.size()); - sycl::range<1> localRange(2); - sycl::buffer buf(dataVec.data(), numOfItems); - sycl::program prg(q.get_context()); - assert(prg.get_state() == sycl::program_state::none); - prg.build_with_kernel_type(); - assert(prg.get_state() == sycl::program_state::linked); - assert(prg.has_kernel()); - sycl::kernel krn = prg.get_kernel(); - assert(krn.get_context() == q.get_context()); - assert(krn.get_program() == prg); + { + sycl::range<1> numOfItems(dataVec.size()); + sycl::range<1> localRange(2); + sycl::buffer buf(dataVec.data(), numOfItems); + sycl::program prg(q.get_context()); + assert(prg.get_state() == sycl::program_state::none); + prg.build_with_kernel_type(); + assert(prg.get_state() == sycl::program_state::linked); + assert(prg.has_kernel()); + sycl::kernel krn = prg.get_kernel(); + assert(krn.get_context() == q.get_context()); + assert(krn.get_program() == prg); - q.submit([&](sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - sycl::accessor - localAcc(localRange, cgh); + q.submit([&](sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + sycl::accessor + localAcc(localRange, cgh); - cgh.parallel_for( - krn, sycl::nd_range<1>(numOfItems, localRange), - [=](sycl::nd_item<1> item) { - size_t idx = item.get_global_linear_id(); - int pos = idx & 1; - int opp = pos ^ 1; - localAcc[pos] = acc[item.get_global_linear_id()]; + cgh.parallel_for( + krn, sycl::nd_range<1>(numOfItems, localRange), + [=](sycl::nd_item<1> item) { + size_t idx = item.get_global_linear_id(); + int pos = idx & 1; + int opp = pos ^ 1; + localAcc[pos] = acc[item.get_global_linear_id()]; - item.barrier(sycl::access::fence_space::local_space); + item.barrier(sycl::access::fence_space::local_space); - acc[idx] = localAcc[opp]; - }); - }); - } - q.wait(); - for (size_t i = 0; i < dataVec.size(); ++i) { - assert(dataVec[i] == (i ^ 1)); - } + acc[idx] = localAcc[opp]; + }); + }); + } + q.wait(); + for (size_t i = 0; i < dataVec.size(); ++i) { + assert(dataVec[i] == (i ^ 1)); } } } diff --git a/SYCL/DeprecatedFeatures/kernel_interop.cpp b/SYCL/DeprecatedFeatures/kernel_interop.cpp index 9c3bcfd923..0da7668ac1 100644 --- a/SYCL/DeprecatedFeatures/kernel_interop.cpp +++ b/SYCL/DeprecatedFeatures/kernel_interop.cpp @@ -26,9 +26,6 @@ using namespace sycl; int main() { queue Queue; - if (Queue.is_host()) - return 0; - context Context = Queue.get_context(); cl_context ClContext = Context.get(); diff --git a/SYCL/DeprecatedFeatures/opencl-interop.cpp b/SYCL/DeprecatedFeatures/opencl-interop.cpp index d89cb9f9f7..3df0513441 100644 --- a/SYCL/DeprecatedFeatures/opencl-interop.cpp +++ b/SYCL/DeprecatedFeatures/opencl-interop.cpp @@ -25,10 +25,6 @@ using namespace sycl; int main(int argc, char *argv[]) { queue q; - if (q.is_host()) { - std::cout << "Skipping test\n"; - return 0; - } // Compute expected answer. const uint32_t dimA = 2; diff --git a/SYCL/DeprecatedFeatures/platform.cpp b/SYCL/DeprecatedFeatures/platform.cpp index 71372eda78..b6700dfd38 100644 --- a/SYCL/DeprecatedFeatures/platform.cpp +++ b/SYCL/DeprecatedFeatures/platform.cpp @@ -1,5 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_INTERNAL_API %s -o %t.out -// RUN: env SYCL_DEVICE_FILTER=host,%sycl_be %t.out +// RUN: env SYCL_DEVICE_FILTER=%sycl_be %t.out //==--------------- platform.cpp - SYCL platform test ----------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. @@ -18,12 +18,9 @@ int main() { int i = 1; std::vector openclPlatforms; for (const auto &plt : platform::get_platforms()) { - std::cout << "Platform " << i++ - << " is available: " << ((plt.is_host()) ? "host: " : "OpenCL: ") - << std::hex - << ((plt.is_host() || plt.get_backend() != sycl::backend::opencl) - ? nullptr - : plt.get()) + std::cout << "Platform " << i++ << " is available: OpenCL: " << std::hex + << ((plt.get_backend() != sycl::backend::opencl) ? nullptr + : plt.get()) << std::endl; } @@ -36,9 +33,7 @@ int main() { size_t hash = std::hash()(Platform); platform MovedPlatform(std::move(Platform)); assert(hash == std::hash()(MovedPlatform)); - assert(platformA.is_host() == MovedPlatform.is_host()); - if (!platformA.is_host() && - platformA.get_backend() == sycl::backend::opencl) { + if (platformA.get_backend() == sycl::backend::opencl) { assert(MovedPlatform.get() != nullptr); } } @@ -49,9 +44,7 @@ int main() { platform WillMovedPlatform(platformB); WillMovedPlatform = std::move(Platform); assert(hash == std::hash()(WillMovedPlatform)); - assert(platformA.is_host() == WillMovedPlatform.is_host()); - if (!platformA.is_host() && - platformA.get_backend() == sycl::backend::opencl) { + if (platformA.get_backend() == sycl::backend::opencl) { assert(WillMovedPlatform.get() != nullptr); } } @@ -63,7 +56,6 @@ int main() { assert(hash == std::hash()(Platform)); assert(hash == std::hash()(PlatformCopy)); assert(Platform == PlatformCopy); - assert(Platform.is_host() == PlatformCopy.is_host()); } { std::cout << "copy assignment operator" << std::endl; @@ -74,6 +66,5 @@ int main() { assert(hash == std::hash()(Platform)); assert(hash == std::hash()(WillPlatformCopy)); assert(Platform == WillPlatformCopy); - assert(Platform.is_host() == WillPlatformCopy.is_host()); } } diff --git a/SYCL/DeprecatedFeatures/program_link.cpp b/SYCL/DeprecatedFeatures/program_link.cpp index 06b7f8edae..40cf793625 100644 --- a/SYCL/DeprecatedFeatures/program_link.cpp +++ b/SYCL/DeprecatedFeatures/program_link.cpp @@ -25,9 +25,6 @@ class DUMMY { int main(void) { default_selector s; platform p(s); - if (p.is_host()) { - return 0; - } context c(p); queue Q(c, s); program prog1(c); diff --git a/SYCL/DeprecatedFeatures/queue_old_interop.cpp b/SYCL/DeprecatedFeatures/queue_old_interop.cpp index a1d520ad6e..e6f0e234ec 100644 --- a/SYCL/DeprecatedFeatures/queue_old_interop.cpp +++ b/SYCL/DeprecatedFeatures/queue_old_interop.cpp @@ -17,15 +17,13 @@ using namespace sycl; std::string get_type(const device &dev) { - return ((dev.is_host()) ? "host" - : (dev.is_gpu() ? "OpenCL.GPU" : "OpenCL.CPU")); + return dev.is_gpu() ? "OpenCL.GPU" : "OpenCL.CPU"; } void print_queue_info(const queue &q) { std::cout << "ID=" << std::hex - << ((q.get_device().is_host() || - q.get_context().get_platform().get_backend() != - sycl::backend::opencl) + << ((q.get_context().get_platform().get_backend() != + sycl::backend::opencl) ? nullptr : q.get()) << std::endl; @@ -51,9 +49,7 @@ int main() { size_t hash = std::hash()(Queue); queue MovedQueue(std::move(Queue)); assert(hash == std::hash()(MovedQueue)); - assert(deviceA.is_host() == MovedQueue.is_host()); - if (!deviceA.is_host() && - deviceA.get_platform().get_backend() == sycl::backend::opencl) { + if (deviceA.get_platform().get_backend() == sycl::backend::opencl) { assert(MovedQueue.get() != nullptr); } } @@ -64,9 +60,7 @@ int main() { queue WillMovedQueue(deviceB); WillMovedQueue = std::move(Queue); assert(hash == std::hash()(WillMovedQueue)); - assert(deviceA.is_host() == WillMovedQueue.is_host()); - if (!deviceA.is_host() && - deviceA.get_platform().get_backend() == sycl::backend::opencl) { + if (deviceA.get_platform().get_backend() == sycl::backend::opencl) { assert(WillMovedQueue.get() != nullptr); } } @@ -78,7 +72,6 @@ int main() { assert(hash == std::hash()(Queue)); assert(hash == std::hash()(QueueCopy)); assert(Queue == QueueCopy); - assert(Queue.is_host() == QueueCopy.is_host()); } { std::cout << "copy assignment operator" << std::endl; @@ -89,7 +82,6 @@ int main() { assert(hash == std::hash()(Queue)); assert(hash == std::hash()(WillQueueCopy)); assert(Queue == WillQueueCopy); - assert(Queue.is_host() == WillQueueCopy.is_host()); } { diff --git a/SYCL/DeprecatedFeatures/set_arg_interop.cpp b/SYCL/DeprecatedFeatures/set_arg_interop.cpp index e18804b682..853ead8161 100644 --- a/SYCL/DeprecatedFeatures/set_arg_interop.cpp +++ b/SYCL/DeprecatedFeatures/set_arg_interop.cpp @@ -13,111 +13,109 @@ using namespace sycl; int main() { queue Queue; - if (!Queue.is_host()) { - context Context = Queue.get_context(); - - cl_context ClContext = Context.get(); - - const size_t CountSources = 3; - const char *Sources[CountSources] = { - "kernel void foo1(global float* Array, global int* Value) { *Array = " - "42; *Value = 1; }\n", - "kernel void foo2(global float* Array) { int id = get_global_id(0); " - "Array[id] = id; }\n", - "kernel void foo3(global float* Array, local float* LocalArray) { " - "(void)LocalArray; (void)Array; }\n", - }; - - cl_int Err; - cl_program ClProgram = clCreateProgramWithSource(ClContext, CountSources, - Sources, nullptr, &Err); - assert(Err == CL_SUCCESS); - - Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr); - assert(Err == CL_SUCCESS); - - cl_kernel FirstCLKernel = clCreateKernel(ClProgram, "foo1", &Err); - assert(Err == CL_SUCCESS); - - cl_kernel SecondCLKernel = clCreateKernel(ClProgram, "foo2", &Err); - assert(Err == CL_SUCCESS); - - cl_kernel ThirdCLKernel = clCreateKernel(ClProgram, "foo3", &Err); - assert(Err == CL_SUCCESS); - - const size_t Count = 100; - float Array[Count]; - - kernel FirstKernel(FirstCLKernel, Context); - kernel SecondKernel(SecondCLKernel, Context); - kernel ThirdKernel(ThirdCLKernel, Context); - int Value; - { - buffer FirstBuffer(Array, range<1>(1)); - buffer SecondBuffer(&Value, range<1>(1)); - Queue.submit([&](handler &CGH) { - CGH.set_arg(0, FirstBuffer.get_access(CGH)); - CGH.set_arg(1, SecondBuffer.get_access(CGH)); - CGH.single_task(FirstKernel); - }); - } - Queue.wait_and_throw(); + context Context = Queue.get_context(); + + cl_context ClContext = Context.get(); + + const size_t CountSources = 3; + const char *Sources[CountSources] = { + "kernel void foo1(global float* Array, global int* Value) { *Array = " + "42; *Value = 1; }\n", + "kernel void foo2(global float* Array) { int id = get_global_id(0); " + "Array[id] = id; }\n", + "kernel void foo3(global float* Array, local float* LocalArray) { " + "(void)LocalArray; (void)Array; }\n", + }; + + cl_int Err; + cl_program ClProgram = clCreateProgramWithSource(ClContext, CountSources, + Sources, nullptr, &Err); + assert(Err == CL_SUCCESS); + + Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr); + assert(Err == CL_SUCCESS); + + cl_kernel FirstCLKernel = clCreateKernel(ClProgram, "foo1", &Err); + assert(Err == CL_SUCCESS); + + cl_kernel SecondCLKernel = clCreateKernel(ClProgram, "foo2", &Err); + assert(Err == CL_SUCCESS); + + cl_kernel ThirdCLKernel = clCreateKernel(ClProgram, "foo3", &Err); + assert(Err == CL_SUCCESS); + + const size_t Count = 100; + float Array[Count]; + + kernel FirstKernel(FirstCLKernel, Context); + kernel SecondKernel(SecondCLKernel, Context); + kernel ThirdKernel(ThirdCLKernel, Context); + int Value; + { + buffer FirstBuffer(Array, range<1>(1)); + buffer SecondBuffer(&Value, range<1>(1)); + Queue.submit([&](handler &CGH) { + CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.set_arg(1, SecondBuffer.get_access(CGH)); + CGH.single_task(FirstKernel); + }); + } + Queue.wait_and_throw(); + + assert(Array[0] == 42); + assert(Value == 1); + + { + buffer FirstBuffer(Array, range<1>(Count)); + Queue.submit([&](handler &CGH) { + auto Acc = FirstBuffer.get_access(CGH); + CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.parallel_for(range<1>{Count}, SecondKernel); + }); + } + Queue.wait_and_throw(); + + for (size_t I = 0; I < Count; ++I) { + assert(Array[I] == I); + } - assert(Array[0] == 42); - assert(Value == 1); + { + auto dev = Queue.get_device(); + auto ctxt = Queue.get_context(); + if (dev.get_info()) { + float *data = + static_cast(malloc_shared(Count * sizeof(float), dev, ctxt)); - { - buffer FirstBuffer(Array, range<1>(Count)); Queue.submit([&](handler &CGH) { - auto Acc = FirstBuffer.get_access(CGH); - CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.set_arg(0, data); CGH.parallel_for(range<1>{Count}, SecondKernel); }); - } - Queue.wait_and_throw(); + Queue.wait_and_throw(); - for (size_t I = 0; I < Count; ++I) { - assert(Array[I] == I); - } - - { - auto dev = Queue.get_device(); - auto ctxt = Queue.get_context(); - if (dev.get_info()) { - float *data = static_cast( - malloc_shared(Count * sizeof(float), dev, ctxt)); - - Queue.submit([&](handler &CGH) { - CGH.set_arg(0, data); - CGH.parallel_for(range<1>{Count}, SecondKernel); - }); - Queue.wait_and_throw(); - - for (size_t I = 0; I < Count; ++I) { - assert(data[I] == I); - } - free(data, ctxt); + for (size_t I = 0; I < Count; ++I) { + assert(data[I] == I); } + free(data, ctxt); } + } - { - buffer FirstBuffer(Array, range<1>(Count)); - Queue.submit([&](handler &CGH) { - auto Acc = FirstBuffer.get_access(CGH); - CGH.set_arg(0, FirstBuffer.get_access(CGH)); - CGH.set_arg(1, sycl::accessor( - sycl::range<1>(Count), CGH)); - CGH.parallel_for(range<1>{Count}, ThirdKernel); - }); - } - Queue.wait_and_throw(); - - clReleaseContext(ClContext); - clReleaseKernel(FirstCLKernel); - clReleaseKernel(SecondCLKernel); - clReleaseKernel(ThirdCLKernel); - clReleaseProgram(ClProgram); + { + buffer FirstBuffer(Array, range<1>(Count)); + Queue.submit([&](handler &CGH) { + auto Acc = FirstBuffer.get_access(CGH); + CGH.set_arg(0, FirstBuffer.get_access(CGH)); + CGH.set_arg(1, sycl::accessor( + sycl::range<1>(Count), CGH)); + CGH.parallel_for(range<1>{Count}, ThirdKernel); + }); } + Queue.wait_and_throw(); + + clReleaseContext(ClContext); + clReleaseKernel(FirstCLKernel); + clReleaseKernel(SecondCLKernel); + clReleaseKernel(ThirdCLKernel); + clReleaseProgram(ClProgram); return 0; }