From b68a96b830175fae0cf3d3158a407a2a768ffb99 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 6 Jun 2025 14:11:22 -0700 Subject: [PATCH 1/3] ensure flags passed to opencl kerel_compiler are passed all the way through Signed-off-by: Chris Perkins --- sycl/source/detail/device_image_impl.hpp | 24 ++++-- sycl/test-e2e/KernelCompiler/opencl_flags.cpp | 73 +++++++++++++++++++ 2 files changed, 89 insertions(+), 8 deletions(-) create mode 100644 sycl/test-e2e/KernelCompiler/opencl_flags.cpp diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index dc7fc89f23cd9..1c82f759aa3e3 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -747,7 +747,7 @@ class device_image_impl { if (!FetchedFromCache) UrProgram = createProgramFromSource(Devices, BuildOptions, LogPtr); - std::string XsFlags = extractXsFlags(BuildOptions); + std::string XsFlags = extractXsFlags(BuildOptions, MRTCBinInfo->MLanguage); auto Res = Adapter->call_nocheck( UrProgram, DeviceVec.size(), DeviceVec.data(), XsFlags.c_str()); if (Res == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { @@ -850,15 +850,23 @@ class device_image_impl { } static std::string - extractXsFlags(const std::vector &BuildOptions) { + extractXsFlags(const std::vector &BuildOptions, + syclex::source_language lang) { std::stringstream SS; for (sycl::detail::string_view Option : BuildOptions) { - std::string_view OptionSV{Option}; - auto Where = OptionSV.find("-Xs"); - if (Where != std::string_view::npos) { - Where += 3; - std::string_view Flags = OptionSV.substr(Where); - SS << trimXsFlags(Flags) << " "; + if (lang == syclex::source_language::sycl) { + // If the option starts with -Xs, we need to trim it. + // This is a workaround for the fact that the kernel compiler + // does not support -Xs options. + std::string_view OptionSV{Option}; + auto Where = OptionSV.find("-Xs"); + if (Where != std::string_view::npos) { + Where += 3; + std::string_view Flags = OptionSV.substr(Where); + SS << trimXsFlags(Flags) << " "; + } + } else { + SS << std::string_view{Option} << " "; } } return SS.str(); diff --git a/sycl/test-e2e/KernelCompiler/opencl_flags.cpp b/sycl/test-e2e/KernelCompiler/opencl_flags.cpp new file mode 100644 index 0000000000000..5f16538a4146d --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/opencl_flags.cpp @@ -0,0 +1,73 @@ +// REQUIRES: ocloc && (opencl || level_zero) +// UNSUPPORTED: accelerator +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// This test ensures that the Kernel Compiler build option flags +// are passed all the way through to the final binary when using OpenCL C +// source. + +#include + +#include +#include +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +const int N = 8; +const char *kernelCLSource = "__kernel void sqrt_test(__global float* A) {" + " __private int x = get_global_id(0);" + " __private int y = get_global_id(1);" + " __private int w = get_global_size(1);" + " __private int address = x * w + y;" + " A[address] = sqrt(A[address]);" + "}"; + +int main(void) { + // only one device is supported at this time, so we limit the queue and + // context to that + sycl::device d{sycl::default_selector_v}; + sycl::context ctx{d}; + sycl::queue q{ctx, d}; + + bool ok = + q.get_device().ext_oneapi_can_build(syclex::source_language::opencl); + if (!ok) { + std::cout << "Apparently this device does not support OpenCL C source " + "kernel bundle extension: " + << q.get_device().get_info() + << std::endl; + return 0; + } + + auto kb_src = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, kernelCLSource); + auto kb_exe = + syclex::build(kb_src, syclex::properties{syclex::build_options( + "-cl-fp32-correctly-rounded-divide-sqrt")}); + sycl::kernel sqrt_test = kb_exe.ext_oneapi_get_kernel("sqrt_test"); + + float *A = sycl::malloc_shared(N, q); + for (int i = 0; i < N; i++) + A[i] = static_cast(i) / N; + + q.submit([&](sycl::handler &cgh) { + cgh.set_args(A); + sycl::nd_range ndr{{N}, {1}}; + cgh.parallel_for(ndr, sqrt_test); + }).wait(); + + for (int i = 0; i < N; i++) { + float diff = A[i] - std::sqrt(static_cast(i) / N); + if (diff != 0.0) { + printf("i:%d diff:%.2e\n", i, diff); + return 1; // Error + } + } + sycl::free(A, q); + + return 0; +} \ No newline at end of file From 142cbdba90f2fd197bdcf032ada5ca2480b2a844 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 6 Jun 2025 14:28:04 -0700 Subject: [PATCH 2/3] cleanup Signed-off-by: Chris Perkins --- sycl/source/detail/device_image_impl.hpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index 1c82f759aa3e3..83af5b246683a 100644 --- a/sycl/source/detail/device_image_impl.hpp +++ b/sycl/source/detail/device_image_impl.hpp @@ -855,9 +855,6 @@ class device_image_impl { std::stringstream SS; for (sycl::detail::string_view Option : BuildOptions) { if (lang == syclex::source_language::sycl) { - // If the option starts with -Xs, we need to trim it. - // This is a workaround for the fact that the kernel compiler - // does not support -Xs options. std::string_view OptionSV{Option}; auto Where = OptionSV.find("-Xs"); if (Where != std::string_view::npos) { From aac580ed6a51411c325dda380629768356393db1 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 9 Jun 2025 10:06:11 -0700 Subject: [PATCH 3/3] reviewer feedback --- sycl/test-e2e/KernelCompiler/opencl_flags.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/opencl_flags.cpp b/sycl/test-e2e/KernelCompiler/opencl_flags.cpp index 5f16538a4146d..c724f24b00b19 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_flags.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_flags.cpp @@ -18,7 +18,7 @@ namespace syclex = sycl::ext::oneapi::experimental; const int N = 8; -const char *kernelCLSource = "__kernel void sqrt_test(__global float* A) {" +const char *KernelCLSource = "__kernel void sqrt_test(__global float* A) {" " __private int x = get_global_id(0);" " __private int y = get_global_id(1);" " __private int w = get_global_size(1);" @@ -27,8 +27,8 @@ const char *kernelCLSource = "__kernel void sqrt_test(__global float* A) {" "}"; int main(void) { - // only one device is supported at this time, so we limit the queue and - // context to that + // Only one device is supported at this time, so we limit the queue and + // context to that. sycl::device d{sycl::default_selector_v}; sycl::context ctx{d}; sycl::queue q{ctx, d}; @@ -44,7 +44,7 @@ int main(void) { } auto kb_src = syclex::create_kernel_bundle_from_source( - ctx, syclex::source_language::opencl, kernelCLSource); + ctx, syclex::source_language::opencl, KernelCLSource); auto kb_exe = syclex::build(kb_src, syclex::properties{syclex::build_options( "-cl-fp32-correctly-rounded-divide-sqrt")});