Skip to content

[clang][SYCL] Disable address space optimization for kernel args #21781

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
KornevNikita merged 5 commits into from
Apr 23, 2026
Merged

[clang][SYCL] Disable address space optimization for kernel args #21781

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
a536e6a
[clang][SYCL] Disable address space optimization for kernel args
Fznamznon Apr 14, 2026
27ff0e5
Merge branch 'sycl' into disable-as-type-punning
Fznamznon Apr 22, 2026
1cf995d
Check the CHECK checks
Fznamznon Apr 22, 2026
47ea2a6
Rename the option
Fznamznon Apr 22, 2026
d51534e
Fix tests after option renaming
Fznamznon Apr 22, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions clang/include/clang/Basic/LangOptions.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -296,6 +296,7 @@ ENUM_LANGOPT(SYCLRangeRounding, SYCLRangeRoundingPreference, 2,
LANGOPT(SYCLExperimentalRangeRounding, 1, 0, NotCompatible, "Use experimental parallel for range rounding")
LANGOPT(SYCLEnableIntHeaderDiags, 1, 0, NotCompatible, "Enable diagnostics that require the SYCL integration header")
LANGOPT(SYCLRTCMode, 1, 0, NotCompatible, "Compile in RTC mode")
LANGOPT(SYCLForceGlobalASInKernelArgs, 1, 0, NotCompatible, "Force global address space in SYCL kernel arguments")

LANGOPT(HIPUseNewLaunchAPI, 1, 0, NotCompatible, "Use new kernel launching API for HIP")
LANGOPT(OffloadUniformBlock, 1, 0, NotCompatible, "Assume that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)")
Expand Down
3 changes: 3 additions & 0 deletions clang/include/clang/Options/Options.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9749,6 +9749,9 @@ defm offload_use_alloca_addrspace_for_srets : BoolFOption<"offload-use-alloca-ad
DefaultTrue,
PosFlag<SetTrue, [], [CC1Option], "Use alloca address space for sret arguments for offloading targets">,
NegFlag<SetFalse>>;
def fsycl_force_global_as_in_kernel_args : Flag<["-"], "fsycl-force-global-as-in-kernel-args">,
HelpText<"Force global address space for USM pointers within SYCL kernel arguments">,
MarshallingInfoFlag<LangOpts<"SYCLForceGlobalASInKernelArgs">>;

def fsycl_remangle_libspirv
: Flag<["-"], "fsycl-remangle-libspirv">,
Expand Down
15 changes: 10 additions & 5 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2333,7 +2333,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
SemaSYCLRef.getASTContext()));
PointerStack.pop_back();
} else if (PointerStack.pop_back_val()) {
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>())
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>() &&
SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs)
RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
SemaSYCLRef.getASTContext()));
}
Expand Down Expand Up @@ -2369,7 +2370,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
PointerStack.pop_back();
} else if (PointerStack.pop_back_val()) {
PointerStack.back() = true;
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>())
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>() &&
SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs)
RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
SemaSYCLRef.getASTContext()));
}
Expand All @@ -2388,7 +2390,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
PointerStack.pop_back();
} else if (PointerStack.pop_back_val()) {
PointerStack.back() = true;
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>())
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>() &&
SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs)
RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
SemaSYCLRef.getASTContext()));
}
Expand Down Expand Up @@ -2418,7 +2421,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
PointerStack.pop_back();
} else if (PointerStack.pop_back_val()) {
PointerStack.back() = true;
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>())
if (!RD->hasAttr<SYCLGenerateNewTypeAttr>() &&
SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs)
RD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
SemaSYCLRef.getASTContext()));
}
Expand Down Expand Up @@ -2452,7 +2456,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
CollectionStack.back() = true;
PointerStack.pop_back();
} else if (PointerStack.pop_back_val()) {
if (!FD->hasAttr<SYCLGenerateNewTypeAttr>())
if (!FD->hasAttr<SYCLGenerateNewTypeAttr>() &&
SemaSYCLRef.getLangOpts().SYCLForceGlobalASInKernelArgs)
FD->addAttr(SYCLGenerateNewTypeAttr::CreateImplicit(
SemaSYCLRef.getASTContext()));
PointerStack.back() = true;
Expand Down
124 changes: 85 additions & 39 deletions clang/test/CodeGenSYCL/array-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,16 +1,17 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 5
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

#include "sycl.hpp"

// CHECK: %struct.__wrapper_class = type { [2 x i32] }
// CHECK: %class.anon = type { [2 x i32] }
// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(4)] }
// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] }

sycl::queue myQueue;

using namespace sycl;

// CHECK: %struct.__wrapper_class = type { [2 x i32] }
// CHECK: %class.anon = type { [2 x i32] }
// CHECK: %struct.__wrapper_class.0 = type { [2 x ptr addrspace(1)] }
// CHECK: %class.anon.1 = type { [2 x ptr addrspace(4)] }

int main() {
int Array[2];
Expand All @@ -21,24 +22,6 @@ int main() {
});
});

