From 13ef6ca1dc1604ab1552e2a7b3a106844ea2f664 Mon Sep 17 00:00:00 2001 From: Kamil Felskowski Date: Mon, 24 Nov 2025 11:49:58 +0100 Subject: [PATCH 1/3] Multiply chance of reproducing CUDA issue --- llvm/utils/lit/lit/cl_arguments.py | 2 +- .../KernelCompiler/sycl_device_globals1.cpp | 173 ++++++++++++++++++ .../KernelCompiler/sycl_device_globals2.cpp | 173 ++++++++++++++++++ .../KernelCompiler/sycl_device_globals3.cpp | 173 ++++++++++++++++++ .../KernelCompiler/sycl_device_globals4.cpp | 173 ++++++++++++++++++ 5 files changed, 693 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp create mode 100644 sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp create mode 100644 sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp create mode 100644 sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp diff --git a/llvm/utils/lit/lit/cl_arguments.py b/llvm/utils/lit/lit/cl_arguments.py index 8238bc42395af..8ab328a5bca9f 100644 --- a/llvm/utils/lit/lit/cl_arguments.py +++ b/llvm/utils/lit/lit/cl_arguments.py @@ -293,7 +293,7 @@ def parse_args(): metavar="REGEX", type=_case_insensitive_regex, help="Only run tests with paths matching the given regular expression", - default=os.environ.get("LIT_FILTER", ".*"), + default="sycl_device_globals", ) selection_group.add_argument( "--filter-out", diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp new file mode 100644 index 0000000000000..5354f37659e03 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp @@ -0,0 +1,173 @@ +//==--- sycl_device_globals.cpp --- kernel_compiler extension tests --------==// +// +// 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: sycl-jit, aspect-usm_device_allocations + +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 + +// RUN: %{build} -o %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{l0_leak_check} %{run} %t.out + +// XFAIL: target-native_cpu +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142 + +#include +#include +#include + +auto constexpr DGSource = R"===( +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +syclex::device_global DG; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_dg_adder(int val) { + DG += val; +} + +syclex::device_global DG_DIS; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_swap(int64_t *val) { + int64_t tmp = DG_DIS; + DG_DIS = *val; + *val = tmp; +} + +)==="; + +int test_device_global() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support `sycl` source kernel " + "bundle extension: " + << d.get_info() << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + + + exe_kb kbExe1 = syclex::build(kbSrc); + auto addK = kbExe1.ext_oneapi_get_kernel("ff_dg_adder"); + // Check presence of device globals. + assert(kbExe1.ext_oneapi_has_device_global("DG")); + // Querying a non-existing device global shall not crash. + assert(!kbExe1.ext_oneapi_has_device_global("bogus_DG")); + + void *dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + size_t dgSize = kbExe1.ext_oneapi_get_device_global_size("DG"); + assert(dgSize == 4); + + int32_t val; + auto checkVal = [&](int32_t expected) { + val = -1; + q.memcpy(&val, dgAddr, dgSize).wait(); + std::cout << "test" << '\n'; + std::cerr << "val: " << val << " == " << expected << '\n'; + assert(val == expected); + }; + + // Device globals are zero-initialized. + checkVal(0); + + // Set the DG. + val = 123; + q.memcpy(dgAddr, &val, dgSize).wait(); + checkVal(123); + + // Run a kernel using it. + val = -17; + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, val); + CGH.single_task(addK); + }); + q.wait(); + checkVal(123 - 17); + + // Test that each bundle has its distinct set of globals. + exe_kb kbExe2 = syclex::build(kbSrc); + dgAddr = kbExe2.ext_oneapi_get_device_global_address("DG", d); + checkVal(0); + + dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + checkVal(123 - 17); + + // Test global with `device_image_scope`. We currently cannot read/write these + // from the host, but they should work device-only. + auto swapK = kbExe2.ext_oneapi_get_kernel("ff_swap"); + int64_t *valBuf = sycl::malloc_shared(1, q); + *valBuf = -1; + auto doSwap = [&]() { + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, valBuf); + CGH.single_task(swapK); + }); + q.wait(); + }; + + doSwap(); + assert(*valBuf == 0); + doSwap(); + assert(*valBuf == -1); + + sycl::free(valBuf, q); + return 0; +} + +int test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + exe_kb kbExe = syclex::build(kbSrc); + + try { + kbExe.ext_oneapi_get_device_global_address("DG_DIS", d); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "Cannot query USM pointer for device global with " + "'device_image_scope' property") != std::string::npos); + } + return 0; +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + + return test_device_global() || test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp new file mode 100644 index 0000000000000..5354f37659e03 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp @@ -0,0 +1,173 @@ +//==--- sycl_device_globals.cpp --- kernel_compiler extension tests --------==// +// +// 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: sycl-jit, aspect-usm_device_allocations + +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 + +// RUN: %{build} -o %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{l0_leak_check} %{run} %t.out + +// XFAIL: target-native_cpu +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142 + +#include +#include +#include + +auto constexpr DGSource = R"===( +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +syclex::device_global DG; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_dg_adder(int val) { + DG += val; +} + +syclex::device_global DG_DIS; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_swap(int64_t *val) { + int64_t tmp = DG_DIS; + DG_DIS = *val; + *val = tmp; +} + +)==="; + +int test_device_global() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support `sycl` source kernel " + "bundle extension: " + << d.get_info() << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + + + exe_kb kbExe1 = syclex::build(kbSrc); + auto addK = kbExe1.ext_oneapi_get_kernel("ff_dg_adder"); + // Check presence of device globals. + assert(kbExe1.ext_oneapi_has_device_global("DG")); + // Querying a non-existing device global shall not crash. + assert(!kbExe1.ext_oneapi_has_device_global("bogus_DG")); + + void *dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + size_t dgSize = kbExe1.ext_oneapi_get_device_global_size("DG"); + assert(dgSize == 4); + + int32_t val; + auto checkVal = [&](int32_t expected) { + val = -1; + q.memcpy(&val, dgAddr, dgSize).wait(); + std::cout << "test" << '\n'; + std::cerr << "val: " << val << " == " << expected << '\n'; + assert(val == expected); + }; + + // Device globals are zero-initialized. + checkVal(0); + + // Set the DG. + val = 123; + q.memcpy(dgAddr, &val, dgSize).wait(); + checkVal(123); + + // Run a kernel using it. + val = -17; + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, val); + CGH.single_task(addK); + }); + q.wait(); + checkVal(123 - 17); + + // Test that each bundle has its distinct set of globals. + exe_kb kbExe2 = syclex::build(kbSrc); + dgAddr = kbExe2.ext_oneapi_get_device_global_address("DG", d); + checkVal(0); + + dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + checkVal(123 - 17); + + // Test global with `device_image_scope`. We currently cannot read/write these + // from the host, but they should work device-only. + auto swapK = kbExe2.ext_oneapi_get_kernel("ff_swap"); + int64_t *valBuf = sycl::malloc_shared(1, q); + *valBuf = -1; + auto doSwap = [&]() { + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, valBuf); + CGH.single_task(swapK); + }); + q.wait(); + }; + + doSwap(); + assert(*valBuf == 0); + doSwap(); + assert(*valBuf == -1); + + sycl::free(valBuf, q); + return 0; +} + +int test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + exe_kb kbExe = syclex::build(kbSrc); + + try { + kbExe.ext_oneapi_get_device_global_address("DG_DIS", d); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "Cannot query USM pointer for device global with " + "'device_image_scope' property") != std::string::npos); + } + return 0; +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + + return test_device_global() || test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp new file mode 100644 index 0000000000000..5354f37659e03 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp @@ -0,0 +1,173 @@ +//==--- sycl_device_globals.cpp --- kernel_compiler extension tests --------==// +// +// 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: sycl-jit, aspect-usm_device_allocations + +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 + +// RUN: %{build} -o %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{l0_leak_check} %{run} %t.out + +// XFAIL: target-native_cpu +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142 + +#include +#include +#include + +auto constexpr DGSource = R"===( +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +syclex::device_global DG; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_dg_adder(int val) { + DG += val; +} + +syclex::device_global DG_DIS; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_swap(int64_t *val) { + int64_t tmp = DG_DIS; + DG_DIS = *val; + *val = tmp; +} + +)==="; + +int test_device_global() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support `sycl` source kernel " + "bundle extension: " + << d.get_info() << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + + + exe_kb kbExe1 = syclex::build(kbSrc); + auto addK = kbExe1.ext_oneapi_get_kernel("ff_dg_adder"); + // Check presence of device globals. + assert(kbExe1.ext_oneapi_has_device_global("DG")); + // Querying a non-existing device global shall not crash. + assert(!kbExe1.ext_oneapi_has_device_global("bogus_DG")); + + void *dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + size_t dgSize = kbExe1.ext_oneapi_get_device_global_size("DG"); + assert(dgSize == 4); + + int32_t val; + auto checkVal = [&](int32_t expected) { + val = -1; + q.memcpy(&val, dgAddr, dgSize).wait(); + std::cout << "test" << '\n'; + std::cerr << "val: " << val << " == " << expected << '\n'; + assert(val == expected); + }; + + // Device globals are zero-initialized. + checkVal(0); + + // Set the DG. + val = 123; + q.memcpy(dgAddr, &val, dgSize).wait(); + checkVal(123); + + // Run a kernel using it. + val = -17; + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, val); + CGH.single_task(addK); + }); + q.wait(); + checkVal(123 - 17); + + // Test that each bundle has its distinct set of globals. + exe_kb kbExe2 = syclex::build(kbSrc); + dgAddr = kbExe2.ext_oneapi_get_device_global_address("DG", d); + checkVal(0); + + dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + checkVal(123 - 17); + + // Test global with `device_image_scope`. We currently cannot read/write these + // from the host, but they should work device-only. + auto swapK = kbExe2.ext_oneapi_get_kernel("ff_swap"); + int64_t *valBuf = sycl::malloc_shared(1, q); + *valBuf = -1; + auto doSwap = [&]() { + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, valBuf); + CGH.single_task(swapK); + }); + q.wait(); + }; + + doSwap(); + assert(*valBuf == 0); + doSwap(); + assert(*valBuf == -1); + + sycl::free(valBuf, q); + return 0; +} + +int test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + exe_kb kbExe = syclex::build(kbSrc); + + try { + kbExe.ext_oneapi_get_device_global_address("DG_DIS", d); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "Cannot query USM pointer for device global with " + "'device_image_scope' property") != std::string::npos); + } + return 0; +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + + return test_device_global() || test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp new file mode 100644 index 0000000000000..5354f37659e03 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp @@ -0,0 +1,173 @@ +//==--- sycl_device_globals.cpp --- kernel_compiler extension tests --------==// +// +// 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: sycl-jit, aspect-usm_device_allocations + +// UNSUPPORTED: opencl && gpu +// UNSUPPORTED-TRACKER: GSD-4287 + +// RUN: %{build} -o %t.out +// RUN: %if hip %{ env SYCL_JIT_AMDGCN_PTX_TARGET_CPU=%{amd_arch} %} %{l0_leak_check} %{run} %t.out + +// XFAIL: target-native_cpu +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142 + +#include +#include +#include + +auto constexpr DGSource = R"===( +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +syclex::device_global DG; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_dg_adder(int val) { + DG += val; +} + +syclex::device_global DG_DIS; + +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (syclex::single_task_kernel)) void ff_swap(int64_t *val) { + int64_t tmp = DG_DIS; + DG_DIS = *val; + *val = tmp; +} + +)==="; + +int test_device_global() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this device does not support `sycl` source kernel " + "bundle extension: " + << d.get_info() << std::endl; + return -1; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + + + exe_kb kbExe1 = syclex::build(kbSrc); + auto addK = kbExe1.ext_oneapi_get_kernel("ff_dg_adder"); + // Check presence of device globals. + assert(kbExe1.ext_oneapi_has_device_global("DG")); + // Querying a non-existing device global shall not crash. + assert(!kbExe1.ext_oneapi_has_device_global("bogus_DG")); + + void *dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + size_t dgSize = kbExe1.ext_oneapi_get_device_global_size("DG"); + assert(dgSize == 4); + + int32_t val; + auto checkVal = [&](int32_t expected) { + val = -1; + q.memcpy(&val, dgAddr, dgSize).wait(); + std::cout << "test" << '\n'; + std::cerr << "val: " << val << " == " << expected << '\n'; + assert(val == expected); + }; + + // Device globals are zero-initialized. + checkVal(0); + + // Set the DG. + val = 123; + q.memcpy(dgAddr, &val, dgSize).wait(); + checkVal(123); + + // Run a kernel using it. + val = -17; + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, val); + CGH.single_task(addK); + }); + q.wait(); + checkVal(123 - 17); + + // Test that each bundle has its distinct set of globals. + exe_kb kbExe2 = syclex::build(kbSrc); + dgAddr = kbExe2.ext_oneapi_get_device_global_address("DG", d); + checkVal(0); + + dgAddr = kbExe1.ext_oneapi_get_device_global_address("DG", d); + checkVal(123 - 17); + + // Test global with `device_image_scope`. We currently cannot read/write these + // from the host, but they should work device-only. + auto swapK = kbExe2.ext_oneapi_get_kernel("ff_swap"); + int64_t *valBuf = sycl::malloc_shared(1, q); + *valBuf = -1; + auto doSwap = [&]() { + q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, valBuf); + CGH.single_task(swapK); + }); + q.wait(); + }; + + doSwap(); + assert(*valBuf == 0); + doSwap(); + assert(*valBuf == -1); + + sycl::free(valBuf, q); + return 0; +} + +int test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + sycl::device d = q.get_device(); + + bool ok = d.ext_oneapi_can_build(syclex::source_language::sycl); + if (!ok) { + return 0; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, DGSource); + exe_kb kbExe = syclex::build(kbSrc); + + try { + kbExe.ext_oneapi_get_device_global_address("DG_DIS", d); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + assert(std::string(e.what()).find( + "Cannot query USM pointer for device global with " + "'device_image_scope' property") != std::string::npos); + } + return 0; +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + + return test_device_global() || test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} From 0caeeacd1d3d9a39034c74a0c8d63519ba9115d2 Mon Sep 17 00:00:00 2001 From: Kamil Felskowski Date: Mon, 24 Nov 2025 11:53:15 +0100 Subject: [PATCH 2/3] Replace cout with cerr --- sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp | 2 +- sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp | 1 - sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp | 1 - sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp | 1 - sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp | 1 - 5 files changed, 1 insertion(+), 5 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp index 7c14538dbe7e3..478f59a0c0897 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals.cpp @@ -80,7 +80,7 @@ int test_device_global() { auto checkVal = [&](int32_t expected) { val = -1; q.memcpy(&val, dgAddr, dgSize).wait(); - std::cout << "val: " << val << " == " << expected << '\n'; + std::cerr << "val: " << val << " == " << expected << '\n'; assert(val == expected); }; diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp index 5354f37659e03..2f8fb933bb9d3 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals1.cpp @@ -80,7 +80,6 @@ int test_device_global() { auto checkVal = [&](int32_t expected) { val = -1; q.memcpy(&val, dgAddr, dgSize).wait(); - std::cout << "test" << '\n'; std::cerr << "val: " << val << " == " << expected << '\n'; assert(val == expected); }; diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp index 5354f37659e03..2f8fb933bb9d3 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals2.cpp @@ -80,7 +80,6 @@ int test_device_global() { auto checkVal = [&](int32_t expected) { val = -1; q.memcpy(&val, dgAddr, dgSize).wait(); - std::cout << "test" << '\n'; std::cerr << "val: " << val << " == " << expected << '\n'; assert(val == expected); }; diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp index 5354f37659e03..2f8fb933bb9d3 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals3.cpp @@ -80,7 +80,6 @@ int test_device_global() { auto checkVal = [&](int32_t expected) { val = -1; q.memcpy(&val, dgAddr, dgSize).wait(); - std::cout << "test" << '\n'; std::cerr << "val: " << val << " == " << expected << '\n'; assert(val == expected); }; diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp index 5354f37659e03..2f8fb933bb9d3 100644 --- a/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp +++ b/sycl/test-e2e/KernelCompiler/sycl_device_globals4.cpp @@ -80,7 +80,6 @@ int test_device_global() { auto checkVal = [&](int32_t expected) { val = -1; q.memcpy(&val, dgAddr, dgSize).wait(); - std::cout << "test" << '\n'; std::cerr << "val: " << val << " == " << expected << '\n'; assert(val == expected); }; From b31eeb263952c125919900b6871c71c80503c4c0 Mon Sep 17 00:00:00 2001 From: Kamil Felskowski Date: Mon, 24 Nov 2025 12:15:47 +0100 Subject: [PATCH 3/3] Set filter only for e2e tests --- llvm/utils/lit/lit/cl_arguments.py | 2 +- sycl/test-e2e/CMakeLists.txt | 2 ++ 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/llvm/utils/lit/lit/cl_arguments.py b/llvm/utils/lit/lit/cl_arguments.py index 8ab328a5bca9f..8238bc42395af 100644 --- a/llvm/utils/lit/lit/cl_arguments.py +++ b/llvm/utils/lit/lit/cl_arguments.py @@ -293,7 +293,7 @@ def parse_args(): metavar="REGEX", type=_case_insensitive_regex, help="Only run tests with paths matching the given regular expression", - default="sycl_device_globals", + default=os.environ.get("LIT_FILTER", ".*"), ) selection_group.add_argument( "--filter-out", diff --git a/sycl/test-e2e/CMakeLists.txt b/sycl/test-e2e/CMakeLists.txt index a430d8475f193..ad8e9f5b95639 100644 --- a/sycl/test-e2e/CMakeLists.txt +++ b/sycl/test-e2e/CMakeLists.txt @@ -2,6 +2,8 @@ cmake_minimum_required(VERSION 3.20.0) message("Configuring SYCL End-to-End Tests") +set(ENV{LIT_FILTER} "sycl_device_globals") + option(SYCL_E2E_LIT_ALLOW_UNKNOWN_ARCH "Allow unknown architectures when configuring e2e tests" Off)