Skip to content

Commit 5208484

Browse files
sommerlukassommerlukas
authored
[SYCL] Lift restrictions on free-function kernels when compiling at runtime (#15892)
In order to be able to generate correct and complete information for the integration header, the current implementation places some restrictions on free-function kernels and their parameters. For example, parameters of free function kernels need to be forward-declarable. However, when compiling SYCL code at runtime (RTC), e.g., through the `kernel_compiler` extension, host code is typically not relevant, so the integration header is not as relevant and some restrictions on free-function kernels can be lifted. This PR introduces a `-fsycl-rtc-mode` flag (and it's negative equivalent) to deactivate some restrictions on free-function kernels and omit some information for free-function kernels from the integration header. --------- Signed-off-by: Lukas Sommer <lukas.sommer@codeplay.com>
1 parent 4d00087 commit 5208484

File tree

6 files changed

+133
-0
lines changed

6 files changed

+133
-0
lines changed

‎clang/include/clang/Basic/LangOptions.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -326,6 +326,7 @@ LANGOPT(SYCLExperimentalRangeRounding, 1, 0, "Use experimental parallel for rang
326326
LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, "Enable diagnostics that require the "
327327
"SYCL integration header")
328328
LANGOPT(SYCLIsNativeCPU , 1, 0, "Generate code for SYCL Native CPU")
329+
LANGOPT(SYCLRTCMode, 1, 0, "Compile in RTC mode")
329330

330331
LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")
331332
LANGOPT(OffloadUniformBlock, 1, 0, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)")

‎clang/include/clang/Driver/Options.td

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6877,6 +6877,11 @@ defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-me
68776877
NegFlag<SetFalse, [], [ClangOption, CLOption], "Do not enforce using "
68786878
"stateless memory accesses.">,
68796879
BothFlags<[], [ClangOption, CLOption, CC1Option], "">>;
6880+
defm sycl_rtc_mode: BoolFOption<"sycl-rtc-mode",
6881+
LangOpts<"SYCLRTCMode">, DefaultFalse,
6882+
PosFlag<SetTrue, [], [ClangOption], "Enable">,
6883+
NegFlag<SetFalse, [], [ClangOption], "Disable">,
6884+
BothFlags<[HelpHidden], [ClangOption, CC1Option], " RTC mode in SYCL.">>;
68806885
// TODO: Remove this option once ESIMD headers are updated to
68816886
// guard vectors to be device only.
68826887
def fno_sycl_esimd_build_host_code : Flag<["-"], "fno-sycl-esimd-build-host-code">,

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5589,6 +5589,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
55895589
Args.AddLastArg(CmdArgs, options::OPT_fsycl_decompose_functor,
55905590
options::OPT_fno_sycl_decompose_functor);
55915591

5592+
Args.AddLastArg(CmdArgs, options::OPT_fsycl_rtc_mode,
5593+
options::OPT_fno_sycl_rtc_mode);
5594+
55925595
// Forward -fsycl-instrument-device-code option to cc1. This option will
55935596
// only be used for SPIR/SPIR-V based targets.
55945597
if (Triple.isSPIROrSPIRV())

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2040,6 +2040,11 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
20402040
}
20412041

20422042
bool handleStructType(ParmVarDecl *PD, QualType ParamTy) final {
2043+
if (SemaSYCLRef.getLangOpts().SYCLRTCMode) {
2044+
// When compiling in RTC mode, the restriction regarding forward
2045+
// declarations doesn't apply, as we don't need the integration header.
2046+
return isValid();
2047+
}
20432048
CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl();
20442049
// For free functions all struct/class kernel arguments are forward declared
20452050
// in integration header, that adds additional restrictions for kernel
@@ -6453,6 +6458,13 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
64536458
O << "} // namespace _V1\n";
64546459
O << "} // namespace sycl\n";
64556460

6461+
// The rest of this function only applies to free-function kernels. However,
6462+
// in RTC mode, we do not need integration header information for
6463+
// free-function kernels, so we can return early here.
6464+
if (S.getLangOpts().SYCLRTCMode) {
6465+
return;
6466+
}
6467+
64566468
unsigned ShimCounter = 1;
64576469
int FreeFunctionCount = 0;
64586470
for (const KernelDesc &K : KernelDescs) {

‎clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp

Lines changed: 80 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,80 @@
1+
// 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
2+
// RUN: FileCheck -input-file=%t.rtc.h --check-prefixes=CHECK,CHECK-RTC %s
3+
4+
// 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
5+
// RUN: FileCheck -input-file=%t.nortc.h --check-prefixes=CHECK,CHECK-NORTC %s
6+
7+
// This test checks that free-function kernel information is included or
8+
// excluded from the integration header, depending on the '-fsycl-rtc-mode'
9+
// flag.
10+
11+
#include "sycl.hpp"
12+
13+
[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 2)]]
14+
void free_function_single(int* ptr, int start, int end){
15+
for(int i = start; i < end; ++i){
16+
ptr[i] = start + 66;
17+
}
18+
}
19+
20+
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
21+
void free_function_nd_range(int* ptr, int start, int end){
22+
for(int i = start; i < end; ++i){
23+
ptr[i] = start + 66;
24+
}
25+
}
26+
27+
template<typename KernelName, typename KernelFunc>
28+
__attribute__((sycl_kernel)) void kernel(const KernelFunc &kernelFunc){
29+
kernelFunc();
30+
}
31+
32+
int main(){
33+
sycl::accessor<int, 1, sycl::access::mode::read_write> accessorA;
34+
kernel<class Kernel_Function>(
35+
[=]() {
36+
accessorA.use();
37+
});
38+
return 0;
39+
}
40+
41+
42+
// CHECK: const char* const kernel_names[] = {
43+
// CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_singlePiii",
44+
// CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_nd_rangePiii",
45+
// CHECK-NEXT: "{{.*}}Kernel_Function",
46+
47+
48+
// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; }
49+
// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; }
50+
// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; }
51+
52+
// CHECK-RTC-NOT: free_function_single_kernel
53+
// CHECK-RTC-NOT: free_function_nd_range
54+
55+
// CHECK-NORTC: void free_function_single(int *ptr, int start, int end);
56+
// CHECK-NORTC: static constexpr auto __sycl_shim[[#FIRST:]]()
57+
// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_single;
58+
59+
// CHECK-NORTC: struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#FIRST]]()> {
60+
// CHECK-NORTC-NEXT: static constexpr bool value = true;
61+
62+
// CHECK-NORTC: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim[[#FIRST]]()> {
63+
// CHECK-NORTC-NEXT: static constexpr bool value = true;
64+
65+
66+
// CHECK-NORTC: void free_function_nd_range(int *ptr, int start, int end);
67+
// CHECK-NORTC: static constexpr auto __sycl_shim[[#SECOND:]]() {
68+
// CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_nd_range;
69+
70+
// CHECK-NORTC: struct ext::oneapi::experimental::is_kernel<__sycl_shim[[#SECOND]]()> {
71+
// CHECK-NORTC-NEXT: static constexpr bool value = true;
72+
73+
// CHECK-NORTC: struct ext::oneapi::experimental::is_nd_range_kernel<__sycl_shim2(), 2> {
74+
// CHECK-NORTC-NEXT: static constexpr bool value = true;
75+
76+
// CHECK-NORTC: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#FIRST]]()>() {
77+
// CHECK-NORTC-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_singlePiii"});
78+
79+
// CHECK-NORTC: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim[[#SECOND]]()>() {
80+
// CHECK-NORTC-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"{{.*}}__sycl_kernel_free_function_nd_rangePiii"});

‎clang/test/Driver/sycl-rtc-mode.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
///
2+
/// Perform driver test for SYCL RTC mode.
3+
///
4+
5+
/// Check that the '-fsycl-rtc-mode' is correctly forwarded to the device
6+
/// compilation and only to the device compilation.
7+
8+
// RUN: %clangxx -fsycl -fsycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \
9+
// RUN: | FileCheck %s
10+
11+
// RUN: %clangxx -fsycl -fsycl-rtc-mode --offload-new-driver %s -### 2>&1 \
12+
// RUN: | FileCheck %s
13+
14+
// CHECK: clang{{.*}} "-fsycl-is-device"
15+
// CHECK-SAME: -fsycl-rtc-mode
16+
// CHECK: clang{{.*}} "-fsycl-is-host"
17+
// CHECK-NOT: -fsycl-rtc-mode
18+
19+
20+
/// Check that the '-fno-sycl-rtc-mode' is correctly forwarded to the device
21+
/// compilation and only to the device compilation.
22+
23+
// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --no-offload-new-driver %s -### 2>&1 \
24+
// RUN: | FileCheck %s --check-prefix=NEGATIVE
25+
26+
// RUN: %clangxx -fsycl -fno-sycl-rtc-mode --offload-new-driver %s -### 2>&1 \
27+
// RUN: | FileCheck %s --check-prefix=NEGATIVE
28+
29+
// NEGATIVE: clang{{.*}} "-fsycl-is-device"
30+
// NEGATIVE-SAME: -fno-sycl-rtc-mode
31+
// NEGATIVE: clang{{.*}} "-fsycl-is-host"
32+
// NEGATIVE-NOT: -fsycl-rtc-mode

0 commit comments

Comments
 (0)