// CHECK-LABEL: @{{.*}}IntArray(ptr {{.*}}byval(%struct.__wrapper_class)
// CHECK: %__SYCLKernel = alloca %class.anon, align 4
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
// CHECK: %_arg_Array.ascast = addrspacecast ptr %_arg_Array to ptr addrspace(4)
// CHECK: %Array = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class, ptr addrspace(4) %_arg_Array.ascast, i32 0, i32 0
// CHECK: %arrayinit.begin = getelementptr inbounds [2 x i32], ptr addrspace(4) %Array, i64 0, i64 0
// CHECK: br label %arrayinit.body
// CHECK: arrayinit.body: ; preds = %arrayinit.body, %entry
// CHECK: %arrayinit.index = phi i64 [ 0, %entry ], [ %arrayinit.next, %arrayinit.body ]
// CHECK: %1 = getelementptr inbounds i32, ptr addrspace(4) %arrayinit.begin, i64 %arrayinit.index
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) %0, i64 0, i64 %arrayinit.index
// CHECK: %2 = load i32, ptr addrspace(4) %arrayidx, align 4
// CHECK: store i32 %2, ptr addrspace(4) %1, align 4
// CHECK: %arrayinit.next = add nuw i64 %arrayinit.index, 1
// CHECK: %arrayinit.done = icmp eq i64 %arrayinit.next, 2
// CHECK: br i1 %arrayinit.done, label %arrayinit.end, label %arrayinit.body
// CHECK: arrayinit.end: ; preds = %arrayinit.body

