diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp index dc7fc89f23cd9..83af5b246683a 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,20 @@ 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) { + 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..c724f24b00b19 --- /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