From 7ee30435b26fdb995cd851af7bd7817a5a22a55e Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Tue, 12 Nov 2024 18:18:20 +0100 Subject: [PATCH 1/8] Support dynamic linking on new offloading model --- clang/lib/Driver/ToolChains/Clang.cpp | 7 + .../ClangLinkerWrapper.cpp | 2 + .../clang-linker-wrapper/LinkerWrapperOpts.td | 8 + .../NewOffloadingDriver/Inputs/a.cpp | 13 + .../NewOffloadingDriver/Inputs/a.hpp | 13 + .../NewOffloadingDriver/Inputs/b.cpp | 13 + .../NewOffloadingDriver/Inputs/b.hpp | 13 + .../NewOffloadingDriver/Inputs/c.cpp | 13 + .../NewOffloadingDriver/Inputs/c.hpp | 13 + .../NewOffloadingDriver/Inputs/d.cpp | 11 + .../NewOffloadingDriver/Inputs/d.hpp | 13 + .../NewOffloadingDriver/Inputs/wrapper.cpp | 26 ++ .../NewOffloadingDriver/Inputs/wrapper.hpp | 8 + .../NewOffloadingDriver/dynamic.cpp | 44 +++ .../free_function_kernels.cpp | 285 ++++++++++++++++++ .../NewOffloadingDriver/math_device_lib.cpp | 27 ++ .../NewOffloadingDriver/objects.cpp | 35 +++ .../singleDynamicLibrary.cpp | 24 ++ 18 files changed, 568 insertions(+) create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp create mode 100644 sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index f90ba124e5a09..6fb4f745c58f1 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11486,6 +11486,13 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasArg(options::OPT_fsycl_embed_ir)) CmdArgs.push_back(Args.MakeArgString("-sycl-embed-ir")); + if (Args.hasArg(options::OPT_fsycl_allow_device_image_dependencies)) + CmdArgs.push_back( + Args.MakeArgString("-sycl-allow-device-image-dependencies")); + if (Args.hasArg(options::OPT_fno_sycl_allow_device_image_dependencies)) + CmdArgs.push_back( + Args.MakeArgString("-no-sycl-allow-device-image-dependencies")); + // Formulate and add any offload-wrapper and AOT specific options. These // are additional options passed in via -Xsycl-target-linker and // -Xsycl-target-backend. diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index ec883c1091196..cc682f1278a51 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -676,6 +676,8 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args, if ((!Args.hasFlag(OPT_no_sycl_remove_unused_external_funcs, OPT_sycl_remove_unused_external_funcs, false) && !SYCLNativeCPU) && + !Args.hasFlag(OPT_sycl_allow_device_image_dependencies, + OPT_no_sycl_allow_device_image_dependencies, false) && !Triple.isNVPTX() && !Triple.isAMDGPU()) PostLinkArgs.push_back("-emit-only-kernels-as-entry-points"); diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index 60a13b23ba306..f6159d100115f 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -243,3 +243,11 @@ Flags<[WrapperOnlyOption]>, HelpText<"Embed LLVM IR for runtime kernel fusion"> def sycl_dump_device_code_EQ : Joined<["--", "-"], "sycl-dump-device-code=">, Flags<[WrapperOnlyOption]>, HelpText<"Path to the folder where the tool dumps SPIR-V device code. Other formats aren't dumped.">; + +// Options to enable/disable device dynamic linking. +def sycl_allow_device_image_dependencies : Flag<["--", "-"], "sycl-allow-device-image-dependencies">, + Flags<[WrapperOnlyOption, HelpHidden]>, + HelpText<"Allow dependencies between device code images">; +def no_sycl_allow_device_image_dependencies : Flag<["--", "-"], "no-sycl-allow-device-image-dependencies">, + Flags<[WrapperOnlyOption, HelpHidden]>, + HelpText<"Dno not allow dependencies between device code images (default)">; diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp new file mode 100644 index 0000000000000..34d1e3bb488f8 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp @@ -0,0 +1,13 @@ +#define A_EXPORT +#include "a.hpp" +#include "b.hpp" +#include + +A_DECLSPEC SYCL_EXTERNAL int levelA(int val) { +#ifndef __SYCL_DEVICE_ONLY__ + std::cerr << "Host symbol used" << std::endl; + val ^= 0x1234; +#endif + val = levelB(val); + return val |= (0xA << 0); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp new file mode 100644 index 0000000000000..ca9320c0e1fdd --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp @@ -0,0 +1,13 @@ +#include + +#if defined(MAKE_DLL) +#ifdef A_EXPORT +#define A_DECLSPEC __declspec(dllexport) +#else +#define A_DECLSPEC __declspec(dllimport) +#endif +#else +#define A_DECLSPEC +#endif + +A_DECLSPEC SYCL_EXTERNAL int levelA(int val); diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp new file mode 100644 index 0000000000000..5dddf5b5311d6 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp @@ -0,0 +1,13 @@ +#define B_EXPORT +#include "b.hpp" +#include "c.hpp" +#include + +B_DECLSPEC SYCL_EXTERNAL int levelB(int val) { +#ifndef __SYCL_DEVICE_ONLY__ + std::cerr << "Host symbol used" << std::endl; + val ^= 0x2345; +#endif + val = levelC(val); + return val |= (0xB << 4); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp new file mode 100644 index 0000000000000..019f1ccd19616 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp @@ -0,0 +1,13 @@ +#include + +#if defined(MAKE_DLL) +#ifdef B_EXPORT +#define B_DECLSPEC __declspec(dllexport) +#else +#define B_DECLSPEC __declspec(dllimport) +#endif +#else +#define B_DECLSPEC +#endif + +B_DECLSPEC SYCL_EXTERNAL int levelB(int val); diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp new file mode 100644 index 0000000000000..247be679882e4 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp @@ -0,0 +1,13 @@ +#define C_EXPORT +#include "c.hpp" +#include "d.hpp" +#include + +C_DECLSPEC SYCL_EXTERNAL int levelC(int val) { +#ifndef __SYCL_DEVICE_ONLY__ + std::cerr << "Host symbol used" << std::endl; + val ^= 0x3456; +#endif + val = levelD(val); + return val |= (0xC << 8); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp new file mode 100644 index 0000000000000..bc189ca5cf175 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp @@ -0,0 +1,13 @@ +#include + +#if defined(MAKE_DLL) +#ifdef C_EXPORT +#define C_DECLSPEC __declspec(dllexport) +#else +#define C_DECLSPEC __declspec(dllimport) +#endif +#else +#define C_DECLSPEC +#endif + +C_DECLSPEC SYCL_EXTERNAL int levelC(int val); diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp new file mode 100644 index 0000000000000..ca3dc79e5218e --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp @@ -0,0 +1,11 @@ +#define D_EXPORT +#include "d.hpp" +#include + +D_DECLSPEC SYCL_EXTERNAL int levelD(int val) { +#ifndef __SYCL_DEVICE_ONLY__ + std::cerr << "Host symbol used" << std::endl; + val ^= 0x4567; +#endif + return val |= (0xD << 12); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp new file mode 100644 index 0000000000000..ae865c0ad3a11 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp @@ -0,0 +1,13 @@ +#include + +#if defined(MAKE_DLL) +#ifdef D_EXPORT +#define D_DECLSPEC __declspec(dllexport) +#else +#define D_DECLSPEC __declspec(dllimport) +#endif +#else +#define D_DECLSPEC +#endif + +D_DECLSPEC SYCL_EXTERNAL int levelD(int val); diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp new file mode 100644 index 0000000000000..f2a0859b17477 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp @@ -0,0 +1,26 @@ +#include "a.hpp" +#include +#include +#define EXPORT +#include "wrapper.hpp" + +using namespace sycl; + +class ExeKernel; + +int wrapper() { + int val = 0; + { + buffer buf(&val, range<1>(1)); + queue q; + q.submit([&](handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] = levelA(acc[0]); }); + }); + } + + std::cout << "val=" << std::hex << val << "\n"; + if (val != 0xDCBA) + return (1); + return (0); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp new file mode 100644 index 0000000000000..8c5d4d8a5c123 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp @@ -0,0 +1,8 @@ +#if defined(_WIN32) +#ifdef EXPORT +__declspec(dllexport) +#else +__declspec(dllimport) +#endif +#endif + int wrapper(); diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp new file mode 100644 index 0000000000000..136a3bed85728 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp @@ -0,0 +1,44 @@ +// Test -fsycl-allow-device-image-dependencies with dynamic libraries. + +// UNSUPPORTED: cuda || hip + +// DEFINE: %{dynamic_lib_options} = -fsycl %fPIC %shared_lib -fsycl-allow-device-image-dependencies -I %S/Inputs %if windows %{-DMAKE_DLL %} +// DEFINE: %{dynamic_lib_suffix} = %if windows %{dll%} %else %{so%} + +// RUN: %clangxx --offload-new-driver %{dynamic_lib_options} %S/Inputs/d.cpp -o %T/libdevice_d.%{dynamic_lib_suffix} +// RUN: %clangxx --offload-new-driver %{dynamic_lib_options} %S/Inputs/c.cpp %if windows %{%T/libdevice_d.lib%} -o %T/libdevice_c.%{dynamic_lib_suffix} +// RUN: %clangxx --offload-new-driver %{dynamic_lib_options} %S/Inputs/b.cpp %if windows %{%T/libdevice_c.lib%} -o %T/libdevice_b.%{dynamic_lib_suffix} +// RUN: %clangxx --offload-new-driver %{dynamic_lib_options} %S/Inputs/a.cpp %if windows %{%T/libdevice_b.lib%} -o %T/libdevice_a.%{dynamic_lib_suffix} + +// RUN: %{build} --offload-new-driver -fsycl-allow-device-image-dependencies -I %S/Inputs -o %t.out \ +// RUN: %if windows \ +// RUN: %{%T/libdevice_a.lib%} \ +// RUN: %else \ +// RUN: %{-L%T -ldevice_a -ldevice_b -ldevice_c -ldevice_d -Wl,-rpath=%T%} + +// RUN: %{run} %t.out + +#include "a.hpp" +#include +#include + +using namespace sycl; + +class ExeKernel; + +int main() { + int val = 0; + { + buffer buf(&val, range<1>(1)); + queue q; + q.submit([&](handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] = levelA(acc[0]); }); + }); + } + + std::cout << "val=" << std::hex << val << "\n"; + if (val != 0xDCBA) + return (1); + return (0); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp new file mode 100644 index 0000000000000..9649368328236 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp @@ -0,0 +1,285 @@ +// Ensure -fsycl-allow-device-dependencies can work with free function kernels. + +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out --offload-new-driver -fsycl-allow-device-image-dependencies +// RUN: %{run} %t.out + +// The name mangling for free function kernels currently does not work with PTX. +// UNSUPPORTED: cuda + +// XFAIL: hip_amd +// XFAIL-TRACKER: https://github.com/intel/llvm/issues/15742 + +#include +#include +#include +#include + +using namespace sycl; + +void printUSM(int *usmPtr, int size) { + std::cout << "usmPtr[] = {"; + for (int i = 0; i < size; i++) { + std::cout << usmPtr[i] << ", "; + } + std::cout << "}\n"; +} + +bool checkUSM(int *usmPtr, int size, int *Result) { + bool Pass = true; + for (int i = 0; i < size; i++) { + if (usmPtr[i] != Result[i]) { + Pass = false; + break; + } + } + if (Pass) + return true; + + std::cout << "Expected = {"; + for (int i = 0; i < size; i++) { + std::cout << Result[i] << ", "; + } + std::cout << "}\n"; + std::cout << "Result = {"; + for (int i = 0; i < size; i++) { + std::cout << usmPtr[i] << ", "; + } + std::cout << "}\n"; + return false; +} + +extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (ext::oneapi::experimental::single_task_kernel)) void ff_0(int *ptr, + int start, + int end) { + for (int i = start; i <= end; i++) + ptr[i] = start + end; +} + +bool test_0(queue Queue) { + constexpr int Range = 10; + int *usmPtr = malloc_shared(Range, Queue); + int start = 3; + int end = 5; + int Result[Range] = {0, 0, 0, 8, 8, 8, 0, 0, 0, 0}; + range<1> R1{Range}; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.single_task([=]() { + for (int i = start; i <= end; i++) + usmPtr[i] = start + end; + }); + }); + Queue.wait(); + bool PassA = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 0a: " << (PassA ? "PASS" : "FAIL") << std::endl; + + bool PassB = false; + // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment +#ifndef __SYCL_DEVICE_ONLY__ + kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id(); + kernel Kernel = Bundle.get_kernel(Kernel_id); + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.set_arg(0, usmPtr); + Handler.set_arg(1, start); + Handler.set_arg(2, end); + Handler.single_task(Kernel); + }); + Queue.wait(); + PassB = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 0b: " << (PassB ? "PASS" : "FAIL") << std::endl; + + free(usmPtr, Queue); +#endif + return PassA && PassB; +} + +// Overloaded free function definition. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_1(int *ptr, int start, int end) { + nd_item<1> Item = ext::oneapi::this_work_item::get_nd_item<1>(); + id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = GId.get(0) + start + end; +} + +bool test_1(queue Queue) { + constexpr int Range = 10; + int *usmPtr = malloc_shared(Range, Queue); + int start = 3; + int Result[Range] = {13, 14, 15, 16, 17, 18, 19, 20, 21, 22}; + nd_range<1> R1{{Range}, {1}}; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.parallel_for(R1, [=](nd_item<1> Item) { + id<1> GId = Item.get_global_id(); + usmPtr[GId.get(0)] = GId.get(0) + start + Range; + }); + }); + Queue.wait(); + bool PassA = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 1a: " << (PassA ? "PASS" : "FAIL") << std::endl; + + bool PassB = false; + // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment +#ifndef __SYCL_DEVICE_ONLY__ + kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<( + void (*)(int *, int, int))ff_1>(); + kernel Kernel = Bundle.get_kernel(Kernel_id); + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.set_arg(0, usmPtr); + Handler.set_arg(1, start); + Handler.set_arg(2, Range); + Handler.parallel_for(R1, Kernel); + }); + Queue.wait(); + PassB = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 1b: " << (PassB ? "PASS" : "FAIL") << std::endl; + + free(usmPtr, Queue); +#endif + return PassA && PassB; +} + +// Overloaded free function definition. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (ext::oneapi::experimental::nd_range_kernel<2>)) +void ff_1(int *ptr, int start) { + int(&ptr2D)[4][4] = *reinterpret_cast(ptr); + nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>(); + id<2> GId = Item.get_global_id(); + id<2> LId = Item.get_local_id(); + ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + start; +} + +bool test_2(queue Queue) { + constexpr int Range = 16; + int *usmPtr = malloc_shared(Range, Queue); + int value = 55; + int Result[Range] = {55, 56, 55, 56, 56, 57, 56, 57, + 55, 56, 55, 56, 56, 57, 56, 57}; + nd_range<2> R2{range<2>{4, 4}, range<2>{2, 2}}; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.parallel_for(R2, [=](nd_item<2> Item) { + int(&ptr2D)[4][4] = *reinterpret_cast(usmPtr); + id<2> GId = Item.get_global_id(); + id<2> LId = Item.get_local_id(); + ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + value; + }); + }); + Queue.wait(); + bool PassA = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 2a: " << (PassA ? "PASS" : "FAIL") << std::endl; + + bool PassB = false; + // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment +#ifndef __SYCL_DEVICE_ONLY__ + kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + kernel_id Kernel_id = + ext::oneapi::experimental::get_kernel_id<(void (*)(int *, int))ff_1>(); + kernel Kernel = Bundle.get_kernel(Kernel_id); + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.set_arg(0, usmPtr); + Handler.set_arg(1, value); + Handler.parallel_for(R2, Kernel); + }); + Queue.wait(); + PassB = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 2b: " << (PassB ? "PASS" : "FAIL") << std::endl; + + free(usmPtr, Queue); +#endif + return PassA && PassB; +} + +// Templated free function definition. +template +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (ext::oneapi::experimental::nd_range_kernel<2>)) +void ff_3(T *ptr, T start) { + int(&ptr2D)[4][4] = *reinterpret_cast(ptr); + nd_item<2> Item = ext::oneapi::this_work_item::get_nd_item<2>(); + id<2> GId = Item.get_global_id(); + id<2> LId = Item.get_local_id(); + ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + start; +} + +// Explicit instantiation with "int*". +template void ff_3(int *ptr, int start); + +bool test_3(queue Queue) { + constexpr int Range = 16; + int *usmPtr = malloc_shared(Range, Queue); + int value = 55; + int Result[Range] = {55, 56, 55, 56, 56, 57, 56, 57, + 55, 56, 55, 56, 56, 57, 56, 57}; + nd_range<2> R2{range<2>{4, 4}, range<2>{2, 2}}; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.parallel_for(R2, [=](nd_item<2> Item) { + int(&ptr2D)[4][4] = *reinterpret_cast(usmPtr); + id<2> GId = Item.get_global_id(); + id<2> LId = Item.get_local_id(); + ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + value; + }); + }); + Queue.wait(); + bool PassA = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 3a: " << (PassA ? "PASS" : "FAIL") << std::endl; + + bool PassB = false; + // TODO: Avoid using __SYCL_DEVICE_ONLY__ or give rationale with a comment +#ifndef __SYCL_DEVICE_ONLY__ + kernel_bundle Bundle = + get_kernel_bundle(Queue.get_context()); + kernel_id Kernel_id = ext::oneapi::experimental::get_kernel_id<( + void (*)(int *, int))ff_3>(); + kernel Kernel = Bundle.get_kernel(Kernel_id); + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](handler &Handler) { + Handler.set_arg(0, usmPtr); + Handler.set_arg(1, value); + Handler.parallel_for(R2, Kernel); + }); + Queue.wait(); + PassB = checkUSM(usmPtr, Range, Result); + // TODO: Avoid printing anything if test passes to reduce I/O. + std::cout << "Test 3b: " << (PassB ? "PASS" : "FAIL") << std::endl; + + free(usmPtr, Queue); +#endif + return PassA && PassB; +} + +int main() { + queue Queue; + + bool Pass = true; + Pass &= test_0(Queue); + Pass &= test_1(Queue); + Pass &= test_2(Queue); + Pass &= test_3(Queue); + + return Pass ? 0 : 1; +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp new file mode 100644 index 0000000000000..b8c1ef2b06655 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp @@ -0,0 +1,27 @@ +// REQUIRES: aspect-fp64 +// UNSUPPORTED: hip || cuda + +// DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%} + +// RUN: %{build} --offload-new-driver -fsycl-allow-device-image-dependencies -fsycl-device-lib-jit-link %{mathflags} -o %t.out +// RUN: %{run} %t.out + +#include +#include + +using namespace sycl; + +// Check that device lib dependencies are resolved with +// -fsycl-allow-device-image-dependencies. +// TODO this test will become redundant once +// -fsycl-allow-device-image-dependencies is enabled by default. +int main() { + range<1> Range{1}; + queue q; + buffer buffer1(Range); + q.submit([&](sycl::handler &cgh) { + auto Acc = buffer1.get_access(cgh); + cgh.single_task([=]() { Acc[0] = std::acosh(1.0); }); + }); + return 0; +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp new file mode 100644 index 0000000000000..0c29e50522e33 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp @@ -0,0 +1,35 @@ +// Test -fsycl-allow-device-image-dependencies with objects. + +// UNSUPPORTED: cuda || hip + +// RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/a.cpp -I %S/Inputs -c -o %t_a.o +// RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/b.cpp -I %S/Inputs -c -o %t_b.o +// RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/c.cpp -I %S/Inputs -c -o %t_c.o +// RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/d.cpp -I %S/Inputs -c -o %t_d.o +// RUN: %{build} --offload-new-driver -fsycl-allow-device-image-dependencies %t_a.o %t_b.o %t_c.o %t_d.o -I %S/Inputs -o %t.out +// RUN: %{run} %t.out + +#include "a.hpp" +#include +#include + +using namespace sycl; + +class ExeKernel; + +int main() { + int val = 0; + { + buffer buf(&val, range<1>(1)); + queue q; + q.submit([&](handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] = levelA(acc[0]); }); + }); + } + + std::cout << "val=" << std::hex << val << "\n"; + if (val != 0xDCBA) + return (1); + return (0); +} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp new file mode 100644 index 0000000000000..30966f3d0b2d4 --- /dev/null +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp @@ -0,0 +1,24 @@ +// Test -fsycl-allow-device-image-dependencies with a single dynamic library on Windows +// and Linux. + +// UNSUPPORTED: cuda || hip + +// RUN: %clangxx --offload-new-driver -fsycl %fPIC %shared_lib -fsycl-allow-device-image-dependencies -I %S/Inputs \ +// RUN: %S/Inputs/a.cpp \ +// RUN: %S/Inputs/b.cpp \ +// RUN: %S/Inputs/c.cpp \ +// RUN: %S/Inputs/d.cpp \ +// RUN: %S/Inputs/wrapper.cpp \ +// RUN: -o %if windows %{%T/device_single.dll%} %else %{%T/libdevice_single.so%} + +// RUN: %{build} --offload-new-driver -I%S/Inputs -o %t.out \ +// RUN: %if windows \ +// RUN: %{%T/device_single.lib%} \ +// RUN: %else \ +// RUN: %{-L%T -ldevice_single -Wl,-rpath=%T%} + +// RUN: %{run} %t.out + +#include "wrapper.hpp" + +int main() { return (wrapper()); } From 171b1698be415c88bdde4e1b49cea55af3b43f91 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Tue, 12 Nov 2024 23:10:10 +0100 Subject: [PATCH 2/8] Fix clang-format issue --- .../NewOffloadingDriver/singleDynamicLibrary.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp index 30966f3d0b2d4..2139f29e97087 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp @@ -1,5 +1,5 @@ -// Test -fsycl-allow-device-image-dependencies with a single dynamic library on Windows -// and Linux. +// Test -fsycl-allow-device-image-dependencies with a single dynamic library on +// Windows and Linux. // UNSUPPORTED: cuda || hip From 1a76fd8dbafac4bfd4b010c6585bdf28d4bac898 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Wed, 13 Nov 2024 10:37:05 +0100 Subject: [PATCH 3/8] Rename new test directory. --- .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/a.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/a.hpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/b.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/b.hpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/c.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/c.hpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/d.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/d.hpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/wrapper.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/Inputs/wrapper.hpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/dynamic.cpp | 0 .../free_function_kernels.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/math_device_lib.cpp | 0 .../{NewOffloadingDriver => NewOffloadDriver}/objects.cpp | 0 .../singleDynamicLibrary.cpp | 0 15 files changed, 0 insertions(+), 0 deletions(-) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/a.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/a.hpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/b.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/b.hpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/c.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/c.hpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/d.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/d.hpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/wrapper.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/Inputs/wrapper.hpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/dynamic.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/free_function_kernels.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/math_device_lib.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/objects.cpp (100%) rename sycl/test-e2e/DeviceImageDependencies/{NewOffloadingDriver => NewOffloadDriver}/singleDynamicLibrary.cpp (100%) diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/a.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/a.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/a.hpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/a.hpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/a.hpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/b.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/b.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/b.hpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/b.hpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/b.hpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/c.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/c.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/c.hpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/c.hpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/c.hpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/d.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/d.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/d.hpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/d.hpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/d.hpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/wrapper.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/wrapper.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/wrapper.hpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/Inputs/wrapper.hpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/Inputs/wrapper.hpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/dynamic.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/free_function_kernels.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/math_device_lib.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/objects.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp similarity index 100% rename from sycl/test-e2e/DeviceImageDependencies/NewOffloadingDriver/singleDynamicLibrary.cpp rename to sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp From d0d83374b2f9f686366192fe1b0979b134526b8f Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Wed, 13 Nov 2024 12:40:43 +0100 Subject: [PATCH 4/8] Fix test failure. --- .../DeviceImageDependencies/NewOffloadDriver/dynamic.cpp | 1 + .../NewOffloadDriver/free_function_kernels.cpp | 1 + .../DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp | 1 + .../DeviceImageDependencies/NewOffloadDriver/objects.cpp | 1 + .../NewOffloadDriver/singleDynamicLibrary.cpp | 1 + 5 files changed, 5 insertions(+) diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp index 136a3bed85728..ff75d442cb892 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/dynamic.cpp @@ -1,6 +1,7 @@ // Test -fsycl-allow-device-image-dependencies with dynamic libraries. // UNSUPPORTED: cuda || hip +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. // DEFINE: %{dynamic_lib_options} = -fsycl %fPIC %shared_lib -fsycl-allow-device-image-dependencies -I %S/Inputs %if windows %{-DMAKE_DLL %} // DEFINE: %{dynamic_lib_suffix} = %if windows %{dll%} %else %{so%} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp index 9649368328236..74758a837cd46 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/free_function_kernels.cpp @@ -6,6 +6,7 @@ // The name mangling for free function kernels currently does not work with PTX. // UNSUPPORTED: cuda +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. // XFAIL: hip_amd // XFAIL-TRACKER: https://github.com/intel/llvm/issues/15742 diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp index b8c1ef2b06655..949063d62b0c1 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/math_device_lib.cpp @@ -1,5 +1,6 @@ // REQUIRES: aspect-fp64 // UNSUPPORTED: hip || cuda +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. // DEFINE: %{mathflags} = %if cl_options %{/clang:-fno-fast-math%} %else %{-fno-fast-math%} diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp index 0c29e50522e33..9e71e88a35422 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/objects.cpp @@ -1,6 +1,7 @@ // Test -fsycl-allow-device-image-dependencies with objects. // UNSUPPORTED: cuda || hip +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. // RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/a.cpp -I %S/Inputs -c -o %t_a.o // RUN: %clangxx --offload-new-driver -fsycl %S/Inputs/b.cpp -I %S/Inputs -c -o %t_b.o diff --git a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp index 2139f29e97087..cde5f7cdbaad9 100644 --- a/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp +++ b/sycl/test-e2e/DeviceImageDependencies/NewOffloadDriver/singleDynamicLibrary.cpp @@ -2,6 +2,7 @@ // Windows and Linux. // UNSUPPORTED: cuda || hip +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. // RUN: %clangxx --offload-new-driver -fsycl %fPIC %shared_lib -fsycl-allow-device-image-dependencies -I %S/Inputs \ // RUN: %S/Inputs/a.cpp \ From 76d4f3d6eae54bc0ec70eaa7f894dbc9cecddde2 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Wed, 13 Nov 2024 13:03:08 +0100 Subject: [PATCH 5/8] Address code review feedback. --- clang/lib/Driver/ToolChains/Clang.cpp | 7 +++---- clang/test/Driver/sycl-offload-new-driver.c | 17 +++++++++++++++++ .../clang-linker-wrapper/LinkerWrapperOpts.td | 3 --- 3 files changed, 20 insertions(+), 7 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 6fb4f745c58f1..b11fe86c84450 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11486,12 +11486,11 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasArg(options::OPT_fsycl_embed_ir)) CmdArgs.push_back(Args.MakeArgString("-sycl-embed-ir")); - if (Args.hasArg(options::OPT_fsycl_allow_device_image_dependencies)) + if (Args.hasArg(options::OPT_fsycl_allow_device_image_dependencies, + options::OPT_fno_sycl_allow_device_image_dependencies, + false)) CmdArgs.push_back( Args.MakeArgString("-sycl-allow-device-image-dependencies")); - if (Args.hasArg(options::OPT_fno_sycl_allow_device_image_dependencies)) - CmdArgs.push_back( - Args.MakeArgString("-no-sycl-allow-device-image-dependencies")); // Formulate and add any offload-wrapper and AOT specific options. These // are additional options passed in via -Xsycl-target-linker and diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index fbadad2863360..9772b65e033b0 100644 --- a/clang/test/Driver/sycl-offload-new-driver.c +++ b/clang/test/Driver/sycl-offload-new-driver.c @@ -195,3 +195,20 @@ // RUN: --offload-new-driver 2>&1 \ // RUN: | FileCheck -check-prefix NVPTX_CUDA_PATH %s // NVPTX_CUDA_PATH: clang-linker-wrapper{{.*}} "--cuda-path={{.*}}Inputs/CUDA_80/usr/local/cuda" + +/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool +// RUN: %clangxx -fsycl -### --offload-new-driver \ +// RUN: -fsycl-allow-device-image-dependencies %s 2>&1 \ +// RUN: | FileCheck -check-prefix CHECK_DYNAMIC_LINKING %s +// CHECK_DYNAMIC_LINKING: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies" + +/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool +// RUN: %clangxx -fsycl -### --offload-new-driver \ +// RUN: -fno-sycl-allow-device-image-dependencies %s 2>&1 \ +// RUN: | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s +// CHECK_NO_DYNAMIC_LINKING-NOT: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies" + +/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool +// RUN: %clangxx -fsycl -### --offload-new-driver %s 2>&1 \ +// RUN: | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s +// CHECK_NO_DYNAMIC_LINKING-NOT: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies" diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index f6159d100115f..1885fb430b2bd 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -248,6 +248,3 @@ def sycl_dump_device_code_EQ : Joined<["--", "-"], "sycl-dump-device-code=">, def sycl_allow_device_image_dependencies : Flag<["--", "-"], "sycl-allow-device-image-dependencies">, Flags<[WrapperOnlyOption, HelpHidden]>, HelpText<"Allow dependencies between device code images">; -def no_sycl_allow_device_image_dependencies : Flag<["--", "-"], "no-sycl-allow-device-image-dependencies">, - Flags<[WrapperOnlyOption, HelpHidden]>, - HelpText<"Dno not allow dependencies between device code images (default)">; From 9082959fe2cddd7e59fdf8cbc1e4083ed6837b30 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Wed, 13 Nov 2024 13:45:52 +0100 Subject: [PATCH 6/8] Fix build failure. --- clang/lib/Driver/ToolChains/Clang.cpp | 6 +++--- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 3 +-- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index b11fe86c84450..7d331a1fb2982 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11486,9 +11486,9 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasArg(options::OPT_fsycl_embed_ir)) CmdArgs.push_back(Args.MakeArgString("-sycl-embed-ir")); - if (Args.hasArg(options::OPT_fsycl_allow_device_image_dependencies, - options::OPT_fno_sycl_allow_device_image_dependencies, - false)) + if (Args.hasFlag(options::OPT_fsycl_allow_device_image_dependencies, + options::OPT_fno_sycl_allow_device_image_dependencies, + false)) CmdArgs.push_back( Args.MakeArgString("-sycl-allow-device-image-dependencies")); diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index cc682f1278a51..1655f58169836 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -676,8 +676,7 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args, if ((!Args.hasFlag(OPT_no_sycl_remove_unused_external_funcs, OPT_sycl_remove_unused_external_funcs, false) && !SYCLNativeCPU) && - !Args.hasFlag(OPT_sycl_allow_device_image_dependencies, - OPT_no_sycl_allow_device_image_dependencies, false) && + !Args.hasArg(OPT_sycl_allow_device_image_dependencies) && !Triple.isNVPTX() && !Triple.isAMDGPU()) PostLinkArgs.push_back("-emit-only-kernels-as-entry-points"); From 64308cbc18f17975af14ca6a3b7b56c16ecb1b26 Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Wed, 13 Nov 2024 20:55:04 +0100 Subject: [PATCH 7/8] Remove unnecessary line. --- clang/test/Driver/sycl-offload-new-driver.c | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index 9772b65e033b0..288e8d0d51be9 100644 --- a/clang/test/Driver/sycl-offload-new-driver.c +++ b/clang/test/Driver/sycl-offload-new-driver.c @@ -206,7 +206,6 @@ // RUN: %clangxx -fsycl -### --offload-new-driver \ // RUN: -fno-sycl-allow-device-image-dependencies %s 2>&1 \ // RUN: | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s -// CHECK_NO_DYNAMIC_LINKING-NOT: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies" /// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool // RUN: %clangxx -fsycl -### --offload-new-driver %s 2>&1 \ From 03ca7ee31e1782fc298d4505c76e00f6dc629a5a Mon Sep 17 00:00:00 2001 From: "Maronas, Marcos" Date: Thu, 14 Nov 2024 17:32:35 +0100 Subject: [PATCH 8/8] Update comments. --- clang/test/Driver/sycl-offload-new-driver.c | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c index 288e8d0d51be9..b6732b3e9312e 100644 --- a/clang/test/Driver/sycl-offload-new-driver.c +++ b/clang/test/Driver/sycl-offload-new-driver.c @@ -202,12 +202,12 @@ // RUN: | FileCheck -check-prefix CHECK_DYNAMIC_LINKING %s // CHECK_DYNAMIC_LINKING: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies" -/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool +/// Check that -sycl-allow-device-image-dependencies is not passed to clang-linker-wrapper tool // RUN: %clangxx -fsycl -### --offload-new-driver \ // RUN: -fno-sycl-allow-device-image-dependencies %s 2>&1 \ // RUN: | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s -/// Check for -sycl-allow-device-image-dependencies transmission to clang-linker-wrapper tool +/// Check that -sycl-allow-device-image-dependencies is not passed to clang-linker-wrapper tool // RUN: %clangxx -fsycl -### --offload-new-driver %s 2>&1 \ // RUN: | FileCheck -check-prefix CHECK_NO_DYNAMIC_LINKING %s // CHECK_NO_DYNAMIC_LINKING-NOT: clang-linker-wrapper{{.*}} "-sycl-allow-device-image-dependencies"