From 44ad31f5db62a8bb9aac4f69970c29786ce7ae8e Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Mon, 21 Oct 2024 14:49:36 +0100 Subject: [PATCH 1/7] Omit free-function kernels from integration header Add a flag to compile in "RTC mode", in which we do not care about complete information in the integration header. If the RTC mode is enabled, omit free-function kernels fom the integration header to lift restrictions regarding the ability to forward declare free-function kernel parameters. Signed-off-by: Lukas Sommer --- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Driver/Options.td | 5 +++++ clang/lib/Driver/ToolChains/Clang.cpp | 5 +++++ clang/lib/Sema/SemaSYCL.cpp | 12 ++++++++++++ 4 files changed, 23 insertions(+) diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 2b99cc35e0f1d..8bf06d5a3a8e0 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -326,6 +326,7 @@ LANGOPT(SYCLExperimentalRangeRounding, 1, 0, "Use experimental parallel for rang LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the " "SYCL integration header") LANGOPT(SYCLIsNativeCPU , 1, 0, "Generate code for SYCL Native CPU") +LANGOPT(SYCLRTCMode, 1, 0, "Compile in RTC mode") LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP") LANGOPT(OffloadUniformBlock, 1, 0, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 8ff41673575ac..7ccb73374d107 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -6877,6 +6877,11 @@ defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-me NegFlag, BothFlags<[], [ClangOption, CLOption, CC1Option], "">>; +defm sycl_rtc_mode: BoolFOption<"sycl-rtc-mode", + LangOpts<"SYCLRTCMode">, DefaultFalse, + PosFlag, + NegFlag, + BothFlags<[HelpHidden], [ClangOption, CC1Option], " RTC mode in SYCL.">>; // TODO: Remove this option once ESIMD headers are updated to // guard vectors to be device only. def fno_sycl_esimd_build_host_code : Flag<["-"], "fno-sycl-esimd-build-host-code">, diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 8033cc490adec..ec66c24228bda 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5583,6 +5583,11 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-fsycl-allow-func-ptr"); } + if (Args.hasFlag(options::OPT_fsycl_rtc_mode, + options::OPT_fno_sycl_rtc_mode, false)) { + CmdArgs.push_back("-fsycl-rtc-mode"); + } + Args.AddLastArg(CmdArgs, options::OPT_fsycl_decompose_functor, options::OPT_fno_sycl_decompose_functor); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d8925669b1ae3..31737e5a60b65 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2040,6 +2040,11 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final { + if(this->SemaSYCLRef.getLangOpts().SYCLRTCMode){ + // When compiling in RTC mode, the restriction regarding forward + // declarations doesn't apply, as we don't need the integration header. + return isValid(); + } CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); // For free functions all struct/class kernel arguments are forward declared // in integration header, that adds additional restrictions for kernel @@ -6453,6 +6458,13 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "} // namespace _V1\n"; O << "} // namespace sycl\n"; + // The rest of this function only applies to free-function kernels. However, + // in RTC mode, we do not need integration header information for + // free-function kernels, so we can return early here. + if(this->S.getLangOpts().SYCLRTCMode){ + return; + } + unsigned ShimCounter = 1; int FreeFunctionCount = 0; for (const KernelDesc &K : KernelDescs) { From ef81032fcb80523174a452f2a3ac2ac45e5a9095 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Mon, 28 Oct 2024 13:12:56 +0000 Subject: [PATCH 2/7] Add test for RTC mode Signed-off-by: Lukas Sommer --- .../free_function_int_header_rtc_mode.cpp | 80 +++++++++++++++++++ 1 file changed, 80 insertions(+) create mode 100644 clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp new file mode 100644 index 0000000000000..ef730635b37c4 --- /dev/null +++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp @@ -0,0 +1,80 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-rtc-mode -fsycl-int-header=%t.rtc.h %s +// RUN: FileCheck -input-file=%t.rtc.h --check-prefixes=CHECK,CHECK-RTC %s + +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -sycl-std=2020 -fno-sycl-rtc-mode -fsycl-int-header=%t.nortc.h %s +// RUN: FileCheck -input-file=%t.nortc.h --check-prefixes=CHECK,CHECK-NORTC %s + +// This test checks that free-function kernel information is included or +// excluded from the integration header, depending on the '-fsycl-rtc-mode' +// flag. + +#include "sycl.hpp" + +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]] +void free_function_single(int* ptr, int start, int end){ + for(int i = start; i < end; ++i){ + ptr[i] = start + 66; + } +} + +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]] +void free_function_nd_range(int* ptr, int start, int end){ + for(int i = start; i < end; ++i){ + ptr[i] = start + 66; + } +} + +template +__attribute__((sycl_kernel)) void kernel(const KernelFunc &kernelFunc){ + kernelFunc(); +} + +int main(){ + sycl::accessor accessorA; + kernel( + [=]() { + accessorA.use(); + }); + return 0; +} + + +// CHECK: const char* const kernel_names[] = { +// CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_singlePiii", +// CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_nd_rangePiii", +// CHECK-NEXT: "{{.*}}Kernel_Function", + + +// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; } +// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; } +// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; } + +// CHECK-RTC-NOT: free_function_single_kernel +// CHECK-RTC-NOT: free_function_nd_range + +// CHECK-NORTC: void free_function_single(int *ptr, int start, int end); +// CHECK-NORTC: static constexpr auto __sycl_shim[[#FIRST:]]() +// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_single; + +// CHECK-NORTC: struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#FIRST]]()> { +// CHECK-NORTC-NEXT: static constexpr bool value = true; + +// CHECK-NORTC: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim[[#FIRST]]()> { +// CHECK-NORTC-NEXT: static constexpr bool value = true; + + +// CHECK-NORTC: void free_function_nd_range(int *ptr, int start, int end); +// CHECK-NORTC: static constexpr auto __sycl_shim[[#SECOND:]]() { +// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_nd_range; + +// CHECK-NORTC: struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#SECOND]]()> { +// CHECK-NORTC-NEXT: static constexpr bool value = true; + +// CHECK-NORTC: struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim2(), 2> { +// CHECK-NORTC-NEXT: static constexpr bool value = true; + +// CHECK-NORTC: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() { +// CHECK-NORTC-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_singlePiii"}); + +// CHECK-NORTC: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() { +// CHECK-NORTC-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_nd_rangePiii"}); From b49f05bc2053af5fdcc5e0c53f6dcbb0c01a2dab Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Mon, 28 Oct 2024 14:07:21 +0000 Subject: [PATCH 3/7] Simplify option forwarding Signed-off-by: Lukas Sommer --- clang/lib/Driver/ToolChains/Clang.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index ec66c24228bda..690e18e181979 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5583,10 +5583,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-fsycl-allow-func-ptr"); } - if (Args.hasFlag(options::OPT_fsycl_rtc_mode, - options::OPT_fno_sycl_rtc_mode, false)) { - CmdArgs.push_back("-fsycl-rtc-mode"); - } + Args.AddLastArg(CmdArgs, options::OPT_fsycl_rtc_mode, + options::OPT_fno_sycl_rtc_mode); Args.AddLastArg(CmdArgs, options::OPT_fsycl_decompose_functor, options::OPT_fno_sycl_decompose_functor); From f22abba36c6ed7afe0d065f934c6d516cb0f3474 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Mon, 28 Oct 2024 15:17:30 +0000 Subject: [PATCH 4/7] Code formatting Signed-off-by: Lukas Sommer --- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 31737e5a60b65..aa04be592d41d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2040,7 +2040,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final { - if(this->SemaSYCLRef.getLangOpts().SYCLRTCMode){ + if (this->SemaSYCLRef.getLangOpts().SYCLRTCMode) { // When compiling in RTC mode, the restriction regarding forward // declarations doesn't apply, as we don't need the integration header. return isValid(); @@ -6461,7 +6461,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // The rest of this function only applies to free-function kernels. However, // in RTC mode, we do not need integration header information for // free-function kernels, so we can return early here. - if(this->S.getLangOpts().SYCLRTCMode){ + if (this->S.getLangOpts().SYCLRTCMode) { return; } From 40baf926cc0522d2cd2886dc012bf83ef49a8180 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Tue, 29 Oct 2024 08:17:13 +0000 Subject: [PATCH 5/7] Address PR feedback Signed-off-by: Lukas Sommer --- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index aa04be592d41d..7dd3afdc7ad5e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2040,7 +2040,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final { - if (this->SemaSYCLRef.getLangOpts().SYCLRTCMode) { + if (SemaSYCLRef.getLangOpts().SYCLRTCMode) { // When compiling in RTC mode, the restriction regarding forward // declarations doesn't apply, as we don't need the integration header. return isValid(); @@ -6461,7 +6461,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // The rest of this function only applies to free-function kernels. However, // in RTC mode, we do not need integration header information for // free-function kernels, so we can return early here. - if (this->S.getLangOpts().SYCLRTCMode) { + if (S.getLangOpts().SYCLRTCMode) { return; } From 3ce13f7db271bcfe64fa29fe7bdc100d6ab51168 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Wed, 30 Oct 2024 16:28:17 +0000 Subject: [PATCH 6/7] Address more PR feedback Signed-off-by: Lukas Sommer --- clang/include/clang/Driver/Options.td | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 7ccb73374d107..053bf5e63d9ad 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -6879,8 +6879,8 @@ defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-me BothFlags<[], [ClangOption, CLOption, CC1Option], "">>; defm sycl_rtc_mode: BoolFOption<"sycl-rtc-mode", LangOpts<"SYCLRTCMode">, DefaultFalse, - PosFlag, - NegFlag, + PosFlag, + NegFlag, BothFlags<[HelpHidden], [ClangOption, CC1Option], " RTC mode in SYCL.">>; // TODO: Remove this option once ESIMD headers are updated to // guard vectors to be device only. From bc9d7bb26d9ca5da937aa8f79d141dc52609cdb1 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 1 Nov 2024 15:13:33 +0000 Subject: [PATCH 7/7] Add driver forwarding test Signed-off-by: Lukas Sommer --- clang/test/Driver/sycl-rtc-mode.cpp | 32 +++++++++++++++++++++++++++++ 1 file changed, 32 insertions(+) create mode 100644 clang/test/Driver/sycl-rtc-mode.cpp diff --git a/clang/test/Driver/sycl-rtc-mode.cpp b/clang/test/Driver/sycl-rtc-mode.cpp new file mode 100644 index 0000000000000..c62d3c4d6ade7 --- /dev/null +++ b/clang/test/Driver/sycl-rtc-mode.cpp @@ -0,0 +1,32 @@ +/// +/// Perform driver test for SYCL RTC mode. +/// + +/// Check that the '-fsycl-rtc-mode' is correctly forwarded to the device +/// compilation and only to the device compilation. + +// RUN: %clangxx -fsycl -fsycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \ +// RUN: | FileCheck %s + +// RUN: %clangxx -fsycl -fsycl-rtc-mode --offload-new-driver %s -### 2>&1 \ +// RUN: | FileCheck %s + +// CHECK: clang{{.*}} "-fsycl-is-device" +// CHECK-SAME: -fsycl-rtc-mode +// CHECK: clang{{.*}} "-fsycl-is-host" +// CHECK-NOT: -fsycl-rtc-mode + + +/// Check that the '-fno-sycl-rtc-mode' is correctly forwarded to the device +/// compilation and only to the device compilation. + +// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NEGATIVE + +// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --offload-new-driver %s -### 2>&1 \ +// RUN: | FileCheck %s --check-prefix=NEGATIVE + +// NEGATIVE: clang{{.*}} "-fsycl-is-device" +// NEGATIVE-SAME: -fno-sycl-rtc-mode +// NEGATIVE: clang{{.*}} "-fsycl-is-host" +// NEGATIVE-NOT: -fsycl-rtc-mode