Skip to content

fix merge with upstream kernel attribute refactor #18868

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
merged 3 commits into from
Jun 10, 2025
Merged

fix merge with upstream kernel attribute refactor #18868

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
6ceb309
fix merge with upstream kernel attribute refactor
sarnex Jun 6, 2025
3380a98
Merge remote-tracking branch 'upstream/sycl-web' into sycl-web
sarnex Jun 10, 2025
fabfd02
add explicit native cpu/fallback braches
sarnex Jun 10, 2025
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
2 changes: 1 addition & 1 deletion clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1650,7 +1650,7 @@ def DeviceKernel : DeclOrTypeAttr {
// list, but here we have the same spelling with unscores and without,
// so handle that case manually.
return A.getAttributeSpellingListIndex() == Keyword_kernel ||
A.getAttrName()->getName() == "kernel";
(A.getAttrName() && A.getAttrName()->getName() == "kernel");
Copy link
Contributor

@Fznamznon Fznamznon Jun 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

How is this change specific to this downstream?

Copy link
Contributor Author

@sarnex sarnex Jun 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry it's not but the issue was only exposed downstream, I'll fix it upstream.

Copy link
Contributor Author

@sarnex sarnex Jun 10, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

llvm/llvm-project#143546

}
static inline bool isOpenCLSpelling(const AttributeCommonInfo* A) {
if (!A) return false;
Expand Down
4 changes: 3 additions & 1 deletion clang/lib/Basic/Targets/X86.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -796,7 +796,7 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
case CC_X86RegCall:
return CCCR_OK;
case CC_DeviceKernel:
return IsOpenCL ? CCCR_OK : CCCR_Warning;
return IsOpenCL || IsSYCLDevice ? CCCR_OK : CCCR_Warning;
default:
return CCCR_Warning;
}
Expand Down Expand Up @@ -842,10 +842,12 @@ class LLVM_LIBRARY_VISIBILITY X86_64TargetInfo : public X86TargetInfo {
void adjust(DiagnosticsEngine &Diags, LangOptions &Opts) override {
TargetInfo::adjust(Diags, Opts);
IsOpenCL = Opts.OpenCL;
IsSYCLDevice = Opts.SYCLIsDevice;
}

private:
bool IsOpenCL = false;
bool IsSYCLDevice = false;
};

// x86-64 UEFI target
Expand Down
13 changes: 9 additions & 4 deletions clang/lib/CodeGen/CGCall.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,6 +94,10 @@ unsigned CodeGenTypes::ClangCallConvToLLVMCallConv(CallingConv CC) {
return llvm::CallingConv::AMDGPU_KERNEL;
if (CGM.getTriple().isNVPTX())
return llvm::CallingConv::PTX_Kernel;
if (CGM.getLangOpts().SYCLIsNativeCPU)
return CGM.getTarget().getDefaultCallingConv();
if (CGM.getLangOpts().SYCLIsDevice)
return CGM.getTarget().getDefaultCallingConv();
llvm_unreachable("Unknown kernel calling convention");
}
case CC_PreserveMost:
Expand Down Expand Up @@ -2628,8 +2632,8 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
NumElemsParam);
}

