diff --git a/sycl/test/context-with-multiple-devices/context-with-multiple-devices.cpp b/sycl/test/context-with-multiple-devices/context-with-multiple-devices.cpp deleted file mode 100644 index 8409d77cdf3b2..0000000000000 --- a/sycl/test/context-with-multiple-devices/context-with-multiple-devices.cpp +++ /dev/null @@ -1,45 +0,0 @@ -// REQUIRES: cpu, accelerator, aoc - -// UNSUPPORTED: cuda, level_zero - -// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t1.out -// RUN: %CPU_RUN_PLACEHOLDER CL_CONFIG_CPU_EMULATE_DEVICES=2 %t1.out -// RUN: %CPU_RUN_PLACEHOLDER CL_CONFIG_CPU_EMULATE_DEVICES=4 %t1.out -// RUN: %clangxx -fsycl -fintelfpga -fsycl-unnamed-lambda %s -o %t2.out -// RUN: %ACC_RUN_PLACEHOLDER CL_CONFIG_CPU_EMULATE_DEVICES=2 %t2.out - -#include - -void exceptionHandler(sycl::exception_list exceptions) { - for (std::exception_ptr const &e : exceptions) { - try { - std::rethrow_exception(e); - } catch (sycl::exception const &e) { - std::cerr << "Caught asynchronous SYCL exception:\n" - << e.what() << std::endl; - } - } -} - -int main() { - auto DeviceList = sycl::device::get_devices(); - - // remove host device from the list - DeviceList.erase(std::remove_if(DeviceList.begin(), DeviceList.end(), - [](auto Device) { return Device.is_host(); }), - DeviceList.end()); - - sycl::context Context(DeviceList, &exceptionHandler); - - std::vector QueueList; - for (const auto &Device : Context.get_devices()) { - QueueList.emplace_back(Context, Device, &exceptionHandler); - } - - for (auto &Queue : QueueList) { - Queue.submit( - [&](sycl::handler &cgh) { cgh.parallel_for(100, [=](auto i) {}); }); - } - - return 0; -} diff --git a/sycl/test/function-pointers/fp-as-kernel-arg.cpp b/sycl/test/function-pointers/fp-as-kernel-arg.cpp index dfd8ba1ffb0e0..522c4c05f6516 100644 --- a/sycl/test/function-pointers/fp-as-kernel-arg.cpp +++ b/sycl/test/function-pointers/fp-as-kernel-arg.cpp @@ -1,12 +1,5 @@ -// UNSUPPORTED: windows -// UNSUPPORTED: cuda || level_zero -// CUDA does not support the function pointer as kernel argument extension. -// Hangs on level zero - // RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %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 // corresponding extension is implemented diff --git a/sycl/test/function-pointers/pass-fp-through-buffer.cpp b/sycl/test/function-pointers/pass-fp-through-buffer.cpp index 65b0c7ec6d760..abb83e25dbd4c 100644 --- a/sycl/test/function-pointers/pass-fp-through-buffer.cpp +++ b/sycl/test/function-pointers/pass-fp-through-buffer.cpp @@ -1,12 +1,5 @@ -// UNSUPPORTED: windows -// UNSUPPORTED: cuda || level_zero -// CUDA does not support the function pointer as kernel argument extension. -// Hangs on level zero - // RUN: %clangxx -Xclang -fsycl-allow-func-ptr -std=c++14 -fsycl %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %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 // corresponding extension is implemented diff --git a/sycl/test/functor/functor_inheritance.cpp b/sycl/test/functor/functor_inheritance.cpp index 436e8ce74a7d0..06cb3e9d81524 100644 --- a/sycl/test/functor/functor_inheritance.cpp +++ b/sycl/test/functor/functor_inheritance.cpp @@ -1,8 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp index 9bed0b451ca25..bea05fcc4dcd1 100644 --- a/sycl/test/functor/kernel_functor.cpp +++ b/sycl/test/functor/kernel_functor.cpp @@ -1,8 +1,6 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -o %t.out %s // RUN: cd %T // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out //==--- kernel_functor.cpp - Functors as SYCL kernel test ------------------==// // diff --git a/sycl/test/kernel-and-program/basic-program.cpp b/sycl/test/kernel-and-program/basic-program.cpp deleted file mode 100644 index 202f647ea37c0..0000000000000 --- a/sycl/test/kernel-and-program/basic-program.cpp +++ /dev/null @@ -1,93 +0,0 @@ -// XFAIL: cuda -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out - -//==--- basic-program.cpp - Basic test of program and kernel APIs ----------==// -// -// 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 - -#include - -int main() { - - // Test program and kernel APIs when building a kernel. - { - cl::sycl::queue q; - int data = 0; - { - cl::sycl::buffer buf(&data, cl::sycl::range<1>(1)); - cl::sycl::program prg(q.get_context()); - assert(prg.get_state() == cl::sycl::program_state::none); - prg.build_with_kernel_type(); - assert(prg.get_state() == cl::sycl::program_state::linked); - cl::sycl::vector_class> binaries = - prg.get_binaries(); - assert(prg.has_kernel()); - cl::sycl::kernel krn = prg.get_kernel(); - cl::sycl::string_class name = - krn.get_info(); - assert(prg.has_kernel(name)); - - q.submit([&](cl::sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - cgh.single_task(krn, [=]() { acc[0] = acc[0] + 1; }); - }); - } - assert(data == 1); - } - - // Test program and kernel APIs when compiling / linking a kernel. - { - cl::sycl::queue q; - int data = 0; - { - cl::sycl::buffer buf(&data, cl::sycl::range<1>(1)); - cl::sycl::program prg(q.get_context()); - assert(prg.get_state() == cl::sycl::program_state::none); - prg.compile_with_kernel_type(); - assert(prg.get_state() == cl::sycl::program_state::compiled); - prg.link(); - assert(prg.get_state() == cl::sycl::program_state::linked); - cl::sycl::vector_class> binaries = - prg.get_binaries(); - assert(prg.has_kernel()); - cl::sycl::kernel krn = prg.get_kernel(); - cl::sycl::string_class name = - krn.get_info(); - assert(prg.has_kernel(name)); - - q.submit([&](cl::sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - cgh.single_task(krn, - [=]() { acc[0] = acc[0] + 1; }); - }); - } - assert(data == 1); - } - - { - sycl::context context; - std::vector devices = context.get_devices(); - - sycl::program prg1(context, sycl::property_list{}); - sycl::program prg2( - context, devices, - sycl::property_list{sycl::property::buffer::use_host_ptr{}}); - if (!prg2.has_property()) { - std::cerr << "Line " << __LINE__ << ": Property was not found" - << std::endl; - return 1; - } - - sycl::property::buffer::use_host_ptr Prop = - prg2.get_property(); - } -} diff --git a/sycl/test/kernel-and-program/build-log.cpp b/sycl/test/kernel-and-program/build-log.cpp deleted file mode 100644 index cc3c3bedbe28b..0000000000000 --- a/sycl/test/kernel-and-program/build-log.cpp +++ /dev/null @@ -1,51 +0,0 @@ -// XFAIL: cuda -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out - -//==--- build-log.cpp - Test log message from faild build ----------==// -// -// 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_EXTERNAL -void symbol_that_does_not_exist(); - -void test() { - cl::sycl::queue Queue; - - // Submitting this kernel should result in a compile_program_error exception - // with a message indicating that "symbol_that_does_not_exist" is undefined. - auto Kernel = []() { -#ifdef __SYCL_DEVICE_ONLY__ - symbol_that_does_not_exist(); -#endif - }; - - std::string Msg; - int Result; - - try { - Queue.submit([&](cl::sycl::handler &CGH) { - CGH.single_task(Kernel); - }); - assert(false && "There must be compilation error"); - } catch (const cl::sycl::compile_program_error &e) { - std::string Msg(e.what()); - assert(Msg.find("symbol_that_does_not_exist") != std::string::npos); - } catch (...) { - assert(false && "There must be cl::sycl::compile_program_error"); - } -} - -int main() { - test(); - - return 0; -} diff --git a/sycl/test/kernel-and-program/cache-build-result.cpp b/sycl/test/kernel-and-program/cache-build-result.cpp deleted file mode 100644 index adf2bf2706d61..0000000000000 --- a/sycl/test/kernel-and-program/cache-build-result.cpp +++ /dev/null @@ -1,47 +0,0 @@ -// RUN: %clangxx -fsycl %s -o %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out - -#include - -SYCL_EXTERNAL -void undefined(); - -void test() { - cl::sycl::queue Queue; - - auto Kernel = []() { -#ifdef __SYCL_DEVICE_ONLY__ - undefined(); -#endif - }; - - std::string Msg; - int Result; - - for (int Idx = 0; Idx < 2; ++Idx) { - try { - Queue.submit([&](cl::sycl::handler &CGH) { - CGH.single_task(Kernel); - }); - assert(false && "There must be compilation error"); - } catch (const cl::sycl::compile_program_error &e) { - fprintf(stderr, "Exception: %s, %d\n", e.what(), e.get_cl_code()); - if (Idx == 0) { - Msg = e.what(); - Result = e.get_cl_code(); - } else { - // Exception constantly adds info on its error code in the message - assert(Msg.find_first_of(e.what()) == 0 && "Exception text differs"); - assert(Result == e.get_cl_code() && "Exception code differs"); - } - } catch (...) { - assert(false && "There must be cl::sycl::compile_program_error"); - } - } -} - -int main() { - test(); - - return 0; -} diff --git a/sycl/test/kernel-and-program/get-options.cpp b/sycl/test/kernel-and-program/get-options.cpp deleted file mode 100644 index 6bb889a4c09ec..0000000000000 --- a/sycl/test/kernel-and-program/get-options.cpp +++ /dev/null @@ -1,45 +0,0 @@ -// RUN: %clangxx -fsycl %s -o %t.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out - -#include - -#include - -// Check program::get_compile/link/build_options functions - -class KernelName; -void submitKernel() { - cl::sycl::queue q; - q.submit( - [&](cl::sycl::handler &cgh) { cgh.single_task([]() {}); }); -} - -int main() { - const cl::sycl::string_class CompileOpts{"-cl-opt-disable"}; - const cl::sycl::string_class LinkOpts{"-cl-fast-relaxed-math"}; - const cl::sycl::string_class BuildOpts{ - "-cl-opt-disable -cl-fast-relaxed-math"}; - - cl::sycl::context Ctx; - cl::sycl::program PrgA{Ctx}; - assert(PrgA.get_compile_options().empty()); - assert(PrgA.get_link_options().empty()); - assert(PrgA.get_build_options().empty()); - - 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)); - - cl::sycl::program PrgB{Ctx}; - PrgB.compile_with_kernel_type(CompileOpts); - assert(PrgB.get_compile_options() == (PrgB.is_host() ? "" : CompileOpts)); - assert(PrgB.get_link_options().empty()); - assert(PrgB.get_build_options() == (PrgB.is_host() ? "" : 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)); -} diff --git a/sycl/test/kernel-and-program/kernel-and-program-interop.cpp b/sycl/test/kernel-and-program/kernel-and-program-interop.cpp deleted file mode 100644 index 991a313712fae..0000000000000 --- a/sycl/test/kernel-and-program/kernel-and-program-interop.cpp +++ /dev/null @@ -1,205 +0,0 @@ -// REQUIRES: opencl - -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -L %opencl_libs_dir -lOpenCL -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUNx: %ACC_RUN_PLACEHOLDER %t.out - -//==--- kernel-and-program.cpp - SYCL kernel/program 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 - -#include -#include -#include -#include - -int main() { - // Single task invocation methods - { - cl::sycl::queue q; - int data = 0; - - // OpenCL interoperability kernel invocation - if (!q.is_host()) { - { - cl_int err; - cl::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); - cl::sycl::program prog(ctx); - prog.build_with_source( - "kernel void SingleTask(global int* a) {*a+=1; }\n"); - q.submit([&](cl::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); - } - { - cl::sycl::queue sycl_queue; - cl::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::sycl::buffer bufa(&a, cl::sycl::range<1>(1)); - cl::sycl::buffer bufb(&b, cl::sycl::range<1>(1)); - cl::sycl::buffer bufc(&c, cl::sycl::range<1>(1)); - sycl_queue.submit([&](cl::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); - } - } - { - cl::sycl::queue Queue; - if (!Queue.is_host()) { - cl::sycl::sampler first( - cl::sycl::coordinate_normalization_mode::normalized, - cl::sycl::addressing_mode::clamp, cl::sycl::filtering_mode::linear); - cl::sycl::sampler second( - cl::sycl::coordinate_normalization_mode::unnormalized, - cl::sycl::addressing_mode::clamp_to_edge, - cl::sycl::filtering_mode::nearest); - cl::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"); - cl::sycl::kernel krn = prog.get_kernel("sampler_args"); - - Queue.submit([&](cl::sycl::handler &cgh) { - cgh.set_args(0, first, 2, second, 3); - cgh.single_task(krn); - }); - } - } - } - // Parallel for with range - { - cl::sycl::queue q; - std::vector dataVec(10); - std::iota(dataVec.begin(), dataVec.end(), 0); - - if (!q.is_host()) { - cl_int err; - { - cl::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); - - cl::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([&](cl::sycl::handler &cgh) { - const int value = 1; - auto local_acc = - cl::sycl::accessor( - cl::sycl::range<1>(10), cgh); - cgh.set_args(clBuffer, value, local_acc); - cgh.parallel_for(cl::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); - } - } - } - } - - // Parallel for with nd_range - { - cl::sycl::queue q; - std::vector dataVec(10); - std::iota(dataVec.begin(), dataVec.end(), 0); - - if (!q.is_host()) { - cl_int err; - { - cl::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); - - cl::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? - cl::sycl::kernel krn = prog.get_kernel("ParallelForND"); - clSetKernelArg(krn.get(), 0, sizeof(int) * 2, NULL); - - q.submit([&](cl::sycl::handler &cgh) { - cgh.set_arg(1, clBuffer); - cgh.parallel_for(cl::sycl::nd_range<1>(cl::sycl::range<1>(10), - cl::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/test/kernel-and-program/kernel-and-program.cpp b/sycl/test/kernel-and-program/kernel-and-program.cpp deleted file mode 100644 index 1233ec3664ee6..0000000000000 --- a/sycl/test/kernel-and-program/kernel-and-program.cpp +++ /dev/null @@ -1,137 +0,0 @@ -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -// RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUNx: %ACC_RUN_PLACEHOLDER %t.out - -//==--- kernel-and-program.cpp - SYCL kernel/program 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 - -#include -#include -#include -#include - -int main() { - - // Single task invocation methods - { - cl::sycl::queue q; - int data = 0; - // Precompiled kernel invocation - { - cl::sycl::buffer buf(&data, cl::sycl::range<1>(1)); - cl::sycl::program prg(q.get_context()); - // Test program building - assert(prg.get_state() == cl::sycl::program_state::none); - prg.build_with_kernel_type(); - assert(prg.get_state() == cl::sycl::program_state::linked); - assert(prg.has_kernel()); - cl::sycl::kernel krn = prg.get_kernel(); - assert(krn.get_context() == q.get_context()); - assert(krn.get_program() == prg); - - q.submit([&](cl::sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - cgh.single_task(krn, [=]() { acc[0] = acc[0] + 1; }); - }); - if (!q.is_host()) { - const std::string integrationHeaderKernelName = - cl::sycl::detail::KernelInfo::getName(); - const std::string clKerneName = - krn.get_info(); - assert(integrationHeaderKernelName == clKerneName); - } - } - assert(data == 1); - - } - // Parallel for with range - { - cl::sycl::queue q; - std::vector dataVec(10); - std::iota(dataVec.begin(), dataVec.end(), 0); - // Precompiled kernel invocation - { - cl::sycl::range<1> numOfItems(dataVec.size()); - cl::sycl::buffer buf(dataVec.data(), numOfItems); - cl::sycl::program prg(q.get_context()); - assert(prg.get_state() == cl::sycl::program_state::none); - // Test compiling -> linking - prg.compile_with_kernel_type(); - assert(prg.get_state() == cl::sycl::program_state::compiled); - prg.link(); - assert(prg.get_state() == cl::sycl::program_state::linked); - assert(prg.has_kernel()); - cl::sycl::kernel krn = prg.get_kernel(); - assert(krn.get_context() == q.get_context()); - assert(krn.get_program() == prg); - - q.submit([&](cl::sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - cgh.parallel_for( - krn, numOfItems, - [=](cl::sycl::id<1> wiID) { acc[wiID] = acc[wiID] + 1; }); - }); - } - for (size_t i = 0; i < dataVec.size(); ++i) { - assert(dataVec[i] == i + 1); - } - } - - // Parallel for with nd_range - { - cl::sycl::queue q; - std::vector dataVec(10); - 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()) { - { - cl::sycl::range<1> numOfItems(dataVec.size()); - cl::sycl::range<1> localRange(2); - cl::sycl::buffer buf(dataVec.data(), numOfItems); - cl::sycl::program prg(q.get_context()); - assert(prg.get_state() == cl::sycl::program_state::none); - prg.build_with_kernel_type(); - assert(prg.get_state() == cl::sycl::program_state::linked); - assert(prg.has_kernel()); - cl::sycl::kernel krn = prg.get_kernel(); - assert(krn.get_context() == q.get_context()); - assert(krn.get_program() == prg); - - q.submit([&](cl::sycl::handler &cgh) { - auto acc = buf.get_access(cgh); - cl::sycl::accessor - localAcc(localRange, cgh); - - cgh.parallel_for( - krn, cl::sycl::nd_range<1>(numOfItems, localRange), - [=](cl::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(cl::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)); - } - } - } -} diff --git a/sycl/test/kernel_from_file/hw.cpp b/sycl/test/kernel_from_file/hw.cpp deleted file mode 100644 index b7c4e573be65c..0000000000000 --- a/sycl/test/kernel_from_file/hw.cpp +++ /dev/null @@ -1,51 +0,0 @@ -// UNSUPPORTED: cuda -// CUDA does not support SPIR-V. - -//-fsycl-targets=%sycl_triple -// RUN: %clangxx -fsycl-device-only -fno-sycl-use-bitcode -Xclang -fsycl-int-header=%t.h -c %s -o %t.spv -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -Wno-sycl-strict -// RUN: %clangxx -include %t.h -g %s -o %t.out -lsycl -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -// RUN: env SYCL_BE=%sycl_be SYCL_USE_KERNEL_SPV=%t.spv %t.out | FileCheck %s -// CHECK: Passed - -// TODO: InvalidTargetTriple: Expects spir-unknown-unknown or spir64-unknown-unknown. Actual target triple is x86_64-unknown-linux-gnu - -#include -#include - -using namespace cl::sycl; - -int main(int argc, char **argv) { - int data = 5; - - try { - queue myQueue; - buffer buf(&data, range<1>(1)); - - event e = myQueue.submit([&](handler& cgh) { - auto ptr = buf.get_access(cgh); - - cgh.single_task([=]() { - ptr[0]++; - }); - }); - e.wait_and_throw(); - - } catch (cl::sycl::exception const& e) { - std::cerr << "SYCL exception caught:\n"; - std::cerr << e.what() << "\n"; - return 2; - } - catch (...) { - std::cerr << "unknown exception caught\n"; - return 1; - } - - if (data == 6) { - std::cout << "Passed\n"; - return 0; - } else { - std::cout << "Failed: " << data << "!= 6(gold)\n"; - return 1; - } -} - diff --git a/sycl/test/array_param/array-kernel-param-nested-run.cpp b/sycl/test/kernel_param/array-kernel-param-nested-run.cpp similarity index 97% rename from sycl/test/array_param/array-kernel-param-nested-run.cpp rename to sycl/test/kernel_param/array-kernel-param-nested-run.cpp index 5d08ff010ce32..d36da50b8fca8 100755 --- a/sycl/test/array_param/array-kernel-param-nested-run.cpp +++ b/sycl/test/kernel_param/array-kernel-param-nested-run.cpp @@ -2,9 +2,6 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include diff --git a/sycl/test/array_param/array-kernel-param-run.cpp b/sycl/test/kernel_param/array-kernel-param-run.cpp similarity index 98% rename from sycl/test/array_param/array-kernel-param-run.cpp rename to sycl/test/kernel_param/array-kernel-param-run.cpp index c26ef53a602f4..f7cc3d9391b63 100755 --- a/sycl/test/array_param/array-kernel-param-run.cpp +++ b/sycl/test/kernel_param/array-kernel-param-run.cpp @@ -2,9 +2,6 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include diff --git a/sycl/test/struct_param/non-standard-layout.cpp b/sycl/test/kernel_param/non-standard-layout.cpp similarity index 87% rename from sycl/test/struct_param/non-standard-layout.cpp rename to sycl/test/kernel_param/non-standard-layout.cpp index 7e1ca43cef6b9..1acb9b0aac6dd 100644 --- a/sycl/test/struct_param/non-standard-layout.cpp +++ b/sycl/test/kernel_param/non-standard-layout.cpp @@ -1,8 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include diff --git a/sycl/test/struct_param/struct_kernel_param.cpp b/sycl/test/kernel_param/struct_kernel_param.cpp similarity index 96% rename from sycl/test/struct_param/struct_kernel_param.cpp rename to sycl/test/kernel_param/struct_kernel_param.cpp index 7162cd872c616..1a6490ef2342b 100644 --- a/sycl/test/struct_param/struct_kernel_param.cpp +++ b/sycl/test/kernel_param/struct_kernel_param.cpp @@ -1,9 +1,5 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// TODO: Uncomment once test is fixed on GPU -// RUNx: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out //==-struct_kernel_param.cpp-Checks passing structs as kernel params--------==// // diff --git a/sycl/test/union_param/union_kernel_param.cpp b/sycl/test/kernel_param/union_kernel_param.cpp similarity index 88% rename from sycl/test/union_param/union_kernel_param.cpp rename to sycl/test/kernel_param/union_kernel_param.cpp index adfa818dd1778..4b0980537e4c7 100644 --- a/sycl/test/union_param/union_kernel_param.cpp +++ b/sycl/test/kernel_param/union_kernel_param.cpp @@ -2,9 +2,6 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out #include #include diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp index 3d80ded385303..dea8dd3ee22a4 100644 --- a/sycl/test/multi_ptr/multi_ptr.cpp +++ b/sycl/test/multi_ptr/multi_ptr.cpp @@ -1,13 +1,7 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-dead-args-optimization %s -o %t.out // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// RUN: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out // RUN: %clangxx -DRESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR -fsycl -fsycl-targets=%sycl_triple -fsycl-dead-args-optimization %s -o %t1.out // RUN: env SYCL_DEVICE_TYPE=HOST %t1.out -// RUN: %CPU_RUN_PLACEHOLDER %t1.out -// RUN: %GPU_RUN_PLACEHOLDER %t1.out -// RUN: %ACC_RUN_PLACEHOLDER %t1.out //==--------------- multi_ptr.cpp - SYCL multi_ptr test --------------------==// //