int *ArrayOfPointers[2];
myQueue.submit([&](sycl::handler &h) {
Expand All @@ -47,20 +30,83 @@ int main() {
int local = *ArrayOfPointers[1];
});
});
// CHECK-LABEL: @{{.*}}PtrArray(ptr {{.*}}byval(%struct.__wrapper_class.0)
// CHECK: %__SYCLKernel = alloca %class.anon.1, align 8
// CHECK: %__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
// CHECK: %_arg_ArrayOfPointers.ascast = addrspacecast ptr %_arg_ArrayOfPointers to ptr addrspace(4)
// CHECK: %ArrayOfPointers = getelementptr inbounds nuw %class.anon.1, ptr addrspace(4) %__SYCLKernel.ascast, i32 0, i32 0
// CHECK: %0 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
// CHECK: %arrayidx = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %0, i64 0, i64 0
// CHECK: %1 = load ptr addrspace(1), ptr addrspace(4) %arrayidx, align 8
// CHECK: %2 = addrspacecast ptr addrspace(1) %1 to ptr addrspace(4)
// CHECK: store ptr addrspace(4) %2, ptr addrspace(4) %ArrayOfPointers, align 8
// CHECK: %arrayinit.element = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) %ArrayOfPointers, i64 1
// CHECK: %3 = getelementptr inbounds nuw %struct.__wrapper_class.0, ptr addrspace(4) %_arg_ArrayOfPointers.ascast, i32 0, i32 0
// CHECK: %arrayidx1 = getelementptr inbounds nuw [2 x ptr addrspace(1)], ptr addrspace(4) %3, i64 0, i64 1
// CHECK: %4 = load ptr addrspace(1), ptr addrspace(4) %arrayidx1, align 8
// CHECK: %5 = addrspacecast ptr addrspace(1) %4 to ptr addrspace(4)
// CHECK: store ptr addrspace(4) %5, ptr addrspace(4) %arrayinit.element, align 8
}
// CHECK-LABEL: define weak_odr spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8IntArray(
// CHECK-SAME: ptr noundef byval([[STRUCT___WRAPPER_CLASS:%.*]]) align 4 [[_ARG_ARRAY:%.*]]) #[[ATTR0:[0-9]+]] comdat !srcloc [[META12:![0-9]+]] !kernel_arg_buffer_location [[META13:![0-9]+]] {
// CHECK-NEXT: [[ENTRY:.*]]:
// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 4
// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4)
// CHECK-NEXT: [[_ARG_ARRAY_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ARRAY]] to ptr addrspace(4)
// CHECK-NEXT: [[ARRAY:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___WRAPPER_CLASS]], ptr addrspace(4) [[_ARG_ARRAY_ASCAST]], i32 0, i32 0
// CHECK-NEXT: [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x i32], ptr addrspace(4) [[ARRAY]], i64 0, i64 0
// CHECK-NEXT: br label %[[ARRAYINIT_BODY:.*]]
// CHECK: [[ARRAYINIT_BODY]]:
// CHECK-NEXT: [[ARRAYINIT_INDEX:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[ARRAYINIT_NEXT:%.*]], %[[ARRAYINIT_BODY]] ]
// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds i32, ptr addrspace(4) [[ARRAYINIT_BEGIN]], i64 [[ARRAYINIT_INDEX]]
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x i32], ptr addrspace(4) [[TMP0]], i64 0, i64 [[ARRAYINIT_INDEX]]
// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4
// CHECK-NEXT: store i32 [[TMP2]], ptr addrspace(4) [[TMP1]], align 4
// CHECK-NEXT: [[ARRAYINIT_NEXT]] = add nuw i64 [[ARRAYINIT_INDEX]], 1
// CHECK-NEXT: [[ARRAYINIT_DONE:%.*]] = icmp eq i64 [[ARRAYINIT_NEXT]], 2
// CHECK-NEXT: br i1 [[ARRAYINIT_DONE]], label %[[ARRAYINIT_END:.*]], label %[[ARRAYINIT_BODY]]
// CHECK: [[ARRAYINIT_END]]:
// CHECK-NEXT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(8) [[__SYCLKERNEL_ASCAST]]) #[[ATTR2:[0-9]+]]
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv(
// CHECK-SAME: ptr addrspace(4) noundef align 4 dereferenceable_or_null(8) [[THIS:%.*]]) #[[ATTR1:[0-9]+]] align 2 !srcloc [[META12]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr addrspace(4), align 8
// CHECK-NEXT: [[LOCAL:%.*]] = alloca i32, align 4
// CHECK-NEXT: [[THIS_ADDR_ASCAST:%.*]] = addrspacecast ptr [[THIS_ADDR]] to ptr addrspace(4)
// CHECK-NEXT: [[LOCAL_ASCAST:%.*]] = addrspacecast ptr [[LOCAL]] to ptr addrspace(4)
// CHECK-NEXT: store ptr addrspace(4) [[THIS]], ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[ARRAY:%.*]] = getelementptr inbounds nuw [[CLASS_ANON:%.*]], ptr addrspace(4) [[THIS1]], i32 0, i32 0
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x i32], ptr addrspace(4) [[ARRAY]], i64 0, i64 1
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) [[ARRAYIDX]], align 4
// CHECK-NEXT: store i32 [[TMP0]], ptr addrspace(4) [[LOCAL_ASCAST]], align 4
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define weak_odr spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E8PtrArray(
// CHECK-SAME: ptr noundef byval([[STRUCT___WRAPPER_CLASS_0:%.*]]) align 8 [[_ARG_ARRAYOFPOINTERS:%.*]]) #[[ATTR0]] comdat !srcloc [[META14:![0-9]+]] !kernel_arg_buffer_location [[META13]] {
// CHECK-NEXT: [[ENTRY:.*]]:
// CHECK-NEXT: [[__SYCLKERNEL:%.*]] = alloca [[CLASS_ANON_1:%.*]], align 8
// CHECK-NEXT: [[__SYCLKERNEL_ASCAST:%.*]] = addrspacecast ptr [[__SYCLKERNEL]] to ptr addrspace(4)
// CHECK-NEXT: [[_ARG_ARRAYOFPOINTERS_ASCAST:%.*]] = addrspacecast ptr [[_ARG_ARRAYOFPOINTERS]] to ptr addrspace(4)
// CHECK-NEXT: [[ARRAYOFPOINTERS:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1]], ptr addrspace(4) [[__SYCLKERNEL_ASCAST]], i32 0, i32 0
// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw [[STRUCT___WRAPPER_CLASS_0]], ptr addrspace(4) [[_ARG_ARRAYOFPOINTERS_ASCAST]], i32 0, i32 0
// CHECK-NEXT: [[ARRAYINIT_BEGIN:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[ARRAYOFPOINTERS]], i64 0, i64 0
// CHECK-NEXT: br label %[[ARRAYINIT_BODY:.*]]
// CHECK: [[ARRAYINIT_BODY]]:
// CHECK-NEXT: [[ARRAYINIT_INDEX:%.*]] = phi i64 [ 0, %[[ENTRY]] ], [ [[ARRAYINIT_NEXT:%.*]], %[[ARRAYINIT_BODY]] ]
// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds ptr addrspace(4), ptr addrspace(4) [[ARRAYINIT_BEGIN]], i64 [[ARRAYINIT_INDEX]]
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds nuw [2 x ptr addrspace(4)], ptr addrspace(4) [[TMP0]], i64 0, i64 [[ARRAYINIT_INDEX]]
// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8
// CHECK-NEXT: store ptr addrspace(4) [[TMP2]], ptr addrspace(4) [[TMP1]], align 8
// CHECK-NEXT: [[ARRAYINIT_NEXT]] = add nuw i64 [[ARRAYINIT_INDEX]], 1
// CHECK-NEXT: [[ARRAYINIT_DONE:%.*]] = icmp eq i64 [[ARRAYINIT_NEXT]], 2
// CHECK-NEXT: br i1 [[ARRAYINIT_DONE]], label %[[ARRAYINIT_END:.*]], label %[[ARRAYINIT_BODY]]
// CHECK: [[ARRAYINIT_END]]:
// CHECK-NEXT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[__SYCLKERNEL_ASCAST]]) #[[ATTR2]]
// CHECK-NEXT: ret void
//
//
// CHECK-LABEL: define internal spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv(
// CHECK-SAME: ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[THIS:%.*]]) #[[ATTR1]] align 2 !srcloc [[META14]] {
// CHECK-NEXT: [[ENTRY:.*:]]
// CHECK-NEXT: [[THIS_ADDR:%.*]] = alloca ptr addrspace(4), align 8
// CHECK-NEXT: [[LOCAL:%.*]] = alloca i32, align 4
// CHECK-NEXT: [[THIS_ADDR_ASCAST:%.*]] = addrspacecast ptr [[THIS_ADDR]] to ptr addrspace(4)
// CHECK-NEXT: [[LOCAL_ASCAST:%.*]] = addrspacecast ptr [[LOCAL]] to ptr addrspace(4)
// CHECK-NEXT: store ptr addrspace(4) [[THIS]], ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[THIS1:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[THIS_ADDR_ASCAST]], align 8
// CHECK-NEXT: [[ARRAYOFPOINTERS:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1:%.*]], ptr addrspace(4) [[THIS1]], i32 0, i32 0
// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [2 x ptr addrspace(4)], ptr addrspace(4) [[ARRAYOFPOINTERS]], i64 0, i64 1
// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ARRAYIDX]], align 8
// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[TMP0]], align 4
// CHECK-NEXT: store i32 [[TMP1]], ptr addrspace(4) [[LOCAL_ASCAST]], align 4
// CHECK-NEXT: ret void
26 changes: 15 additions & 11 deletions clang/test/CodeGenSYCL/free_function_kernel_params.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,7 @@
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \
// RUN: -emit-llvm %s -o - | FileCheck %s
// RUN: -emit-llvm %s -o - | FileCheck --check-prefixes=CHECK,GEN-AS %s
// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -triple spir64 \
// RUN: -fsycl-force-global-as-in-kernel-args -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,GLOB-AS
// This test checks parameter IR generation for free functions with parameters
// of non-decomposed struct type, work group memory type, dynamic work group memory type
// and special types.
Expand Down Expand Up @@ -47,16 +49,18 @@ void ff_6(KArgWithPtrArray<ArrSize> KArg) {

template void ff_6(KArgWithPtrArray<TestArrSize> KArg);

// CHECK: %struct.NoPointers = type { i32 }
// CHECK: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
// CHECK: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
// CHECK: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
// CHECK: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.4 }
// CHECK: %struct.__generated_Pointers.4 = type { ptr addrspace(1), ptr addrspace(1) }
// CHECK: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] }
// CHECK: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] }
// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3)
// CHECK: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg)
// GEN-AS: %struct.NoPointers = type { i32 }
// GEN-AS: %struct.Pointers = type { ptr addrspace(4), ptr addrspace(4) }
// GEN-AS: %struct.Agg = type { %struct.NoPointers, i32, ptr addrspace(4), %struct.Pointers }
// GLOB-AS: %struct.__generated_Pointers = type { ptr addrspace(1), ptr addrspace(1) }
// GLOB-AS: %struct.__generated_Agg = type { %struct.NoPointers, i32, ptr addrspace(1), %struct.__generated_Pointers.4 }
// GLOB-AS: %struct.__generated_Pointers.4 = type { ptr addrspace(1), ptr addrspace(1) }
// GLOB-AS: %struct.__generated_KArgWithPtrArray = type { [3 x ptr addrspace(1)], [3 x i32], [3 x i32] }
// GLOB-AS: %struct.KArgWithPtrArray = type { [3 x ptr addrspace(4)], [3 x i32], [3 x i32] }
// GLOB-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.__generated_Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.__generated_Agg) align 8 %__arg_S3)
// GLOB-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.__generated_KArgWithPtrArray) align 8 %__arg_KArg)
// GEN-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel{{.*}}(ptr noundef byval(%struct.NoPointers) align 4 %__arg_S1, ptr noundef byval(%struct.Pointers) align 8 %__arg_S2, ptr noundef byval(%struct.Agg) align 8 %__arg_S3)
// GEN-AS: define dso_local spir_kernel void @{{.*}}__sycl_kernel_ff_6{{.*}}(ptr noundef byval(%struct.KArgWithPtrArray) align 8 %__arg_KArg)

__attribute__((sycl_device))
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]]
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/generated-types-initialization.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -fsycl-force-global-as-in-kernel-args %s -o - | FileCheck %s

// This test checks that compiler generates correct code when kernel arguments
// are structs that contain pointers but not decomposed.
Expand Down
Loading
Loading