if (DeviceKernelAttr::isOpenCLSpelling(
TargetDecl->getAttr<DeviceKernelAttr>()) &&
if (TargetDecl->hasAttr<DeviceKernelAttr>() &&
(getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) &&
CallingConv != CallingConv::CC_C &&
CallingConv != CallingConv::CC_SpirFunction) {
// Check CallingConv to avoid adding uniform-work-group-size attribute to
Expand Down Expand Up @@ -3013,8 +3017,9 @@ void CodeGenModule::ConstructAttributeList(StringRef Name,
// > data type, the OpenCL compiler can assume that the pointee is always
// > appropriately aligned as required by the data type.
if (TargetDecl &&
DeviceKernelAttr::isOpenCLSpelling(
TargetDecl->getAttr<DeviceKernelAttr>()) &&

(TargetDecl->hasAttr<DeviceKernelAttr>() &&
(getLangOpts().OpenCL || getLangOpts().SYCLIsDevice)) &&
ParamType->isPointerType()) {
QualType PTy = ParamType->getPointeeType();
if (!PTy->isIncompleteType() && PTy->isConstantSizeType()) {
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaChecking.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4014,7 +4014,7 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto,
if (FD)
diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc);

if (FD && FD->hasAttr<DeviceKernelAttr>())
if (FD && FD->hasAttr<DeviceKernelAttr>() && getLangOpts().isSYCL())
SYCL().CheckSYCLKernelCall(FD, Args);

// Diagnose variadic calls in SYCL.
Expand Down
3 changes: 2 additions & 1 deletion clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20762,7 +20762,8 @@ Sema::FunctionEmissionStatus Sema::getEmissionStatus(const FunctionDecl *FD,
return FunctionEmissionStatus::TemplateDiscarded;

if (LangOpts.SYCLIsDevice &&
(FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()))
(FD->hasAttr<SYCLDeviceAttr>() || FD->hasAttr<DeviceKernelAttr>()) &&
!FD->hasAttr<ArtificialAttr>())
return FunctionEmissionStatus::Emitted;

// Check whether this function is an externally visible definition.
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5355,7 +5355,7 @@ static void handleCallConvAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
static void handleDeviceKernelAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *FD = dyn_cast_or_null<FunctionDecl>(D);
bool IsFunctionTemplate = FD && FD->getDescribedFunctionTemplate();
if (S.getLangOpts().SYCLIsDevice) {
if (S.getLangOpts().isSYCL()) {
if (!IsFunctionTemplate) {
S.Diag(AL.getLoc(), diag::warn_attribute_wrong_decl_type_str)
<< AL << AL.isRegularKeywordAttribute() << "function templates";
Expand Down
11 changes: 5 additions & 6 deletions clang/lib/Sema/SemaType.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3785,12 +3785,11 @@ static CallingConv getCCForDeclaratorChunk(
}
}
}
if (!S.getLangOpts().isSYCL()) {
for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
if (AL.getKind() == ParsedAttr::AT_DeviceKernel) {
CC = CC_DeviceKernel;
break;
}
for (const ParsedAttr &AL : D.getDeclSpec().getAttributes()) {
if (AL.getKind() == ParsedAttr::AT_DeviceKernel &&
!DeviceKernelAttr::isSYCLSpelling(AL)) {
CC = CC_DeviceKernel;
break;
}
}
return CC;
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/kernel-op-calls.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ int main() {

Q.submit([&](sycl::handler& cgh) {
ESIMDFunctor EF;
// CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// CHECK: define dso_local spir_kernel void @_ZTS12ESIMDFunctor() {{.*}} !sycl_explicit_simd !{{[0-9]+}} !intel_reqd_sub_group_size !{{[0-9]+}} !work_group_size_hint !{{[0-9]+}} !kernel_arg_accessor_ptr !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
cgh.parallel_for(sycl::range<1>(10), EF);
});

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/usm-int-header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,4 +30,4 @@ int main() {
});
}

// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *)'
// CHECK: FunctionDecl {{.*}}usm_test{{.*}} 'void (__global int *, __global float *) __attribute__((device_kernel))'
2 changes: 1 addition & 1 deletion clang/test/SemaOpenCL/sampled_image_overload.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
void __attribute__((overloadable)) foo(__ocl_sampled_image1d_ro_t);
void __attribute__((overloadable)) foo(__ocl_sampled_image2d_ro_t);

// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)'
// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t) __attribute__((device_kernel))'
void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) {
// CHECK: CallExpr
// CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)'
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ int main() {
}

// Check kernel parameters
// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int)'
// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (int, int, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global char *, sycl::range<1>, sycl::range<1>, sycl::id<1>, int) __attribute__((device_kernel))'
// CHECK: ParmVarDecl{{.*}} used _arg_A 'int'
// CHECK: ParmVarDecl{{.*}} used _arg_B 'int'
// CHECK: ParmVarDecl{{.*}} used _arg_AccField '__global char *'
Expand Down
12 changes: 6 additions & 6 deletions clang/test/SemaSYCL/accessors-targets-image.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -72,12 +72,12 @@ int main() {
});
}

// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)'
// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t)'
// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t)'
// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t)'
// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t)'
// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t)'
// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t) __attribute__((device_kernel))'
// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t) __attribute__((device_kernel))'
// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t) __attribute__((device_kernel))'
// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t) __attribute__((device_kernel))'
// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t) __attribute__((device_kernel))'
// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t) __attribute__((device_kernel))'

// TODO: SYCL specific fail - analyze and enable
// XFAIL: target=x86_64-pc-windows-msvc
8 changes: 4 additions & 4 deletions clang/test/SemaSYCL/accessors-targets.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,7 +51,7 @@ int main() {
});
});
}
// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
// CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
// CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
// CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
18 changes: 9 additions & 9 deletions clang/test/SemaSYCL/array-kernel-param.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,7 +114,7 @@ int main() {
}

// Check Kernel_Accessor parameters
// CHECK: FunctionDecl {{.*}}Kernel_Accessor{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: FunctionDecl {{.*}}Kernel_Accessor{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ReadWriteAccessor '__global int *'
// CHECK-NEXT: SYCLAccessorPtrAttr
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ReadWriteAccessor 'sycl::range<1>'
Expand All @@ -131,7 +131,7 @@ int main() {
// CHECK-NEXT: MemberExpr {{.*}}__init

// Check Kernel_Array parameters
// CHECK: FunctionDecl {{.*}}Kernel_Array{{.*}} 'void (__wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_Array{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_Array '__wrapper_class'
// Check Kernel_Array inits
// CHECK-NEXT: CompoundStmt
Expand All @@ -150,7 +150,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_Array' '__wrapper_class'

// Check Kernel_Array_Ptrs parameters
// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class'
// Check Kernel_Array_Ptrs inits
// CHECK-NEXT: CompoundStmt
Expand Down Expand Up @@ -180,7 +180,7 @@ int main() {
// CHECK-NEXT: IntegerLiteral {{.*}} 1

// Check Kernel_StructAccArray parameters
// CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: FunctionDecl {{.*}}Kernel_StructAccArray{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>, __global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc '__global int *'
// CHECK-NEXT: SYCLAccessorPtrAttr
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_member_acc 'sycl::range<1>'
Expand All @@ -207,7 +207,7 @@ int main() {
// CHECK-NEXT: MemberExpr {{.*}}__init

// Check Kernel_TemplatedStructArray parameters
// CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S<int>)'
// CHECK: FunctionDecl {{.*}}Kernel_TemplatedStructArray{{.*}} 'void (S<int>) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_s 'S<int>'
// CHECK-NEXT: CompoundStmt
// CHECK-NEXT: DeclStmt
Expand All @@ -218,7 +218,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} 'S<int>' lvalue ParmVar {{.*}} '_arg_s' 'S<int>'

// Check Kernel_Array_2D parameters
// CHECK: FunctionDecl {{.*}}Kernel_Array_2D{{.*}} 'void (__wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_Array_2D{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_array_2D '__wrapper_class'
// Check Kernel_Array_2D inits
// CHECK-NEXT: CompoundStmt
Expand Down Expand Up @@ -250,7 +250,7 @@ int main() {
// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned

// Check Kernel_NonDecomposedStruct parameters.
// CHECK: FunctionDecl {{.*}}Kernel_NonDecomposedStruct{{.*}} 'void (__wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_NonDecomposedStruct{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_NonDecompStructArray '__wrapper_class'
// Check Kernel_NonDecomposedStruct inits
// CHECK-NEXT: CompoundStmt
Expand All @@ -271,7 +271,7 @@ int main() {
// CHECK-NEXT: ArrayInitIndexExpr {{.*}} 'unsigned

// Check Kernel_StructWithPointers parameters.
// CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_StructWithPointers{{.*}} 'void (__wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_StructWithPointersArray '__wrapper_class'
// Check Kernel_StructWithPointers inits
// CHECK-NEXT: CompoundStmt
Expand Down Expand Up @@ -303,7 +303,7 @@ int main() {
// CHECK-NEXT: IntegerLiteral {{.*}} 1

// Check Kernel_Array_Ptrs_2D parameters
// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class)'
// CHECK: FunctionDecl {{.*}}Kernel_Array_Ptrs_2D{{.*}} 'void (__wrapper_class, __wrapper_class) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers_2D '__wrapper_class'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_ArrayOfPointers '__wrapper_class'

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@ int main() {

// Check declaration of the kernel

// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
// CHECK: FunctionDecl {{.*}}kernel_wrapper{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>) __attribute__((device_kernel))'

// Check parameters of the kernel

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaSYCL/binding_decl_lambda_nullptr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,5 +18,5 @@ void foo() {
h.single_task<class C>(Lambda);
}

// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int)'
// CHECK: FunctionDecl {{.*}}foo{{.*}} 'void (int) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_bind_x 'int'
10 changes: 5 additions & 5 deletions clang/test/SemaSYCL/built-in-type-kernel-arg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -77,7 +77,7 @@ int main() {
return 0;
}
// Check kernel parameters
// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)'
// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int) __attribute__((device_kernel))'
// CHECK: ParmVarDecl {{.*}} used _arg_some_const 'const int'

// Check that lambda field of const built-in type is initialized
Expand All @@ -87,7 +87,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_some_const' 'const int'

// Check kernel parameters
// CHECK: {{.*}}kernel_int{{.*}} 'void (int)'
// CHECK: {{.*}}kernel_int{{.*}} 'void (int) __attribute__((device_kernel))'
// CHECK: ParmVarDecl {{.*}} used _arg_data 'int'

// Check that lambda field of built-in type is initialized
Expand All @@ -97,7 +97,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_data' 'int'

// Check kernel parameters
// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct)'
// CHECK: {{.*}}kernel_struct{{.*}} 'void (__generated_test_struct) __attribute__((device_kernel))'
// CHECK: ParmVarDecl {{.*}} used _arg_s '__generated_test_struct'

// Check that lambda field of struct type is initialized
Expand All @@ -111,7 +111,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} '__generated_test_struct' lvalue ParmVar {{.*}} '_arg_s'

// Check kernel parameters
// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class)'
// CHECK: {{.*}}kernel_pointer{{.*}} 'void (__global int *, __global int *, __wrapper_class) __attribute__((device_kernel))'
// CHECK: ParmVarDecl {{.*}} used _arg_new_data_addr '__global int *'
// CHECK: ParmVarDecl {{.*}} used _arg_data_addr '__global int *'
// CHECK: ParmVarDecl {{.*}} used _arg_ptr_array '__wrapper_class'
Expand Down Expand Up @@ -145,7 +145,7 @@ int main() {
// CHECK-NEXT: DeclRefExpr {{.*}} '__wrapper_class' lvalue ParmVar {{.*}} '_arg_ptr_array'
// CHECK-NEXT: IntegerLiteral {{.*}} 1

// CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple)'
// CHECK: FunctionDecl {{.*}}kernel_nns{{.*}} 'void (__generated_test_struct_simple) __attribute__((device_kernel))'
// CHECK-NEXT: ParmVarDecl {{.*}} used _arg_tds '__generated_test_struct_simple'

// CHECK: VarDecl {{.*}} used __SYCLKernel
Expand Down
Loading