From abdbf8905d324f9b935b34bbc97c508ede5ac028 Mon Sep 17 00:00:00 2001 From: "Chittireddy, Sindhu" Date: Fri, 16 May 2025 08:51:06 -0700 Subject: [PATCH 01/33] Add sycl_external attribute --- clang/include/clang/Basic/Attr.td | 20 ++++++- clang/include/clang/Basic/AttrDocs.td | 11 ++++ .../clang/Basic/DiagnosticSemaKinds.td | 5 ++ clang/include/clang/Sema/SemaSYCL.h | 1 + clang/lib/AST/ASTContext.cpp | 7 +++ clang/lib/Sema/SemaDeclAttr.cpp | 3 ++ clang/lib/Sema/SemaSYCL.cpp | 11 ++++ .../test/SemaSYCL/sycl-external-attribute.cpp | 52 +++++++++++++++++++ 8 files changed, 109 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaSYCL/sycl-external-attribute.cpp diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index ccd13a4cca4dd..1c13d0eb23f3b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -143,6 +143,7 @@ def SharedVar : SubsetSubjecthasGlobalStorage()}], "global variables">; + def ExternalGlobalVar : SubsetSubjecthasGlobalStorage() && S->getStorageClass()!=StorageClass::SC_Static && @@ -408,10 +409,14 @@ class SubjectList subjects, SubjectDiag diag = WarnDiag, string CustomDiag = customDiag; } -class LangOpt { +class LangOpt { // The language option to test; ignored when custom code is supplied. string Name = name; + // If set to 1, the attribute is accepted but is silently ignored. This is + // useful in multi-compilation situations like SYCL. + bit SilentlyIgnore = silentlyIgnore; + // A custom predicate, written as an expression evaluated in a context with // "LangOpts" bound. code CustomCode = customCode; @@ -422,6 +427,7 @@ def CUDA : LangOpt<"CUDA">; def HIP : LangOpt<"HIP">; def SYCLHost : LangOpt<"SYCLIsHost">; def SYCLDevice : LangOpt<"SYCLIsDevice">; +def SilentlyIgnoreSYCLHost : LangOpt<"SYCLIsHost", "", 1>; def COnly : LangOpt<"", "!LangOpts.CPlusPlus">; def CPlusPlus : LangOpt<"CPlusPlus">; def OpenCL : LangOpt<"OpenCL">; @@ -1545,6 +1551,18 @@ def SYCLKernel : InheritableAttr { let Documentation = [SYCLKernelDocs]; } +def GlobalStorageNonLocalVar : SubsetSubjecthasGlobalStorage() && + !S->isLocalVarDeclOrParm()}], + "global variables">; + +def SYCLExternal : InheritableAttr { + let Spellings = [GNU<"sycl_external">]; + let Subjects = SubjectList<[Function, GlobalStorageNonLocalVar]>; + let LangOpts = [SYCLDevice, SilentlyIgnoreSYCLHost]; + let Documentation = [SYCLExternalDocs]; +} + def SYCLKernelEntryPoint : InheritableAttr { let Spellings = [Clang<"sycl_kernel_entry_point">]; let Args = [ diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 5fb5f16680b41..2eef46a1348f3 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -472,6 +472,17 @@ The SYCL kernel in the previous code sample meets these expectations. }]; } +def SYCLExternalDocs : Documentation { + let Category = DocCatFunction; + let Heading = "sycl_external"; + let Content = [{ +The ``sycl_external`` attribute (or the ``SYCL_EXTERNAL`` macro) can only be applied to +functions, and indicates that the function must be treated as a device function and +must be emitted even if it has no direct uses from other device functions. +All ``sycl_external`` function callees implicitly inherit this attribute. + }]; +} + def SYCLKernelEntryPointDocs : Documentation { let Category = DocCatFunction; let Content = [{ diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3efe9593b8633..9228d388bc10b 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12746,6 +12746,11 @@ def err_sycl_special_type_num_init_method : Error< "types with 'sycl_special_class' attribute must have one and only one '__init' " "method defined">; +//SYCL external attribute diagnostics +def err_sycl_attribute_internal_decl + : Error<"%0 attribute cannot be applied to a %select{function|variable}1" + " without external linkage">; + // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< "'sycl_kernel_entry_point' attribute cannot be applied to a" diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index b47b2f155ef93..099cc56b0ef92 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -62,6 +62,7 @@ class SemaSYCL : public SemaBase { ParsedType ParsedTy); void handleKernelAttr(Decl *D, const ParsedAttr &AL); + void handleSYCLExternalAttr(Decl *D, const ParsedAttr &AL); void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL); void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD); diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index c58cd2c93fb60..e767b79d0b25f 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12909,6 +12909,9 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (D->hasAttr()) return false; + if (LangOpts.SYCLIsDevice && !D->hasAttr()) + return false; + // Aliases and used decls are required. if (D->hasAttr() || D->hasAttr()) return true; @@ -12926,6 +12929,10 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { // FIXME: Functions declared with SYCL_EXTERNAL are required during // device compilation. + // Functions definitions with sycl_external attribute are required during + // device compilation. + if (LangOpts.SYCLIsDevice && FD->hasAttr()) + return true; // Constructors and destructors are required. if (FD->hasAttr() || FD->hasAttr()) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 377595639bef1..271cf417a12f3 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7113,6 +7113,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, case ParsedAttr::AT_SYCLKernel: S.SYCL().handleKernelAttr(D, AL); break; + case ParsedAttr::AT_SYCLExternal: + S.SYCL().handleSYCLExternalAttr(D, AL); + break; case ParsedAttr::AT_SYCLKernelEntryPoint: S.SYCL().handleKernelEntryPointAttr(D, AL); break; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 1969d7b0ba837..a08e2f076fb12 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -202,6 +202,17 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) { handleSimpleAttribute(*this, D, AL); } +void SemaSYCL::handleSYCLExternalAttr(Decl *D, const ParsedAttr &AL) { + auto *ND = cast(D); + if (!ND->isExternallyVisible()) { + Diag(AL.getLoc(), diag::err_sycl_attribute_internal_decl) + << AL << !isa(ND); + return; + } + + handleSimpleAttribute(*this, D, AL); +} + void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) { ParsedType PT = AL.getTypeArg(); TypeSourceInfo *TSI = nullptr; diff --git a/clang/test/SemaSYCL/sycl-external-attribute.cpp b/clang/test/SemaSYCL/sycl-external-attribute.cpp new file mode 100644 index 0000000000000..2260dba386a73 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attribute.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -DSYCL %s +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify -DHOST %s +// RUN: %clang_cc1 -verify %s + +// Semantic tests for sycl_external attribute + +#ifdef SYCL + +__attribute__((sycl_external(3))) // expected-error {{'sycl_external' attribute takes no arguments}} +void bar() {} + +__attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} +static void func1() {} + +namespace { + __attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} + void func2() {} + + struct UnnX {}; +} + +__attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} + void func4(UnnX) {} + +class A { + __attribute__((sycl_external)) + A() {} + + __attribute__((sycl_external)) void func3() {} +}; + +class B { +public: + __attribute__((sycl_external)) virtual void foo() {} + + __attribute__((sycl_external)) virtual void bar() = 0; +}; + +__attribute__((sycl_external)) int *func0() { return nullptr; } + +__attribute__((sycl_external)) void func2(int *) {} + +#elif defined(HOST) + +// expected-no-diagnostics +__attribute__((sycl_external)) void func3() {} + +#else +__attribute__((sycl_external)) // expected-warning {{'sycl_external' attribute ignored}} +void baz() {} + +#endif From f631d7a58e5ce5e80f3562226723410af8b199d6 Mon Sep 17 00:00:00 2001 From: "Chittireddy, Sindhu" Date: Fri, 16 May 2025 13:48:04 -0700 Subject: [PATCH 02/33] Fix test and remove space --- clang/include/clang/Basic/Attr.td | 1 - clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp | 9 ++++----- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 1c13d0eb23f3b..d2fdb1a12c9fb 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -143,7 +143,6 @@ def SharedVar : SubsetSubjecthasGlobalStorage()}], "global variables">; - def ExternalGlobalVar : SubsetSubjecthasGlobalStorage() && S->getStorageClass()!=StorageClass::SC_Static && diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp index b5687523aee36..e41209673c9cc 100644 --- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp +++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp @@ -100,11 +100,10 @@ int main() { // Verify that SYCL kernel caller functions are emitted for each device target. // -// FIXME: The following set of matches are used to skip over the declaration of -// main(). main() shouldn't be emitted in device code, but that pruning isn't -// performed yet. -// CHECK-DEVICE: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone -// CHECK-DEVICE-NEXT: define {{[a-z_ ]*}}noundef i32 @main() #0 +// main() shouldn't be emitted in device code. It is not annotated with +// sycl_kernel_entry_point or sycl_external attributes. +// Function Attrs: convergent mustprogress noinline norecurse nounwind optnone +// CHECK-NOT: define {{[a-z_ ]*}}noundef i32 @main() #0 // IR for the SYCL kernel caller function generated for // single_purpose_kernel_task with single_purpose_kernel_name as the SYCL kernel From 128ab1b1e3030b30d4b157f6657ac1c543fffb4f Mon Sep 17 00:00:00 2001 From: "Chittireddy, Sindhu" Date: Fri, 23 May 2025 10:46:29 -0700 Subject: [PATCH 03/33] Address review comments #1 --- clang/include/clang/Basic/Attr.td | 9 ++------- clang/include/clang/Basic/DiagnosticSemaKinds.td | 5 ++--- clang/include/clang/Sema/SemaSYCL.h | 2 +- clang/lib/AST/ASTContext.cpp | 10 +++++----- clang/lib/Sema/SemaDeclAttr.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 6 files changed, 13 insertions(+), 19 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index d2fdb1a12c9fb..43a7de4257583 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -412,10 +412,6 @@ class LangOpt { // The language option to test; ignored when custom code is supplied. string Name = name; - // If set to 1, the attribute is accepted but is silently ignored. This is - // useful in multi-compilation situations like SYCL. - bit SilentlyIgnore = silentlyIgnore; - // A custom predicate, written as an expression evaluated in a context with // "LangOpts" bound. code CustomCode = customCode; @@ -426,7 +422,6 @@ def CUDA : LangOpt<"CUDA">; def HIP : LangOpt<"HIP">; def SYCLHost : LangOpt<"SYCLIsHost">; def SYCLDevice : LangOpt<"SYCLIsDevice">; -def SilentlyIgnoreSYCLHost : LangOpt<"SYCLIsHost", "", 1>; def COnly : LangOpt<"", "!LangOpts.CPlusPlus">; def CPlusPlus : LangOpt<"CPlusPlus">; def OpenCL : LangOpt<"OpenCL">; @@ -1556,9 +1551,9 @@ def GlobalStorageNonLocalVar : SubsetSubject; def SYCLExternal : InheritableAttr { - let Spellings = [GNU<"sycl_external">]; + let Spellings = [Clang<"sycl_external">]; let Subjects = SubjectList<[Function, GlobalStorageNonLocalVar]>; - let LangOpts = [SYCLDevice, SilentlyIgnoreSYCLHost]; + let LangOpts = [SYCLDevice]; let Documentation = [SYCLExternalDocs]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9228d388bc10b..6ca2e6d6811d1 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12747,9 +12747,8 @@ def err_sycl_special_type_num_init_method : Error< "method defined">; //SYCL external attribute diagnostics -def err_sycl_attribute_internal_decl - : Error<"%0 attribute cannot be applied to a %select{function|variable}1" - " without external linkage">; +def err_sycl_attribute_invalid_linkage : Error< + "'sycl_external' can only be applied to functions with external linkage">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index 099cc56b0ef92..b1c94ee17abb2 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -62,7 +62,7 @@ class SemaSYCL : public SemaBase { ParsedType ParsedTy); void handleKernelAttr(Decl *D, const ParsedAttr &AL); - void handleSYCLExternalAttr(Decl *D, const ParsedAttr &AL); + void handleExternalAttr(Decl *D, const ParsedAttr &AL); void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL); void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD); diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index e767b79d0b25f..7db8a07cad603 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12909,7 +12909,8 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (D->hasAttr()) return false; - if (LangOpts.SYCLIsDevice && !D->hasAttr()) + if (LangOpts.SYCLIsDevice && + (!D->hasAttr() || !D->hasAttr())) return false; // Aliases and used decls are required. @@ -12927,10 +12928,9 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (LangOpts.SYCLIsDevice && FD->hasAttr()) return true; - // FIXME: Functions declared with SYCL_EXTERNAL are required during - // device compilation. - // Functions definitions with sycl_external attribute are required during - // device compilation. + // Functions definitions with the sycl_external attribute are required + // during device compilation regardless of whether they are reachable from + // a SYCL kernel. if (LangOpts.SYCLIsDevice && FD->hasAttr()) return true; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 271cf417a12f3..a6996fe0ce3d1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7114,7 +7114,7 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, S.SYCL().handleKernelAttr(D, AL); break; case ParsedAttr::AT_SYCLExternal: - S.SYCL().handleSYCLExternalAttr(D, AL); + S.SYCL().handleExternalAttr(D, AL); break; case ParsedAttr::AT_SYCLKernelEntryPoint: S.SYCL().handleKernelEntryPointAttr(D, AL); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a08e2f076fb12..c12ed801ccc97 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -202,10 +202,10 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) { handleSimpleAttribute(*this, D, AL); } -void SemaSYCL::handleSYCLExternalAttr(Decl *D, const ParsedAttr &AL) { +void SemaSYCL::handleExternalAttr(Decl *D, const ParsedAttr &AL) { auto *ND = cast(D); if (!ND->isExternallyVisible()) { - Diag(AL.getLoc(), diag::err_sycl_attribute_internal_decl) + Diag(AL.getLoc(), diag::err_sycl_attribute_invalid_linkage) << AL << !isa(ND); return; } From 118656c84eba4bcc608d7c19bb7a41fa0b0fabc8 Mon Sep 17 00:00:00 2001 From: "Chittireddy, Sindhu" Date: Wed, 28 May 2025 12:40:52 -0700 Subject: [PATCH 04/33] Fix conditional and failing tests --- clang/lib/AST/ASTContext.cpp | 2 +- .../CodeGenSYCL/address-space-deduction.cpp | 2 +- .../CodeGenSYCL/address-space-mangling.cpp | 2 +- .../debug-info-kernel-variables.cpp | 2 +- .../CodeGenSYCL/field-annotate-addr-space.cpp | 2 +- .../CodeGenSYCL/functionptr-addrspace.cpp | 2 +- .../test/SemaSYCL/sycl-external-attribute.cpp | 36 ++++++------------- 7 files changed, 17 insertions(+), 31 deletions(-) diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 7db8a07cad603..c02686800b890 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12910,7 +12910,7 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { return false; if (LangOpts.SYCLIsDevice && - (!D->hasAttr() || !D->hasAttr())) + !D->hasAttr() && !D->hasAttr()) return false; // Aliases and used decls are required. diff --git a/clang/test/CodeGenSYCL/address-space-deduction.cpp b/clang/test/CodeGenSYCL/address-space-deduction.cpp index 5910ec3bfc305..0fb5c41d630cc 100644 --- a/clang/test/CodeGenSYCL/address-space-deduction.cpp +++ b/clang/test/CodeGenSYCL/address-space-deduction.cpp @@ -85,7 +85,7 @@ // CHECK-NEXT: store ptr addrspace(4) addrspacecast (ptr addrspace(1) @.str.1 to ptr addrspace(4)), ptr addrspace(4) [[SELECT_STR_TRIVIAL2_ASCAST]], align 8 // CHECK-NEXT: ret void // -void test() { +[[clang::sycl_external]] void test() { static const int foo = 0x42; diff --git a/clang/test/CodeGenSYCL/address-space-mangling.cpp b/clang/test/CodeGenSYCL/address-space-mangling.cpp index 868bf8ccbdcf8..ecc2d4b43a159 100644 --- a/clang/test/CodeGenSYCL/address-space-mangling.cpp +++ b/clang/test/CodeGenSYCL/address-space-mangling.cpp @@ -18,7 +18,7 @@ void foo(int *); // X86: declare void @_Z3fooPU9SYprivatei(ptr noundef) #1 // X86: declare void @_Z3fooPi(ptr noundef) #1 -void test() { +[[clang::sycl_external]] void test() { __attribute__((opencl_global)) int *glob; __attribute__((opencl_local)) int *loc; __attribute__((opencl_private)) int *priv; diff --git a/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp b/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp index 96c0dcfdb75b6..551c4e7e2b8b4 100644 --- a/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp +++ b/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp @@ -18,7 +18,7 @@ KERNEL void parallel_for(const KernelType &KernelFunc) { KernelFunc(); } -void my_kernel(int my_param) { +[[clang::sycl_external]] void my_kernel(int my_param) { int my_local = 0; my_local = my_param; } diff --git a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp index 26bfda8185112..fe7a160900a54 100644 --- a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp +++ b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp @@ -9,7 +9,7 @@ struct HasField { int *a; }; -void foo(int *b) { +[[clang::sycl_external]] void foo(int *b) { struct HasField f; // CHECK: %[[A:.+]] = getelementptr inbounds nuw %struct.HasField, ptr addrspace(4) %{{.+}} // CHECK: %[[CALL:.+]] = call ptr addrspace(4) @llvm.ptr.annotation.p4.p1(ptr addrspace(4) %[[A]], ptr addrspace(1) [[ANNOT]] diff --git a/clang/test/CodeGenSYCL/functionptr-addrspace.cpp b/clang/test/CodeGenSYCL/functionptr-addrspace.cpp index a477b4c7d03ab..060104a29b60a 100644 --- a/clang/test/CodeGenSYCL/functionptr-addrspace.cpp +++ b/clang/test/CodeGenSYCL/functionptr-addrspace.cpp @@ -8,7 +8,7 @@ __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { } // CHECK: define dso_local spir_func{{.*}}invoke_function{{.*}}(ptr noundef %fptr, ptr addrspace(4) noundef %ptr) -void invoke_function(int (*fptr)(), int *ptr) {} +[[clang::sycl_external]] void invoke_function(int (*fptr)(), int *ptr) {} int f() { return 0; } diff --git a/clang/test/SemaSYCL/sycl-external-attribute.cpp b/clang/test/SemaSYCL/sycl-external-attribute.cpp index 2260dba386a73..4bf40c37400ea 100644 --- a/clang/test/SemaSYCL/sycl-external-attribute.cpp +++ b/clang/test/SemaSYCL/sycl-external-attribute.cpp @@ -1,52 +1,38 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify -DSYCL %s -// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify -DHOST %s -// RUN: %clang_cc1 -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s // Semantic tests for sycl_external attribute -#ifdef SYCL - -__attribute__((sycl_external(3))) // expected-error {{'sycl_external' attribute takes no arguments}} +[[clang::sycl_external(3)]] // expected-error {{'sycl_external' attribute takes no arguments}} void bar() {} -__attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} +[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} static void func1() {} namespace { - __attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} + [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} void func2() {} struct UnnX {}; } -__attribute__((sycl_external)) // expected-error {{'sycl_external' attribute cannot be applied to a function without external linkage}} +[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} void func4(UnnX) {} class A { - __attribute__((sycl_external)) + [[clang::sycl_external]] A() {} - __attribute__((sycl_external)) void func3() {} + [[clang::sycl_external]] void func3() {} }; class B { public: - __attribute__((sycl_external)) virtual void foo() {} + [[clang::sycl_external]] virtual void foo() {} - __attribute__((sycl_external)) virtual void bar() = 0; + [[clang::sycl_external]] virtual void bar() = 0; }; -__attribute__((sycl_external)) int *func0() { return nullptr; } - -__attribute__((sycl_external)) void func2(int *) {} - -#elif defined(HOST) - -// expected-no-diagnostics -__attribute__((sycl_external)) void func3() {} +[[clang::sycl_external]] int *func0() { return nullptr; } -#else -__attribute__((sycl_external)) // expected-warning {{'sycl_external' attribute ignored}} -void baz() {} +[[clang::sycl_external]] void func2(int *) {} -#endif From 7c592a42b8814e17f48bac3df29a2883ac70c23f Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 3 Jun 2025 20:48:07 -0700 Subject: [PATCH 05/33] Fix the remaining six failing tests --- .../CodeGenSYCL/address-space-conversions.cpp | 170 +++++++++--------- .../amd-address-space-conversions.cpp | 160 ++++++++--------- .../cuda-address-space-conversions.cpp | 152 ++++++++-------- clang/test/CodeGenSYCL/function-attrs.cpp | 21 +-- clang/test/CodeGenSYCL/unique_stable_name.cpp | 126 +++++++------ .../unique_stable_name_windows_diff.cpp | 38 +++- 6 files changed, 345 insertions(+), 322 deletions(-) diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp index ee3183b74e038..fa7acb0d99433 100644 --- a/clang/test/CodeGenSYCL/address-space-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp @@ -1,143 +1,143 @@ // RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s void bar(int &Data) {} -// CHECK: define{{.*}} spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) % +// CHECK-DAG: define{{.*}} spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) % void bar2(int &Data) {} -// CHECK: define{{.*}} spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) % +// CHECK-DAG: define{{.*}} spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef align 4 dereferenceable(4) % void bar(__attribute__((opencl_local)) int &Data) {} -// CHECK: define{{.*}} spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % +// CHECK-DAG: define{{.*}} spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % void foo(int *Data) {} -// CHECK: define{{.*}} spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef % +// CHECK-DAG: define{{.*}} spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef % void foo2(int *Data) {} -// CHECK: define{{.*}} spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef % +// CHECK-DAG: define{{.*}} spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr addrspace(4) noundef % void foo(__attribute__((opencl_local)) int *Data) {} -// CHECK: define{{.*}} spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % +// CHECK-DAG: define{{.*}} spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % template void tmpl(T t) {} // See Check Lines below. -void usages() { +[[clang::sycl_external]] void usages() { int *NoAS; - // CHECK: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr addrspace(4) + // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr addrspace(4) __attribute__((opencl_global)) int *GLOB; - // CHECK: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1) + // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1) __attribute__((opencl_local)) int *LOC; - // CHECK: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3) + // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3) __attribute__((opencl_private)) int *PRIV; - // CHECK: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr + // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr __attribute__((opencl_global_device)) int *GLOBDEVICE; - // CHECK: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5) + // CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5) __attribute__((opencl_global_host)) int *GLOBHOST; - // CHECK: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(6) + // CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(6) - // CHECK: [[NoAS]].ascast = addrspacecast ptr [[NoAS]] to ptr addrspace(4) - // CHECK: [[GLOB]].ascast = addrspacecast ptr [[GLOB]] to ptr addrspace(4) - // CHECK: [[LOC]].ascast = addrspacecast ptr [[LOC]] to ptr addrspace(4) - // CHECK: [[PRIV]].ascast = addrspacecast ptr [[PRIV]] to ptr addrspace(4) + // CHECK-DAG: [[NoAS]].ascast = addrspacecast ptr [[NoAS]] to ptr addrspace(4) + // CHECK-DAG: [[GLOB]].ascast = addrspacecast ptr [[GLOB]] to ptr addrspace(4) + // CHECK-DAG: [[LOC]].ascast = addrspacecast ptr [[LOC]] to ptr addrspace(4) + // CHECK-DAG: [[PRIV]].ascast = addrspacecast ptr [[PRIV]] to ptr addrspace(4) LOC = nullptr; - // CHECK: store ptr addrspace(3) null, ptr addrspace(4) [[LOC]].ascast, align 8 + // CHECK-DAG: store ptr addrspace(3) null, ptr addrspace(4) [[LOC]].ascast, align 8 GLOB = nullptr; - // CHECK: store ptr addrspace(1) null, ptr addrspace(4) [[GLOB]].ascast, align 8 + // CHECK-DAG: store ptr addrspace(1) null, ptr addrspace(4) [[GLOB]].ascast, align 8 // Explicit conversions // From named address spaces to default address space - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr addrspace(4) - // CHECK: store ptr addrspace(4) [[GLOB_CAST]], ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr addrspace(4) + // CHECK-DAG: store ptr addrspace(4) [[GLOB_CAST]], ptr addrspace(4) [[NoAS]].ascast NoAS = (int *)GLOB; - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr addrspace(4) - // CHECK: store ptr addrspace(4) [[LOC_CAST]], ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr addrspace(4) + // CHECK-DAG: store ptr addrspace(4) [[LOC_CAST]], ptr addrspace(4) [[NoAS]].ascast NoAS = (int *)LOC; - // CHECK: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast - // CHECK: [[PRIV_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[PRIV_LOAD]] to ptr addrspace(4) - // CHECK: store ptr addrspace(4) [[PRIV_CAST]], ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast + // CHECK-DAG: [[PRIV_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[PRIV_LOAD]] to ptr addrspace(4) + // CHECK-DAG: store ptr addrspace(4) [[PRIV_CAST]], ptr addrspace(4) [[NoAS]].ascast NoAS = (int *)PRIV; // From default address space to named address space - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(1) - // CHECK: store ptr addrspace(1) [[NoAS_CAST]], ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(1) + // CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr addrspace(4) [[GLOB]].ascast GLOB = (__attribute__((opencl_global)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(3) - // CHECK: store ptr addrspace(3) [[NoAS_CAST]], ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr addrspace(3) + // CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr addrspace(4) [[LOC]].ascast LOC = (__attribute__((opencl_local)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr - // CHECK: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(4) [[NoAS_LOAD]] to ptr + // CHECK-DAG: store ptr [[NoAS_CAST]], ptr addrspace(4) [[PRIV]].ascast PRIV = (__attribute__((opencl_private)) int *)NoAS; // From opencl_global_[host/device] address spaces to opencl_global - // CHECK: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast - // CHECK: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1) - // CHECK: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr addrspace(4) [[GLOB_DEVICE]].ascast + // CHECK-DAG: [[GLOBDEVICE_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[GLOBDEVICE_LOAD]] to ptr addrspace(1) + // CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_CAST]], ptr addrspace(4) [[GLOB]].ascast GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE; - // CHECK: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(6), ptr addrspace(4) [[GLOB_HOST]].ascast - // CHECK: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(6) [[GLOBHOST_LOAD]] to ptr addrspace(1) - // CHECK: store ptr addrspace(1) [[GLOBHOST_CAST]], ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOBHOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(6), ptr addrspace(4) [[GLOB_HOST]].ascast + // CHECK-DAG: [[GLOBHOST_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(6) [[GLOBHOST_LOAD]] to ptr addrspace(1) + // CHECK-DAG: store ptr addrspace(1) [[GLOBHOST_CAST]], ptr addrspace(4) [[GLOB]].ascast GLOB = (__attribute__((opencl_global)) int *)GLOBHOST; bar(*GLOB); - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_REF]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[GLOB_CAST]]) + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_REF]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[GLOB_CAST]]) bar2(*GLOB); - // CHECK: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: [[GLOB_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD2]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[GLOB_CAST2]]) + // CHECK-DAG: [[GLOB_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOB_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD2]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[GLOB_CAST2]]) bar(*LOC); - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: call spir_func void [[LOC_REF]](ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: call spir_func void [[LOC_REF]](ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) bar2(*LOC); - // CHECK: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[LOC_CAST2]]) + // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[LOC_CAST2]]) bar(*NoAS); - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: call spir_func void @[[RAW_REF]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[NoAS_LOAD]]) + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: call spir_func void @[[RAW_REF]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[NoAS_LOAD]]) bar2(*NoAS); - // CHECK: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[NoAS_LOAD2]]) + // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: call spir_func void @[[RAW_REF2]](ptr addrspace(4) noundef align 4 dereferenceable(4) [[NoAS_LOAD2]]) foo(GLOB); - // CHECK: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_PTR]](ptr addrspace(4) noundef [[GLOB_CAST3]]) + // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_PTR]](ptr addrspace(4) noundef [[GLOB_CAST3]]) foo2(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[GLOB_CAST4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[GLOB_CAST4]]) foo(LOC); - // CHECK: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: call spir_func void [[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) + // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: call spir_func void [[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) foo2(LOC); - // CHECK: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr addrspace(4) - // CHECK: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[LOC_CAST4]]) + // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr addrspace(4) + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[LOC_CAST4]]) foo(NoAS); - // CHECK: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: call spir_func void @[[RAW_PTR]](ptr addrspace(4) noundef [[NoAS_LOAD3]]) + // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: call spir_func void @[[RAW_PTR]](ptr addrspace(4) noundef [[NoAS_LOAD3]]) foo2(NoAS); - // CHECK: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[NoAS_LOAD4]]) + // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: call spir_func void @[[RAW_PTR2]](ptr addrspace(4) noundef [[NoAS_LOAD4]]) // Ensure that we still get 3 different template instantiations. tmpl(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast - // CHECK: call spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast + // CHECK-DAG: call spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) tmpl(LOC); - // CHECK: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast - // CHECK: call spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) + // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast + // CHECK-DAG: call spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) tmpl(PRIV); - // CHECK: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast - // CHECK: call spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef [[PRIV_LOAD5]]) + // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast + // CHECK-DAG: call spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef [[PRIV_LOAD5]]) tmpl(NoAS); - // CHECK: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast - // CHECK: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef [[NoAS_LOAD5]]) + // CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast + // CHECK-DAG: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef [[NoAS_LOAD5]]) } -// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef % -// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef % -// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef % -// CHECK: define linkonce_odr spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef % +// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef % +// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef % +// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef % +// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef % diff --git a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp index d316f22096d3d..17a98195318ad 100644 --- a/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp +++ b/clang/test/CodeGenSYCL/amd-address-space-conversions.cpp @@ -1,128 +1,128 @@ // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s void bar(int &Data) {} -// CHECK: define dso_local void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % void bar2(int &Data) {} -// CHECK: define dso_local void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % void bar(__attribute__((opencl_local)) int &Data) {} -// CHECK: define dso_local void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % void foo(int *Data) {} -// CHECK: define dso_local void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef % +// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef % void foo2(int *Data) {} -// CHECK: define dso_local void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef % +// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef % void foo(__attribute__((opencl_local)) int *Data) {} -// CHECK: define dso_local void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % +// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % template -void tmpl(T t); +void tmpl(T t) {} // See Check Lines below. -void usages() { +[[clang::sycl_external]] void usages() { int *NoAS; - // CHECK: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8, addrspace(5) + // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8, addrspace(5) __attribute__((opencl_global)) int *GLOB; - // CHECK: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) + // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) __attribute__((opencl_local)) int *LOC; - // CHECK: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 4, addrspace(5) + // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 4, addrspace(5) __attribute__((opencl_private)) int *PRIV; - // CHECK: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5), align 4, addrspace(5) + // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr addrspace(5), align 4, addrspace(5) __attribute__((opencl_global_device)) int *GLOBDEVICE; - // CHECK: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) + // CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) __attribute__((opencl_global_host)) int *GLOBHOST; - // CHECK: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) + // CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8, addrspace(5) LOC = nullptr; - // CHECK: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]].ascast, align 4 GLOB = nullptr; - // CHECK: store ptr addrspace(1) null, ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: store ptr addrspace(1) null, ptr [[GLOB]].ascast, align 8 NoAS = (int *)GLOB; - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: store ptr [[GLOB_CAST]], ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: store ptr [[GLOB_CAST]], ptr [[NoAS]].ascast, align 8 NoAS = (int *)LOC; - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr - // CHECK: store ptr [[LOC_CAST]], ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr + // CHECK-DAG: store ptr [[LOC_CAST]], ptr [[NoAS]].ascast, align 8 NoAS = (int *)PRIV; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[NoAS_LOAD]] to ptr - // CHECK: store ptr %5, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(5) [[NoAS_LOAD]] to ptr + // CHECK-DAG: store ptr %5, ptr [[NoAS]].ascast, align 8 GLOB = (__attribute__((opencl_global)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr %6 to ptr addrspace(1) - // CHECK: store ptr addrspace(1) %7, ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr %6 to ptr addrspace(1) + // CHECK-DAG: store ptr addrspace(1) %7, ptr [[GLOB]].ascast, align 8 LOC = (__attribute__((opencl_local)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3) - // CHECK: store ptr addrspace(3) %9, ptr [[LOC]].ascast, align 4 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3) + // CHECK-DAG: store ptr addrspace(3) %9, ptr [[LOC]].ascast, align 4 PRIV = (__attribute__((opencl_private)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(5) - // CHECK: store ptr addrspace(5) [[NoAS_CAST]], ptr [[PRIV]].ascast, align 4 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(5) + // CHECK-DAG: store ptr addrspace(5) [[NoAS_CAST]], ptr [[PRIV]].ascast, align 4 GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]DEVICE.ascast, align 8 - // CHECK: store ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]DEVICE.ascast, align 8 + // CHECK-DAG: store ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8 GLOB = (__attribute__((opencl_global)) int *)GLOBHOST; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]HOST.ascast, align 8 - // CHECK: tore ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]]HOST.ascast, align 8 + // CHECK-DAG: tore ptr addrspace(1) [[NoAS_LOAD]], ptr [[GLOB]].ascast, align 8 bar(*GLOB); - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) bar2(*GLOB); - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) bar(*LOC); - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: call void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: call void @_Z3barRU3AS3i(ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) bar2(*LOC); - // CHECK: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr - // CHECK: call void @_Z4bar2Ri(ptr noundef nonnull align 4 dereferenceable(4) [[LOC_CAST2]]) + // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr + // CHECK-DAG: call void @_Z4bar2Ri(ptr noundef nonnull align 4 dereferenceable(4) [[LOC_CAST2]]) bar(*NoAS); - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: call void @_Z3barRi(ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD]]) + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: call void @_Z3barRi(ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD]]) bar2(*NoAS); - // CHECK: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: call void @_Z4bar2Ri(ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD2]]) + // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: call void @_Z4bar2Ri(ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD2]]) foo(GLOB); - // CHECK: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr - // CHECK: call void @[[RAW_PTR]](ptr noundef [[GLOB_CAST3]]) + // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr + // CHECK-DAG: call void @[[RAW_PTR]](ptr noundef [[GLOB_CAST3]]) foo2(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[GLOB_CAST4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[GLOB_CAST4]]) foo(LOC); - // CHECK: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: call void @[[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) + // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: call void @[[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) foo2(LOC); - // CHECK: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[LOC_CAST4]]) + // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[LOC_CAST4]]) foo(NoAS); - // CHECK: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: call void @[[RAW_PTR]](ptr noundef [[NoAS_LOAD3]]) + // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: call void @[[RAW_PTR]](ptr noundef [[NoAS_LOAD3]]) foo2(NoAS); - // CHECK: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[NoAS_LOAD4]]) + // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[NoAS_LOAD4]]) // Ensure that we still get 3 different template instantiations. tmpl(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 - // CHECK: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8 + // CHECK-DAG: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) tmpl(LOC); - // CHECK: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 - // CHECK: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) + // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4 + // CHECK-DAG: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) tmpl(PRIV); - // CHECK: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4 - // CHECK: call void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef [[PRIV_LOAD5]]) + // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4 + // CHECK-DAG: call void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef [[PRIV_LOAD5]]) tmpl(NoAS); - // CHECK: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 - // CHECK: call void @_Z4tmplIPiEvT_(ptr noundef [[NoAS_LOAD5]]) + // CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8 + // CHECK-DAG: call void @_Z4tmplIPiEvT_(ptr noundef [[NoAS_LOAD5]]) } -// CHECK: declare void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef) -// CHECK: declare void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef) -// CHECK: declare void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef) -// CHECK: declare void @_Z4tmplIPiEvT_(ptr noundef) +// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef % +// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef % +// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef % +// CHECK-DAG: define linkonce_odr void @_Z4tmplIPiEvT_(ptr noundef % diff --git a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp index 1875029de0856..ffb601e62c118 100644 --- a/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp +++ b/clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp @@ -1,122 +1,122 @@ // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s void bar(int &Data) {} -// CHECK: define dso_local void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[RAW_REF:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % void bar2(int &Data) {} -// CHECK: define dso_local void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[RAW_REF2:[a-zA-Z0-9_]+]](ptr noundef nonnull align 4 dereferenceable(4) % void bar(__attribute__((opencl_local)) int &Data) {} -// CHECK: define dso_local void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % +// CHECK-DAG: define {{.*}} void @[[LOCAL_REF:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef align 4 dereferenceable(4) % void foo(int *Data) {} -// CHECK: define dso_local void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef % +// CHECK-DAG: define {{.*}} void @[[RAW_PTR:[a-zA-Z0-9_]+]](ptr noundef % void foo2(int *Data) {} -// CHECK: define dso_local void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef % +// CHECK-DAG: define {{.*}} void @[[RAW_PTR2:[a-zA-Z0-9_]+]](ptr noundef % void foo(__attribute__((opencl_local)) int *Data) {} -// CHECK: define dso_local void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % +// CHECK-DAG: define {{.*}} void @[[LOC_PTR:[a-zA-Z0-9_]+]](ptr addrspace(3) noundef % template void tmpl(T t); // See Check Lines below. -void usages() { +[[clang::sycl_external]] void usages() { int *NoAS; - // CHECK: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8 + // CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca ptr, align 8 __attribute__((opencl_global)) int *GLOB; - // CHECK: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 + // CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 __attribute__((opencl_local)) int *LOC; - // CHECK: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 8 + // CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca ptr addrspace(3), align 8 __attribute__((opencl_private)) int *PRIV; - // CHECK: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr, align 8 + // CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca ptr, align 8 __attribute__((opencl_global_device)) int *GLOBDEVICE; - // CHECK: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 + // CHECK-DAG: [[GLOB_DEVICE:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 __attribute__((opencl_global_host)) int *GLOBHOST; - // CHECK: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 + // CHECK-DAG: [[GLOB_HOST:%[a-zA-Z0-9]+]] = alloca ptr addrspace(1), align 8 LOC = nullptr; - // CHECK: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]], align 8 + // CHECK-DAG: store ptr addrspace(3) addrspacecast (ptr null to ptr addrspace(3)), ptr [[LOC]], align 8 GLOB = nullptr; - // CHECK: store ptr addrspace(1) null, ptr [[GLOB]], align 8 + // CHECK-DAG: store ptr addrspace(1) null, ptr [[GLOB]], align 8 NoAS = (int *)GLOB; - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: store ptr [[GLOB_CAST]], ptr [[NoAS]], align 8 + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: store ptr [[GLOB_CAST]], ptr [[NoAS]], align 8 NoAS = (int *)LOC; - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr - // CHECK: store ptr [[LOC_CAST]], ptr [[NoAS]], align 8 + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: [[LOC_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD]] to ptr + // CHECK-DAG: store ptr [[LOC_CAST]], ptr [[NoAS]], align 8 NoAS = (int *)PRIV; - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8 - // CHECK: store ptr [[LOC_LOAD]], ptr [[NoAS]], align 8 + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8 + // CHECK-DAG: store ptr [[LOC_LOAD]], ptr [[NoAS]], align 8 GLOB = (__attribute__((opencl_global)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(1) - // CHECK: store ptr addrspace(1) [[NoAS_CAST]], ptr [[GLOB]], align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(1) + // CHECK-DAG: store ptr addrspace(1) [[NoAS_CAST]], ptr [[GLOB]], align 8 LOC = (__attribute__((opencl_local)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3) - // CHECK: store ptr addrspace(3) [[NoAS_CAST]], ptr [[LOC]], align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: [[NoAS_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr [[NoAS_LOAD]] to ptr addrspace(3) + // CHECK-DAG: store ptr addrspace(3) [[NoAS_CAST]], ptr [[LOC]], align 8 PRIV = (__attribute__((opencl_private)) int *)NoAS; - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: store ptr [[NoAS_LOAD]], ptr [[PRIV]], align 8 + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: store ptr [[NoAS_LOAD]], ptr [[PRIV]], align 8 GLOB = (__attribute__((opencl_global)) int *)GLOBDEVICE; - // CHECK: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_DEVICE]], align 8 - // CHECK: store ptr addrspace(1) [[GLOBDEVICE_LOAD]], ptr %GLOB, align 8 + // CHECK-DAG: [[GLOBDEVICE_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_DEVICE]], align 8 + // CHECK-DAG: store ptr addrspace(1) [[GLOBDEVICE_LOAD]], ptr %GLOB, align 8 GLOB = (__attribute__((opencl_global)) int *)GLOBHOST; - // CHECK: [[GLOB_HOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_HOST]], align 8 - // CHECK: store ptr addrspace(1) [[GLOB_HOST_LOAD]], ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_HOST_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB_HOST]], align 8 + // CHECK-DAG: store ptr addrspace(1) [[GLOB_HOST_LOAD]], ptr [[GLOB]], align 8 bar(*GLOB); - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) bar2(*GLOB); - // CHECK: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr - // CHECK: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) + // CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD]] to ptr + // CHECK-DAG: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[GLOB_CAST]]) bar(*LOC); - // CHECK: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: call void @[[LOCAL_REF]](ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) + // CHECK-DAG: [[LOC_LOAD:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: call void @[[LOCAL_REF]](ptr addrspace(3) noundef align 4 dereferenceable(4) [[LOC_LOAD]]) bar2(*LOC); - // CHECK: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr - // CHECK: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[LOC_CAST2]]) + // CHECK-DAG: [[LOC_LOAD2:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: [[LOC_CAST2:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD2]] to ptr + // CHECK-DAG: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[LOC_CAST2]]) bar(*NoAS); - // CHECK: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD]]) + // CHECK-DAG: [[NoAS_LOAD:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: call void @[[RAW_REF]](ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD]]) bar2(*NoAS); - // CHECK: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD2]]) + // CHECK-DAG: [[NoAS_LOAD2:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: call void @[[RAW_REF2]](ptr noundef nonnull align 4 dereferenceable(4) [[NoAS_LOAD2]]) foo(GLOB); - // CHECK: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr - // CHECK: call void @[[RAW_PTR]](ptr noundef [[GLOB_CAST3]]) + // CHECK-DAG: [[GLOB_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_CAST3:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD3]] to ptr + // CHECK-DAG: call void @[[RAW_PTR]](ptr noundef [[GLOB_CAST3]]) foo2(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[GLOB_CAST4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: [[GLOB_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(1) [[GLOB_LOAD4]] to ptr + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[GLOB_CAST4]]) foo(LOC); - // CHECK: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: call void @[[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) + // CHECK-DAG: [[LOC_LOAD3:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: call void @[[LOC_PTR]](ptr addrspace(3) noundef [[LOC_LOAD3]]) foo2(LOC); - // CHECK: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[LOC_CAST4]]) + // CHECK-DAG: [[LOC_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: [[LOC_CAST4:%[a-zA-Z0-9]+]] = addrspacecast ptr addrspace(3) [[LOC_LOAD4]] to ptr + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[LOC_CAST4]]) foo(NoAS); - // CHECK: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: call void @[[RAW_PTR]](ptr noundef [[NoAS_LOAD3]]) + // CHECK-DAG: [[NoAS_LOAD3:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: call void @[[RAW_PTR]](ptr noundef [[NoAS_LOAD3]]) foo2(NoAS); - // CHECK: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 - // CHECK: call void @[[RAW_PTR2]](ptr noundef [[NoAS_LOAD4]]) + // CHECK-DAG: [[NoAS_LOAD4:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]], align 8 + // CHECK-DAG: call void @[[RAW_PTR2]](ptr noundef [[NoAS_LOAD4]]) tmpl(GLOB); - // CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 - // CHECK: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) + // CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]], align 8 + // CHECK-DAG: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]]) tmpl(LOC); - // CHECK: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 - // CHECK: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) + // CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]], align 8 + // CHECK-DAG: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]]) tmpl(PRIV); - // CHECK: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8 - // CHECK: call void @_Z4tmplIPiEvT_(ptr noundef [[PRIV_LOAD5]]) + // CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[PRIV]], align 8 + // CHECK-DAG: call void @_Z4tmplIPiEvT_(ptr noundef [[PRIV_LOAD5]]) tmpl(NoAS); -// CHECK: %33 = load ptr, ptr %NoAS, align 8 -// CHECK: call void @_Z4tmplIPiEvT_(ptr noundef %33) +// CHECK-DAG: %33 = load ptr, ptr %NoAS, align 8 +// CHECK-DAG: call void @_Z4tmplIPiEvT_(ptr noundef %33) } -// CHECK: declare void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef) -// CHECK: declare void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef) -// CHECK: declare void @_Z4tmplIPiEvT_(ptr noundef) +// CHECK-DAG: void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef +// CHECK-DAG: void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef +// CHECK-DAG: void @_Z4tmplIPiEvT_(ptr noundef diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp index 83a77a617240a..f70c918c66c7b 100644 --- a/clang/test/CodeGenSYCL/function-attrs.cpp +++ b/clang/test/CodeGenSYCL/function-attrs.cpp @@ -5,11 +5,11 @@ int foo(); // CHECK-LABEL: define dso_local spir_func void @_Z3barv( -// CHECK-SAME: ) #[[ATTR0:[0-9]+]] { +// CHECK-SAME: ) #[[ATTR2:[0-9]+]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) -// CHECK-NEXT: [[CALL:%.*]] = call spir_func noundef i32 @_Z3foov() #[[ATTR1:[0-9]+]] +// CHECK-NEXT: [[CALL:%.*]] = call spir_func noundef i32 @_Z3foov() #[[ATTR3:[0-9]+]] // CHECK-NEXT: store i32 [[CALL]], ptr addrspace(4) [[A_ASCAST]], align 4 // CHECK-NEXT: ret void // @@ -18,7 +18,7 @@ void bar() { } // CHECK-LABEL: define dso_local spir_func noundef i32 @_Z3foov( -// CHECK-SAME: ) #[[ATTR0]] { +// CHECK-SAME: ) #[[ATTR2]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr [[RETVAL]] to ptr addrspace(4) @@ -29,21 +29,10 @@ int foo() { } template -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { +[[clang::sycl_external]] void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } -// CHECK-LABEL: define dso_local noundef i32 @main( -// CHECK-SAME: ) #[[ATTR0]] { -// CHECK-NEXT: entry: -// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 -// CHECK-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON:%.*]], align 1 -// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr [[RETVAL]] to ptr addrspace(4) -// CHECK-NEXT: [[REF_TMP_ASCAST:%.*]] = addrspacecast ptr [[REF_TMP]] to ptr addrspace(4) -// CHECK-NEXT: store i32 0, ptr addrspace(4) [[RETVAL_ASCAST]], align 4 -// CHECK-NEXT: call spir_func void @_Z18kernel_single_taskIZ4mainE11fake_kernelZ4mainEUlvE_EvRKT0_(ptr addrspace(4) noundef align 1 dereferenceable(1) [[REF_TMP_ASCAST]]) #[[ATTR1]] -// CHECK-NEXT: ret i32 0 -// int main() { kernel_single_task([] { bar(); }); return 0; @@ -52,5 +41,5 @@ int main() { // CHECK: attributes #0 = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } // CHECK: attributes #1 = { convergent nounwind } //. -// CHECK: !0 = !{i32 1, !"wchar_size", i32 4} +// CHECK: !{{[0-9]+}} = !{i32 1, !"wchar_size", i32 4} //. diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp index cc9dd61f435d7..3ab7e3b8f2e7a 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp @@ -1,22 +1,22 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -// CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE_E_\00" -// CHECK: @[[INT1:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00" -// CHECK: @[[STRING:[^\w]+]] = private unnamed_addr addrspace(1) constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00", -// CHECK: @[[INT2:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE]] c"_ZTSi\00" -// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00" -// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00" -// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00" -// CHECK: @{{.*}} = private unnamed_addr addrspace(1) constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1 -// CHECK: @{{.*}} = private unnamed_addr addrspace(1) constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE3_\00", align 1 -// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE4_\00" -// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr addrspace(1) constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE5_\00" -// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr addrspace(1) constant [[INT_SIZE]] c"_ZTSi\00" -// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr addrspace(1) constant [[LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00" -// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00", -// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_EvvEUlvE_\00", -// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr addrspace(1) constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00", -// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00", -// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr addrspace(1) constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00", +// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE_E_\00" +// CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00" +// CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00", +// CHECK: @[[INT2:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00" +// CHECK: @[[LAMBDA_X:[^\w]+]] = private unnamed_addr constant [[LAMBDA_X_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00" +// CHECK: @[[MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE0_\00" +// CHECK: @[[MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE1_\00" +// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE2_\00", align 1 +// CHECK: @{{.*}} = private unnamed_addr constant [32 x i8] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE3_\00", align 1 +// CHECK: @[[MACRO_MACRO_X:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE4_\00" +// CHECK: @[[MACRO_MACRO_Y:[^\w]+]] = private unnamed_addr constant [[MACRO_MACRO_SIZE]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE5_\00" +// CHECK: @[[INT3:[^\w]+]] = private unnamed_addr constant [[INT_SIZE]] c"_ZTSi\00" +// CHECK: @[[LAMBDA:[^\w]+]] = private unnamed_addr constant [[LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZZ4mainENKUlvE0_clEvEUlvE_\00" +// CHECK: @[[LAMBDA_IN_DEP_INT:[^\w]+]] = private unnamed_addr constant [[DEP_INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIiEvvEUlvE_\00", +// CHECK: @[[LAMBDA_IN_DEP_X:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_EvvEUlvE_\00", +// CHECK: @[[LAMBDA_NO_DEP:[^\w]+]] = private unnamed_addr constant [[NO_DEP_LAMBDA_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ13lambda_no_depIidEvT_T0_EUlidE_\00", +// CHECK: @[[LAMBDA_TWO_DEP:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA1_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_EvvEUlvE_\00", +// CHECK: @[[LAMBDA_TWO_DEP2:[^\w]+]] = private unnamed_addr constant [[DEP_LAMBDA2_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_EvvEUlvE_\00", extern "C" void puts(const char *) {} @@ -65,95 +65,105 @@ template kernelFunc(); } +template +void unnamed_kernel_single_task(KernelType kernelFunc) { + kernel_single_task(kernelFunc); +} + +template +void not_kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} + int main() { - kernel_single_task(func); - // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) + not_kernel_single_task(func); + // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) auto l1 = []() { return 1; }; auto l2 = [](decltype(l1) *l = nullptr) { return 2; }; - kernel_single_task(l2); + kernel_single_task(l2); puts(__builtin_sycl_unique_stable_name(decltype(l2))); - // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_ - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_KERNEL3]] to ptr addrspace(4))) + // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ + // CHECK: call void @puts(ptr noundef @[[LAMBDA_KERNEL3]]) constexpr const char str[] = "lalala"; static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling"); int i = 0; puts(__builtin_sycl_unique_stable_name(decltype(i++))); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT1]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[INT1]]) // FIXME: Ensure that j is incremented because VLAs are terrible. int j = 55; puts(__builtin_sycl_unique_stable_name(int[++j])); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[STRING]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[STRING]]) - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_ - // CHECK: declare spir_func noundef ptr addrspace(4) @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_ - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_ + // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_ + // CHECK: declare noundef ptr @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_ - kernel_single_task( + unnamed_kernel_single_task( []() { puts(__builtin_sycl_unique_stable_name(int)); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT2]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[INT2]]) auto x = []() {}; puts(__builtin_sycl_unique_stable_name(decltype(x))); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_X]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[LAMBDA_X]]) DEF_IN_MACRO(); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_X]] to ptr addrspace(4))) - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_Y]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[MACRO_X]]) + // CHECK: call void @puts(ptr noundef @[[MACRO_Y]]) MACRO_CALLS_MACRO(); - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_X]] to ptr addrspace(4))) - // CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[MACRO_MACRO_Y]] to ptr addrspace(4))) + // CHECK: call void @puts(ptr noundef @[[MACRO_MACRO_X]]) + // CHECK: call void @puts(ptr noundef @[[MACRO_MACRO_Y]]) template_param(); - // CHECK: call spir_func void @_Z14template_paramIiEvv + // CHECK: call void @_Z14template_paramIiEvv template_param(); - // CHECK: call spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_in_dependent_function(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIiEvv + // CHECK: call void @_Z28lambda_in_dependent_functionIiEvv lambda_in_dependent_function(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_no_dep(3, 5.5); - // CHECK: call spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 noundef 3, double noundef 5.500000e+00) + // CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 noundef 3, double noundef 5.500000e+00) int a = 5; double b = 10.7; auto y = [](int a) { return a; }; auto z = [](double b) { return b; }; lambda_two_dep(); - // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv + // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv lambda_two_dep(); - // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv + // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv }); } -// CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[INT3]] to ptr addrspace(4))) +// CHECK: define linkonce_odr void @_Z14template_paramIiEvv +// CHECK: call void @puts(ptr noundef @[[INT3]]) -// CHECK: define internal spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA]] to ptr addrspace(4))) +// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv +// CHECK: call void @puts(ptr noundef @[[LAMBDA]]) -// CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_IN_DEP_INT]] to ptr addrspace(4))) +// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv +// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_INT]]) -// CHECK: define internal spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_IN_DEP_X]] to ptr addrspace(4))) +// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv +// CHECK: call void @puts(ptr noundef @[[LAMBDA_IN_DEP_X]]) -// CHECK: define linkonce_odr spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 noundef %a, double noundef %b) -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_NO_DEP]] to ptr addrspace(4))) +// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 noundef %a, double noundef %b) +// CHECK: call void @puts(ptr noundef @[[LAMBDA_NO_DEP]]) -// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_TWO_DEP]] to ptr addrspace(4))) +// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv +// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP]]) -// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv -// CHECK: call spir_func void @puts(ptr addrspace(4) noundef addrspacecast (ptr addrspace(1) @[[LAMBDA_TWO_DEP2]] to ptr addrspace(4))) +// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv +// CHECK: call void @puts(ptr noundef @[[LAMBDA_TWO_DEP2]]) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 7dd08a0c89255..4603c5e14e09f 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -1,22 +1,37 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' +// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' // RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' template -__attribute__((sycl_kernel)) void kernel(Func F){ +[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){ F(); } +template +void kernel_wrapper(Func F) { + kernel(F); +} + template -__attribute__((sycl_kernel)) void kernel2(Func F){ +[[clang::sycl_kernel_entry_point(KN)]] void kernel2(Func F){ F(1); } +template +void kernel2_wrapper(Func F) { + kernel2(F); +} + template -__attribute__((sycl_kernel)) void kernel3(Func F){ +[[clang::sycl_kernel_entry_point(KN)]] void kernel3(Func F){ F(1.1); } +template +void kernel3_wrapper(Func F) { + kernel3(F); +} + int main() { int i; double d; @@ -25,15 +40,17 @@ int main() { auto lambda2 = [](int){}; auto lambda3 = [](double){}; - kernel(lambda1); - kernel2(lambda2); - kernel3(lambda3); + kernel_wrapper(lambda1); + kernel2_wrapper(lambda2); + kernel3_wrapper(lambda3); // Ensure the kernels are named the same between the device and host // invocations. + kernel_wrapper([](){ (void)__builtin_sycl_unique_stable_name(decltype(lambda1)); (void)__builtin_sycl_unique_stable_name(decltype(lambda2)); (void)__builtin_sycl_unique_stable_name(decltype(lambda3)); + }); // Make sure the following 3 are the same between the host and device compile. // Note that these are NOT the same value as each other, they differ by the @@ -41,4 +58,11 @@ int main() { // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUlvE_\00" // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUliE_\00" // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUldE_\00" + + // On Windows, ensure that we haven't broken the 'lambda numbering' for thex + // lambda itself. + // WIN: define internal void @"??R Date: Wed, 4 Jun 2025 14:42:30 -0700 Subject: [PATCH 06/33] Fix formatting --- clang/lib/AST/ASTContext.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index c02686800b890..865c7efe5e263 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12909,8 +12909,8 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (D->hasAttr()) return false; - if (LangOpts.SYCLIsDevice && - !D->hasAttr() && !D->hasAttr()) + if (LangOpts.SYCLIsDevice && !D->hasAttr() && + !D->hasAttr()) return false; // Aliases and used decls are required. From a0071d179af4a62677b29153078e31738a674c0f Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 9 Jun 2025 13:17:04 -0700 Subject: [PATCH 07/33] Remove sycl_external attribute support to variables. Support for functions is sufficient for SYCL 2020 spec conformance. --- clang/include/clang/Basic/Attr.td | 9 ++------- 1 file changed, 2 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index ecf22915f7c30..cdef4c14e94b1 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -416,7 +416,7 @@ class SubjectList subjects, SubjectDiag diag = WarnDiag, string CustomDiag = customDiag; } -class LangOpt { +class LangOpt { // The language option to test; ignored when custom code is supplied. string Name = name; @@ -1599,14 +1599,9 @@ def DeviceKernel : DeclOrTypeAttr { }]; } -def GlobalStorageNonLocalVar : SubsetSubjecthasGlobalStorage() && - !S->isLocalVarDeclOrParm()}], - "global variables">; - def SYCLExternal : InheritableAttr { let Spellings = [Clang<"sycl_external">]; - let Subjects = SubjectList<[Function, GlobalStorageNonLocalVar]>; + let Subjects = SubjectList<[Function]>; let LangOpts = [SYCLDevice]; let Documentation = [SYCLExternalDocs]; } From d20382c11a28f6d51c22017452a1299c7fa7c4ea Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 9 Jun 2025 13:52:50 -0700 Subject: [PATCH 08/33] Rename test file --- .../{sycl-external-attribute.cpp => sycl-external-attr.cpp} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename clang/test/SemaSYCL/{sycl-external-attribute.cpp => sycl-external-attr.cpp} (100%) diff --git a/clang/test/SemaSYCL/sycl-external-attribute.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp similarity index 100% rename from clang/test/SemaSYCL/sycl-external-attribute.cpp rename to clang/test/SemaSYCL/sycl-external-attr.cpp From 65262baa19581f1bf23a76244f950629c25c201e Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 9 Jun 2025 13:53:10 -0700 Subject: [PATCH 09/33] Add tests for sycl_external attribute --- .../test/SemaSYCL/sycl-external-attr-grammar.cpp | 16 ++++++++++++++++ .../test/SemaSYCL/sycl-external-attr-ignore.cpp | 12 ++++++++++++ 2 files changed, 28 insertions(+) create mode 100644 clang/test/SemaSYCL/sycl-external-attr-grammar.cpp create mode 100644 clang/test/SemaSYCL/sycl-external-attr-ignore.cpp diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp new file mode 100644 index 0000000000000..f63434bb809fe --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -0,0 +1,16 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// expected-warning@+1{{'sycl_external' attribute only applies to functions}} +[[clang::sycl_external]] int a; + + +// expected-warning@+2{{'sycl_external' attribute only applies to functions}} +struct s { +[[clang::sycl_external]] int b; +}; + +// FIXME: The first declaration of a function is required to have the attribute. +// The attribute may be optionally present on subsequent declarations +int foo(int c); + +[[clang::sycl_external]] void foo(); diff --git a/clang/test/SemaSYCL/sycl-external-attr-ignore.cpp b/clang/test/SemaSYCL/sycl-external-attr-ignore.cpp new file mode 100644 index 0000000000000..d1da8b28417d0 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-ignore.cpp @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -fsyntax-only -verify %s + +// expected-warning@+1{{'sycl_external' attribute ignored}} +[[clang::sycl_external]] void bar() {} + +// expected-warning@+1{{'sycl_external' attribute ignored}} +[[clang::sycl_external]] int a; + +// expected-warning@+2{{'sycl_external' attribute ignored}} +namespace not_sycl { +[[clang::sycl_external]] void foo() {} +} From 770c65e8f5253f3593266348cd063baae40f8a1e Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 10 Jun 2025 09:13:49 -0700 Subject: [PATCH 10/33] Add code examples to sycl_external documentation --- clang/include/clang/Basic/AttrDocs.td | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index ae7ae5ca3b5fd..5d8ca155a7dc6 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -484,6 +484,22 @@ The ``sycl_external`` attribute (or the ``SYCL_EXTERNAL`` macro) can only be app functions, and indicates that the function must be treated as a device function and must be emitted even if it has no direct uses from other device functions. All ``sycl_external`` function callees implicitly inherit this attribute. + +The following examples demonstrate the use of this attribute: + +.. code-block:: c++ +#include + +SYCL_EXTERNAL void Foo(); + +SYCL_EXTERNAL void Bar() { /* ... */ } + +SYCL_EXTERNAL extern void Baz(); + +[[nodiscard]] SYCL_EXTERNAL void Important(); + +SYCL_EXTERNAL [[nodiscard]] void AlsoImportant(); + }]; } From aab6f7df49ed524fd6850ca72abd27a1840f291e Mon Sep 17 00:00:00 2001 From: schittir Date: Tue, 10 Jun 2025 15:21:10 -0400 Subject: [PATCH 11/33] Update clang/lib/Sema/SemaDeclAttr.cpp Co-authored-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaDeclAttr.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 9ead5f7481721..7a265822d85d4 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7210,9 +7210,6 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, case ParsedAttr::AT_EnumExtensibility: handleEnumExtensibilityAttr(S, D, AL); break; - case ParsedAttr::AT_SYCLKernel: - S.SYCL().handleKernelAttr(D, AL); - break; case ParsedAttr::AT_SYCLExternal: S.SYCL().handleExternalAttr(D, AL); break; From 385ea37e83b5126173696d3d88088b4a91693fc7 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 10 Jun 2025 14:09:56 -0700 Subject: [PATCH 12/33] Address review comments -2 --- clang/include/clang/Basic/AttrDocs.td | 18 ++++++++---------- clang/test/CodeGenSYCL/function-attrs.cpp | 2 +- .../unique_stable_name_windows_diff.cpp | 2 +- 3 files changed, 10 insertions(+), 12 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 267ca07c67c85..ddee7afbc217d 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -480,25 +480,23 @@ def SYCLExternalDocs : Documentation { let Category = DocCatFunction; let Heading = "sycl_external"; let Content = [{ -The ``sycl_external`` attribute (or the ``SYCL_EXTERNAL`` macro) can only be applied to -functions, and indicates that the function must be treated as a device function and -must be emitted even if it has no direct uses from other device functions. -All ``sycl_external`` function callees implicitly inherit this attribute. +The ``sycl_external`` attribute can only be applied to external functions, and +indicates that the function must be treated as a device function and must be +emitted even if it has no direct uses from other device functions. The following examples demonstrate the use of this attribute: .. code-block:: c++ -#include -SYCL_EXTERNAL void Foo(); +[[clang::sycl_external]] void Foo(); -SYCL_EXTERNAL void Bar() { /* ... */ } +[[clang::sycl_external]] void Bar() { /* ... */ } -SYCL_EXTERNAL extern void Baz(); +[[clang::sycl_external]] extern void Baz(); -[[nodiscard]] SYCL_EXTERNAL void Important(); +[[nodiscard]] [[clang::sycl_external]] void Important(); -SYCL_EXTERNAL [[nodiscard]] void AlsoImportant(); +[[clang::sycl_external]] [[nodiscard]] void AlsoImportant(); }]; } diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp index f70c918c66c7b..81f893644bc7c 100644 --- a/clang/test/CodeGenSYCL/function-attrs.cpp +++ b/clang/test/CodeGenSYCL/function-attrs.cpp @@ -29,7 +29,7 @@ int foo() { } template -[[clang::sycl_external]] void kernel_single_task(const Func &kernelFunc) { +[[clang::sycl_kernel_entry_point(Name)]] void kernel_single_task(const Func &kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 4603c5e14e09f..37b3a31da4a62 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -59,7 +59,7 @@ int main() { // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUliE_\00" // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUldE_\00" - // On Windows, ensure that we haven't broken the 'lambda numbering' for thex + // On Windows, ensure that we haven't broken the 'lambda numbering' for the // lambda itself. // WIN: define internal void @"??R Date: Tue, 17 Jun 2025 08:09:40 -0700 Subject: [PATCH 13/33] Address review comments -3 --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Basic/AttrDocs.td | 26 +++++++++++++------ .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/AST/ASTContext.cpp | 2 +- 4 files changed, 21 insertions(+), 11 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 0c482fafaf5a7..4a70a281b6b13 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1623,7 +1623,7 @@ def DeviceKernel : DeclOrTypeAttr { def SYCLExternal : InheritableAttr { let Spellings = [Clang<"sycl_external">]; - let Subjects = SubjectList<[Function]>; + let Subjects = SubjectList<[Function], ErrorDiag>; let LangOpts = [SYCLDevice]; let Documentation = [SYCLExternalDocs]; } diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index ddee7afbc217d..86fa43f0842d8 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -480,10 +480,24 @@ def SYCLExternalDocs : Documentation { let Category = DocCatFunction; let Heading = "sycl_external"; let Content = [{ -The ``sycl_external`` attribute can only be applied to external functions, and -indicates that the function must be treated as a device function and must be -emitted even if it has no direct uses from other device functions. - +The ``sycl_external`` attribute indicates that a function defined in another +translation unit may be called by a device function defined in the current +translation unit or, if defined in the current translation unit, the function +may be called by device functions defined in other translation units. +The attribute is intended for use in the implementation of the ``SYCL_EXTERNAL`` +macro as specified in section 5.10.1, "SYCL functions and member functions +linkage", of the SYCL 2020 specification. +The attribute only appertains to functions and only those that meet the +following requirements. +* Has external linkage. +* Is not explicitly defined as deleted (the function may be an explicitly + defaulted function that is defined as deleted). +The attribute shall be present on the first declaration of a function and +may optionally be present on subsequent declarations. +When compiling for a SYCL device target that does not support the generic +address space, the function shall not specify a raw pointer or reference type +as the return type or as a parameter type. +See section 5.9, "Address-space deduction", of the SYCL 2020 specification. The following examples demonstrate the use of this attribute: .. code-block:: c++ @@ -494,10 +508,6 @@ The following examples demonstrate the use of this attribute: [[clang::sycl_external]] extern void Baz(); -[[nodiscard]] [[clang::sycl_external]] void Important(); - -[[clang::sycl_external]] [[nodiscard]] void AlsoImportant(); - }]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index ae15dcf59a941..db4f12528c3e9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12851,7 +12851,7 @@ def err_sycl_special_type_num_init_method : Error< "types with 'sycl_special_class' attribute must have one and only one '__init' " "method defined">; -//SYCL external attribute diagnostics +// SYCL external attribute diagnostics def err_sycl_attribute_invalid_linkage : Error< "'sycl_external' can only be applied to functions with external linkage">; diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index a2aea762035b7..c7517df1669ea 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -12983,7 +12983,7 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (LangOpts.SYCLIsDevice && FD->hasAttr()) return true; - // Functions definitions with the sycl_external attribute are required + // Function definitions with the sycl_external attribute are required // during device compilation regardless of whether they are reachable from // a SYCL kernel. if (LangOpts.SYCLIsDevice && FD->hasAttr()) From 060b24fb85282eb52afe6bd83f6f8e063644b6b6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 17 Jun 2025 08:20:38 -0700 Subject: [PATCH 14/33] Rename test file --- ...cl-external-attr-ignore.cpp => sycl-external-attr-ignored.cpp} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename clang/test/SemaSYCL/{sycl-external-attr-ignore.cpp => sycl-external-attr-ignored.cpp} (100%) diff --git a/clang/test/SemaSYCL/sycl-external-attr-ignore.cpp b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp similarity index 100% rename from clang/test/SemaSYCL/sycl-external-attr-ignore.cpp rename to clang/test/SemaSYCL/sycl-external-attr-ignored.cpp From 625cff21380f777aba6f8aecba479d27466839c8 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 17 Jun 2025 08:56:16 -0700 Subject: [PATCH 15/33] Address review comments -4 --- clang/test/SemaSYCL/sycl-external-attr-grammar.cpp | 7 +++++-- clang/test/SemaSYCL/sycl-external-attr.cpp | 9 ++++++--- 2 files changed, 11 insertions(+), 5 deletions(-) diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp index f63434bb809fe..ae62b50b4700e 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -1,14 +1,17 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s -// expected-warning@+1{{'sycl_external' attribute only applies to functions}} +// expected-error@+1{{'sycl_external' attribute only applies to functions}} [[clang::sycl_external]] int a; -// expected-warning@+2{{'sycl_external' attribute only applies to functions}} +// expected-error@+2{{'sycl_external' attribute only applies to functions}} struct s { [[clang::sycl_external]] int b; }; +// expected-error@+1{{'sycl_external' attribute takes no arguments}} +[[clang::sycl_external(3)]] void bar() {} + // FIXME: The first declaration of a function is required to have the attribute. // The attribute may be optionally present on subsequent declarations int foo(int c); diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 4bf40c37400ea..622425f0b8669 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -2,9 +2,6 @@ // Semantic tests for sycl_external attribute -[[clang::sycl_external(3)]] // expected-error {{'sycl_external' attribute takes no arguments}} -void bar() {} - [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} static void func1() {} @@ -18,6 +15,12 @@ namespace { [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} void func4(UnnX) {} +// FIXME: The first declaration of a function is required to have the attribute. +// The attribute may be optionally present on subsequent declarations +int foo(int c); + +[[clang::sycl_external]] void foo(); + class A { [[clang::sycl_external]] A() {} From 4eb05b858793380242e5755e7c57ae0d3d0756f6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 18 Jun 2025 16:19:08 -0700 Subject: [PATCH 16/33] Fix failing tests and address review comments --- clang/include/clang/Basic/AttrDocs.td | 13 ++++++++++--- clang/lib/Sema/SemaSYCL.cpp | 7 +++---- .../Builtins/generic_cast_to_ptr_explicit.c | 6 +++--- .../test/CodeGenSYCL/kernel-caller-entry-point.cpp | 1 - clang/test/Headers/spirv_functions.cpp | 2 +- .../pragma-attribute-supported-attributes-list.test | 1 + 6 files changed, 18 insertions(+), 12 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index b89c53aa49909..a0a0939e13456 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -487,26 +487,33 @@ may be called by device functions defined in other translation units. The attribute is intended for use in the implementation of the ``SYCL_EXTERNAL`` macro as specified in section 5.10.1, "SYCL functions and member functions linkage", of the SYCL 2020 specification. + The attribute only appertains to functions and only those that meet the following requirements. + * Has external linkage. * Is not explicitly defined as deleted (the function may be an explicitly defaulted function that is defined as deleted). + The attribute shall be present on the first declaration of a function and may optionally be present on subsequent declarations. + When compiling for a SYCL device target that does not support the generic address space, the function shall not specify a raw pointer or reference type as the return type or as a parameter type. See section 5.9, "Address-space deduction", of the SYCL 2020 specification. + The following examples demonstrate the use of this attribute: .. code-block:: c++ -[[clang::sycl_external]] void Foo(); + [[clang::sycl_external]] void Foo(); // Ok. + + [[clang::sycl_external]] void Bar() { /* ... */ } // Ok. -[[clang::sycl_external]] void Bar() { /* ... */ } + [[clang::sycl_external]] extern void Baz(); // Ok. -[[clang::sycl_external]] extern void Baz(); + [[clang::sycl_external]] static void Quux() { /* ... */ } // error: Quux() has internal linkage. }]; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 596cc94086e0e..cd784ccb45a5f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -203,10 +203,9 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) { } void SemaSYCL::handleExternalAttr(Decl *D, const ParsedAttr &AL) { - auto *ND = cast(D); - if (!ND->isExternallyVisible()) { - Diag(AL.getLoc(), diag::err_sycl_attribute_invalid_linkage) - << AL << !isa(ND); + auto *FD = cast(D); + if (!FD->isExternallyVisible()) { + Diag(AL.getLoc(), diag::err_sycl_attribute_invalid_linkage); return; } diff --git a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c index 8cfe650f4db10..bdbc134526a7a 100644 --- a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c +++ b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c @@ -8,7 +8,7 @@ // CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr @llvm.spv.generic.cast.to.ptr.explicit.p0(ptr addrspace(4) %p) // CHECK-NEXT: ret ptr [[SPV_CAST]] // -__attribute__((opencl_private)) int* test_cast_to_private(int* p) { +[[clang::sycl_external]] __attribute__((opencl_private)) int* test_cast_to_private(int* p) { return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7); } @@ -18,7 +18,7 @@ __attribute__((opencl_private)) int* test_cast_to_private(int* p) { // CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(1) @llvm.spv.generic.cast.to.ptr.explicit.p1(ptr addrspace(4) %p) // CHECK-NEXT: ret ptr addrspace(1) [[SPV_CAST]] // -__attribute__((opencl_global)) int* test_cast_to_global(int* p) { +[[clang::sycl_external]] __attribute__((opencl_global)) int* test_cast_to_global(int* p) { return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5); } @@ -28,6 +28,6 @@ __attribute__((opencl_global)) int* test_cast_to_global(int* p) { // CHECK-NEXT: [[SPV_CAST:%.*]] = tail call noundef ptr addrspace(3) @llvm.spv.generic.cast.to.ptr.explicit.p3(ptr addrspace(4) %p) // CHECK-NEXT: ret ptr addrspace(3) [[SPV_CAST]] // -__attribute__((opencl_local)) int* test_cast_to_local(int* p) { +[[clang::sycl_external]] __attribute__((opencl_local)) int* test_cast_to_local(int* p) { return __builtin_spirv_generic_cast_to_ptr_explicit(p, 4); } diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp index e41209673c9cc..903e108ecd6d3 100644 --- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp +++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp @@ -102,7 +102,6 @@ int main() { // // main() shouldn't be emitted in device code. It is not annotated with // sycl_kernel_entry_point or sycl_external attributes. -// Function Attrs: convergent mustprogress noinline norecurse nounwind optnone // CHECK-NOT: define {{[a-z_ ]*}}noundef i32 @main() #0 // IR for the SYCL kernel caller function generated for diff --git a/clang/test/Headers/spirv_functions.cpp b/clang/test/Headers/spirv_functions.cpp index ff036b75faf02..fa9e2cea51079 100644 --- a/clang/test/Headers/spirv_functions.cpp +++ b/clang/test/Headers/spirv_functions.cpp @@ -15,7 +15,7 @@ // NV: call noundef ptr @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePvi // NV: addrspacecast ptr %{{.*}} to ptr addrspace(1) // NV: addrspacecast ptr %{{.*}} to ptr addrspace(3) -void test_cast(int* p) { +[[clang::sycl_external]] void test_cast(int* p) { __spirv_GenericCastToPtrExplicit_ToGlobal(p, 5); __spirv_GenericCastToPtrExplicit_ToLocal(p, 4); __spirv_GenericCastToPtrExplicit_ToPrivate(p, 7); diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index 41d00dae3f69a..65cd43bc0bbd0 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -181,6 +181,7 @@ // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter) // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function) // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function) +// CHECK-NEXT: SYCLExternal (SubjectMatchRule_function) // CHECK-NEXT: SYCLKernelEntryPoint (SubjectMatchRule_function) // CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record) // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record) From ab845a281623cad7e1d7853169fb2f432f3e56e5 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 23 Jun 2025 21:19:01 -0700 Subject: [PATCH 17/33] Address review comments -3 --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 ++++ clang/lib/Sema/SemaSYCL.cpp | 9 +++++++++ clang/test/SemaSYCL/sycl-external-attr.cpp | 13 +++++++++++++ 3 files changed, 26 insertions(+) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 999f11fe95034..00a2871c78620 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12859,6 +12859,10 @@ def err_sycl_special_type_num_init_method : Error< // SYCL external attribute diagnostics def err_sycl_attribute_invalid_linkage : Error< "'sycl_external' can only be applied to functions with external linkage">; +def err_sycl_attribute_avoid_main : Error< + "'sycl_external' cannot be applied to main function">; +def err_sycl_attribute_avoid_deleted_function : Error< + "'sycl_external' cannot be applied to explicitly deleted functions">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cd784ccb45a5f..6b9321233f29b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -208,6 +208,15 @@ void SemaSYCL::handleExternalAttr(Decl *D, const ParsedAttr &AL) { Diag(AL.getLoc(), diag::err_sycl_attribute_invalid_linkage); return; } + std::string FunctionName = StringRef(FD->getNameInfo().getAsString()).lower(); + if (FunctionName.find("main") != std::string::npos) { + Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_main); + return; + } + if (FD->isDeleted()) { + Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_deleted_function); + return; + } handleSimpleAttribute(*this, D, AL); } diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 622425f0b8669..413d6b47d037a 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -21,6 +21,19 @@ int foo(int c); [[clang::sycl_external]] void foo(); +class C { + [[clang::sycl_external]] void member(); +}; + +[[clang::sycl_external]] int main() // expected-error {{'sycl_external' cannot be applied to main function}} +{ + return 0; +} + +class D { + [[clang::sycl_external]] void del() = delete; // expected-error {{'sycl_external' cannot be applied to explicitly deleted functions}} +}; + class A { [[clang::sycl_external]] A() {} From 7e76afdcffc5d788b2105b7776a17f3530b6127f Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Jul 2025 06:57:27 -0700 Subject: [PATCH 18/33] Change the second RUN line to use -sycl-is-host --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 37b3a31da4a62..ce45097d0fe87 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -1,5 +1,5 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' -// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' template From e8d26a2d33bcad2c6f2f04e240b59c6df0d1c1aa Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Jul 2025 08:17:55 -0700 Subject: [PATCH 19/33] Switch to using sycl_external attr to pass the failing test There seems to be an issue with the use of sycl_kernel_entry_point attribute vs __builtin_unique_stable_name --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index ce45097d0fe87..5c0a3732239b3 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -3,7 +3,7 @@ template -[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){ +[[clang::sycl_external]] void kernel(Func F){ F(); } @@ -13,7 +13,7 @@ void kernel_wrapper(Func F) { } template -[[clang::sycl_kernel_entry_point(KN)]] void kernel2(Func F){ +[[clang::sycl_external]] void kernel2(Func F){ F(1); } @@ -23,7 +23,7 @@ void kernel2_wrapper(Func F) { } template -[[clang::sycl_kernel_entry_point(KN)]] void kernel3(Func F){ +[[clang::sycl_external]] void kernel3(Func F){ F(1.1); } From 82fa98a69b96ee4a82ea7b3032d186e275e8f315 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 25 Jun 2025 09:34:55 -0700 Subject: [PATCH 20/33] Change diagnostic messages --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 4 ++-- clang/lib/Sema/SemaSYCL.cpp | 2 +- clang/test/SemaSYCL/sycl-external-attr.cpp | 11 ++++++++--- 3 files changed, 11 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index c7970ad145eb7..7589a0e7f0d05 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12872,9 +12872,9 @@ def err_sycl_special_type_num_init_method : Error< def err_sycl_attribute_invalid_linkage : Error< "'sycl_external' can only be applied to functions with external linkage">; def err_sycl_attribute_avoid_main : Error< - "'sycl_external' cannot be applied to main function">; + "'sycl_external' cannot be applied to the 'main' function">; def err_sycl_attribute_avoid_deleted_function : Error< - "'sycl_external' cannot be applied to explicitly deleted functions">; + "'sycl_external' cannot be applied to an explicitly deleted function">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6b9321233f29b..a5a5af0b88f14 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -213,7 +213,7 @@ void SemaSYCL::handleExternalAttr(Decl *D, const ParsedAttr &AL) { Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_main); return; } - if (FD->isDeleted()) { + if (FD->isDeletedAsWritten()) { Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_deleted_function); return; } diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 413d6b47d037a..8cd37513e350c 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -25,13 +25,18 @@ class C { [[clang::sycl_external]] void member(); }; -[[clang::sycl_external]] int main() // expected-error {{'sycl_external' cannot be applied to main function}} +[[clang::sycl_external]] int main() // expected-error {{'sycl_external' cannot be applied to the 'main' function}} { - return 0; + return 0; } class D { - [[clang::sycl_external]] void del() = delete; // expected-error {{'sycl_external' cannot be applied to explicitly deleted functions}} + [[clang::sycl_external]] void del() = delete; // expected-error {{'sycl_external' cannot be applied to an explicitly deleted function}} +}; + +struct NonCopyable { + ~NonCopyable() = delete; + [[clang::sycl_external]] NonCopyable(const NonCopyable&) = default; }; class A { From e4d15eb4774f237610fe848df424316a214ab845 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Jul 2025 10:34:52 -0700 Subject: [PATCH 21/33] Revert RUN line to -fsycl-is-device --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 5c0a3732239b3..3c89943b40e9e 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -1,5 +1,5 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' -// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' template From b38e5786e1159e3ccc029aa25f536072c637338b Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 3 Jul 2025 10:38:21 -0700 Subject: [PATCH 22/33] Revert test change --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 3c89943b40e9e..d1f5d9b9272b4 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' +// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=addrspace(1) ' // RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s '-D$ADDRSPACE=' From d751b434e25ab1556e71be79213384b89516f460 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 8 Jul 2025 08:20:10 -0700 Subject: [PATCH 23/33] Fix conflict resolution errors. --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index a9b566bb1d0f2..f7849ab9ad39a 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -46,10 +46,10 @@ int main() { // Ensure the kernels are named the same between the device and host // invocations. + // Call from host. (void)__builtin_sycl_unique_stable_name(decltype(lambda1)); (void)__builtin_sycl_unique_stable_name(decltype(lambda2)); (void)__builtin_sycl_unique_stable_name(decltype(lambda3)); - }); // Call from device. auto lambda4 = [](){ From 2b22ed241765aee7bf282d202b1a4c7d7770c7b0 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 8 Jul 2025 09:13:16 -0700 Subject: [PATCH 24/33] Remove changes introduced from downstream. --- clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index f7849ab9ad39a..4fb57bf7af4fd 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -65,11 +65,4 @@ int main() { // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUlvE_\00" // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUliE_\00" // CHECK: private unnamed_addr [[$ADDRSPACE]]constant [17 x i8] c"_ZTSZ4mainEUldE_\00" - - // On Windows, ensure that we haven't broken the 'lambda numbering' for the - // lambda itself. - // WIN: define internal void @"??R Date: Tue, 8 Jul 2025 10:42:37 -0700 Subject: [PATCH 25/33] Update diagnostic messages in tests --- clang/test/SemaSYCL/sycl-external-attr-grammar.cpp | 6 +++--- clang/test/SemaSYCL/sycl-external-attr-ignored.cpp | 6 +++--- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp index ae62b50b4700e..4a6a028dbd988 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -1,15 +1,15 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s -// expected-error@+1{{'sycl_external' attribute only applies to functions}} +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} [[clang::sycl_external]] int a; -// expected-error@+2{{'sycl_external' attribute only applies to functions}} +// expected-error@+2{{'clang::sycl_external' attribute only applies to functions}} struct s { [[clang::sycl_external]] int b; }; -// expected-error@+1{{'sycl_external' attribute takes no arguments}} +// expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} [[clang::sycl_external(3)]] void bar() {} // FIXME: The first declaration of a function is required to have the attribute. diff --git a/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp index d1da8b28417d0..21d44bfe73191 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp @@ -1,12 +1,12 @@ // RUN: %clang_cc1 -fsyntax-only -verify %s -// expected-warning@+1{{'sycl_external' attribute ignored}} +// expected-warning@+1{{'clang::sycl_external' attribute ignored}} [[clang::sycl_external]] void bar() {} -// expected-warning@+1{{'sycl_external' attribute ignored}} +// expected-warning@+1{{'clang::sycl_external' attribute ignored}} [[clang::sycl_external]] int a; -// expected-warning@+2{{'sycl_external' attribute ignored}} +// expected-warning@+2{{'clang::sycl_external' attribute ignored}} namespace not_sycl { [[clang::sycl_external]] void foo() {} } From 568b5699638dc733fd6fa42ad436efe046b77d0f Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 8 Jul 2025 11:42:44 -0700 Subject: [PATCH 26/33] Undo more downstream changes --- .../unique_stable_name_windows_diff.cpp | 21 +++---------------- 1 file changed, 3 insertions(+), 18 deletions(-) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 4fb57bf7af4fd..14366a092a1fe 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -7,31 +7,16 @@ template F(); } -template -void kernel_wrapper(Func F) { - kernel(F); -} - template [[clang::sycl_kernel_entry_point(KN)]] void kernel2(Func F){ F(1); } -template -void kernel2_wrapper(Func F) { - kernel2(F); -} - template [[clang::sycl_kernel_entry_point(KN)]] void kernel3(Func F){ F(1.1); } -template -void kernel3_wrapper(Func F) { - kernel3(F); -} - int main() { int i; double d; @@ -40,9 +25,9 @@ int main() { auto lambda2 = [](int){}; auto lambda3 = [](double){}; - kernel_wrapper(lambda1); - kernel2_wrapper(lambda2); - kernel3_wrapper(lambda3); + kernel(lambda1); + kernel2(lambda2); + kernel3(lambda3); // Ensure the kernels are named the same between the device and host // invocations. From 0ab9ac5cd715b9c1d65972673897c6ffcee6c5da Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 9 Jul 2025 13:10:48 -0700 Subject: [PATCH 27/33] Ungroup diagnostics and add test cases --- .../clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/include/clang/Sema/SemaSYCL.h | 1 + clang/lib/Sema/SemaDecl.cpp | 24 ++++++++++++++ clang/lib/Sema/SemaDeclAttr.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 32 ++++++++----------- clang/test/SemaSYCL/sycl-external-attr.cpp | 14 +++++--- 6 files changed, 51 insertions(+), 24 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 68a9ef2ca28d8..85145317b5ba7 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12885,6 +12885,8 @@ def err_sycl_attribute_avoid_main : Error< "'sycl_external' cannot be applied to the 'main' function">; def err_sycl_attribute_avoid_deleted_function : Error< "'sycl_external' cannot be applied to an explicitly deleted function">; +def err_sycl_attribute_missing_on_first_decl + : Error<"'sycl_external' must be applied to the first declaration">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index b1c94ee17abb2..5b1e17b65b0bc 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -65,6 +65,7 @@ class SemaSYCL : public SemaBase { void handleExternalAttr(Decl *D, const ParsedAttr &AL); void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL); + void CheckSYCLExternalFunctionDecl(FunctionDecl *FD); void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD); StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body); }; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 11cbda412667f..69148531262d3 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -4084,6 +4084,19 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S, diag::note_carries_dependency_missing_first_decl) << 0/*Function*/; } + // SYCL spec 2020 + // The first declaration of a function with external linkage must + // specify sycl_external attribute. + // Subsequent declarations may optionally specify this attribute. + if (LangOpts.SYCLIsDevice) { + const SYCLExternalAttr *SEA = New->getAttr(); + if (SEA && !Old->hasAttr()) { + Diag(SEA->getLocation(), diag::err_sycl_attribute_missing_on_first_decl) + << SEA; + Diag(Old->getLocation(), diag::note_previous_declaration); + } + } + // (C++98 8.3.5p3): // All declarations for a function shall agree exactly in both the // return type and the parameter-type-list. @@ -12251,6 +12264,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD, if (NewFD->hasAttr()) SYCL().CheckSYCLEntryPointFunctionDecl(NewFD); + if (NewFD->hasAttr()) + SYCL().CheckSYCLExternalFunctionDecl(NewFD); + // Semantic checking for this function declaration (in isolation). if (getLangOpts().CPlusPlus) { @@ -12439,6 +12455,14 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) { return; } + if (getLangOpts().SYCLIsDevice) { + if (FD->hasAttr()) { + Diag(FD->getLocation(), diag::err_sycl_attribute_avoid_main); + FD->setInvalidDecl(); + return; + } + } + // Functions named main in hlsl are default entries, but don't have specific // signatures they are required to conform to. if (getLangOpts().HLSL) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 3610733835a62..caeb03595d227 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7256,7 +7256,7 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, handleEnumExtensibilityAttr(S, D, AL); break; case ParsedAttr::AT_SYCLExternal: - S.SYCL().handleExternalAttr(D, AL); + handleSimpleAttribute(S, D, AL); break; case ParsedAttr::AT_SYCLKernelEntryPoint: S.SYCL().handleKernelEntryPointAttr(D, AL); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a5a5af0b88f14..7f4089364d8b2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -202,25 +202,6 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) { handleSimpleAttribute(*this, D, AL); } -void SemaSYCL::handleExternalAttr(Decl *D, const ParsedAttr &AL) { - auto *FD = cast(D); - if (!FD->isExternallyVisible()) { - Diag(AL.getLoc(), diag::err_sycl_attribute_invalid_linkage); - return; - } - std::string FunctionName = StringRef(FD->getNameInfo().getAsString()).lower(); - if (FunctionName.find("main") != std::string::npos) { - Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_main); - return; - } - if (FD->isDeletedAsWritten()) { - Diag(AL.getLoc(), diag::err_sycl_attribute_avoid_deleted_function); - return; - } - - handleSimpleAttribute(*this, D, AL); -} - void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) { ParsedType PT = AL.getTypeArg(); TypeSourceInfo *TSI = nullptr; @@ -268,6 +249,19 @@ static bool CheckSYCLKernelName(Sema &S, SourceLocation Loc, return false; } +void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) { + for (auto *SEAttr : FD->specific_attrs()) { + if (!FD->isExternallyVisible()) { + Diag(SEAttr->getLocation(), diag::err_sycl_attribute_invalid_linkage); + return; + } + if (FD->isDeletedAsWritten()) { + Diag(SEAttr->getLocation(), + diag::err_sycl_attribute_avoid_deleted_function); + return; + } + } +} void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) { // Ensure that all attributes present on the declaration are consistent diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 8cd37513e350c..72148fbad8faa 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -15,11 +15,17 @@ namespace { [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} void func4(UnnX) {} -// FIXME: The first declaration of a function is required to have the attribute. -// The attribute may be optionally present on subsequent declarations -int foo(int c); +// The first declaration of a SYCL external function is required to have this attribute. +int foo(); // expected-note {{previous declaration is here}} -[[clang::sycl_external]] void foo(); +[[clang::sycl_external]] int foo(); // expected-error {{'sycl_external' must be applied to the first declaration}} + +// Subsequent declrations of a SYCL external function may optionally specify this attribute. +[[clang::sycl_external]] int boo(); + +[[clang::sycl_external]] int boo(); // OK + +int boo(); // OK class C { [[clang::sycl_external]] void member(); From a70e2df56acb161c56632a64738c411074d4a4b1 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 10 Jul 2025 05:51:13 -0700 Subject: [PATCH 28/33] Fix newly failing tests by adding sycl_external attribute --- .../CodeGenSPIRV/Builtins/ids_and_ranges.c | 24 +++++++++---------- clang/test/Headers/spirv_ids.cpp | 2 +- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c index f71af779ec358..3508a1ebbcdc7 100644 --- a/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c +++ b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c @@ -7,7 +7,7 @@ // CHECK64-NEXT: tail call i64 @llvm.spv.num.workgroups.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.num.workgroups.i32(i32 0) // -unsigned int test_num_workgroups() { +[[clang::sycl_external]] unsigned int test_num_workgroups() { return __builtin_spirv_num_workgroups(0); } @@ -16,7 +16,7 @@ unsigned int test_num_workgroups() { // CHECK64-NEXT: tail call i64 @llvm.spv.workgroup.size.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.workgroup.size.i32(i32 0) // -unsigned int test_workgroup_size() { +[[clang::sycl_external]] unsigned int test_workgroup_size() { return __builtin_spirv_workgroup_size(0); } @@ -25,7 +25,7 @@ unsigned int test_workgroup_size() { // CHECK64-NEXT: tail call i64 @llvm.spv.group.id.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.group.id.i32(i32 0) // -unsigned int test_workgroup_id() { +[[clang::sycl_external]] unsigned int test_workgroup_id() { return __builtin_spirv_workgroup_id(0); } @@ -34,7 +34,7 @@ unsigned int test_workgroup_id() { // CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.in.group.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.in.group.i32(i32 0) // -unsigned int test_local_invocation_id() { +[[clang::sycl_external]] unsigned int test_local_invocation_id() { return __builtin_spirv_local_invocation_id(0); } @@ -43,7 +43,7 @@ unsigned int test_local_invocation_id() { // CHECK64-NEXT: tail call i64 @llvm.spv.thread.id.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.thread.id.i32(i32 0) // -unsigned int test_global_invocation_id() { +[[clang::sycl_external]] unsigned int test_global_invocation_id() { return __builtin_spirv_global_invocation_id(0); } @@ -52,7 +52,7 @@ unsigned int test_global_invocation_id() { // CHECK64-NEXT: tail call i64 @llvm.spv.global.size.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.global.size.i32(i32 0) // -unsigned int test_global_size() { +[[clang::sycl_external]] unsigned int test_global_size() { return __builtin_spirv_global_size(0); } @@ -61,7 +61,7 @@ unsigned int test_global_size() { // CHECK64-NEXT: tail call i64 @llvm.spv.global.offset.i64(i32 0) // CHECK32-NEXT: tail call i32 @llvm.spv.global.offset.i32(i32 0) // -unsigned int test_global_offset() { +[[clang::sycl_external]] unsigned int test_global_offset() { return __builtin_spirv_global_offset(0); } @@ -69,7 +69,7 @@ unsigned int test_global_offset() { // CHECK-NEXT: [[ENTRY:.*:]] // CHECK-NEXT: tail call i32 @llvm.spv.subgroup.size() // -unsigned int test_subgroup_size() { +[[clang::sycl_external]] unsigned int test_subgroup_size() { return __builtin_spirv_subgroup_size(); } @@ -77,7 +77,7 @@ unsigned int test_subgroup_size() { // CHECK-NEXT: [[ENTRY:.*:]] // CHECK-NEXT: tail call i32 @llvm.spv.subgroup.max.size() // -unsigned int test_subgroup_max_size() { +[[clang::sycl_external]] unsigned int test_subgroup_max_size() { return __builtin_spirv_subgroup_max_size(); } @@ -85,7 +85,7 @@ unsigned int test_subgroup_max_size() { // CHECK-NEXT: [[ENTRY:.*:]] // CHECK-NEXT: tail call i32 @llvm.spv.num.subgroups() // -unsigned int test_num_subgroups() { +[[clang::sycl_external]] unsigned int test_num_subgroups() { return __builtin_spirv_num_subgroups(); } @@ -93,7 +93,7 @@ unsigned int test_num_subgroups() { // CHECK-NEXT: [[ENTRY:.*:]] // CHECK-NEXT: tail call i32 @llvm.spv.subgroup.id() // -unsigned int test_subgroup_id() { +[[clang::sycl_external]] unsigned int test_subgroup_id() { return __builtin_spirv_subgroup_id(); } @@ -101,6 +101,6 @@ unsigned int test_subgroup_id() { // CHECK-NEXT: [[ENTRY:.*:]] // CHECK-NEXT: tail call i32 @llvm.spv.subgroup.local.invocation.id() // -unsigned int test_subgroup_local_invocation_id() { +[[clang::sycl_external]] unsigned int test_subgroup_local_invocation_id() { return __builtin_spirv_subgroup_local_invocation_id(); } diff --git a/clang/test/Headers/spirv_ids.cpp b/clang/test/Headers/spirv_ids.cpp index 0cd74dbca53aa..03bf2f802a1d3 100644 --- a/clang/test/Headers/spirv_ids.cpp +++ b/clang/test/Headers/spirv_ids.cpp @@ -80,7 +80,7 @@ // NV: call noundef i32 @_Z18__spirv_SubgroupIdv() #2 // NV: call noundef i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #2 -void test_id_and_range() { +[[clang::sycl_external]] void test_id_and_range() { __spirv_NumWorkgroups(0); __spirv_NumWorkgroups(1); __spirv_NumWorkgroups(2); From 45f7b09b282750d5f3dde549961ad7b5fc0434a5 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 10 Jul 2025 13:56:06 -0700 Subject: [PATCH 29/33] Add constexpr and consteval test cases --- clang/test/SemaSYCL/sycl-external-attr.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 72148fbad8faa..8aefb82eca92d 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -1,5 +1,5 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s - +// RUN: %clang_cc1 -fsycl-is-device -std=c++20 -fsyntax-only -verify -DCPP20 %s // Semantic tests for sycl_external attribute [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} @@ -63,3 +63,8 @@ class B { [[clang::sycl_external]] void func2(int *) {} +[[clang::sycl_external]] constexpr int square(int x); + +#ifdef CPP20 +[[clang::sycl_external]] consteval int func(); +#endif From 4db4101101d3e3f5a3264818d8a0b674135116c9 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 11 Jul 2025 14:27:13 -0700 Subject: [PATCH 30/33] Use existing diagnostic and address other minor comments --- clang/include/clang/Basic/Attr.td | 2 +- clang/include/clang/Basic/DiagnosticSemaKinds.td | 7 +++---- clang/lib/Sema/SemaDecl.cpp | 12 +++++++----- clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp | 3 +-- clang/test/SemaSYCL/sycl-external-attr.cpp | 2 +- 5 files changed, 13 insertions(+), 13 deletions(-) diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 7904186e7ff20..57d119e560b04 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1644,7 +1644,7 @@ def DeviceKernel : DeclOrTypeAttr { def SYCLExternal : InheritableAttr { let Spellings = [Clang<"sycl_external">]; let Subjects = SubjectList<[Function], ErrorDiag>; - let LangOpts = [SYCLDevice]; + let LangOpts = [SYCLHost, SYCLDevice]; let Documentation = [SYCLExternalDocs]; } diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 13457656afbf6..4c724bf3a5796 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12884,10 +12884,9 @@ def err_sycl_attribute_invalid_linkage : Error< "'sycl_external' can only be applied to functions with external linkage">; def err_sycl_attribute_avoid_main : Error< "'sycl_external' cannot be applied to the 'main' function">; -def err_sycl_attribute_avoid_deleted_function : Error< - "'sycl_external' cannot be applied to an explicitly deleted function">; -def err_sycl_attribute_missing_on_first_decl - : Error<"'sycl_external' must be applied to the first declaration">; +def err_sycl_attribute_avoid_deleted_function + : Error<"'sycl_external' cannot be applied to an explicitly deleted " + "function">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 69148531262d3..6e5041927344e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -4084,16 +4084,18 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S, diag::note_carries_dependency_missing_first_decl) << 0/*Function*/; } - // SYCL spec 2020 - // The first declaration of a function with external linkage must - // specify sycl_external attribute. - // Subsequent declarations may optionally specify this attribute. + // SYCL 2020 section 5.10.1, "SYCL functions and member functions linkage": + // When a function is declared with sycl_external, that attribute must be + // used on the first declaration of that function in the translation unit. + // Redeclarations of the function in the same translation unit may + // optionally use sycl_external, but this is not required. if (LangOpts.SYCLIsDevice) { const SYCLExternalAttr *SEA = New->getAttr(); if (SEA && !Old->hasAttr()) { - Diag(SEA->getLocation(), diag::err_sycl_attribute_missing_on_first_decl) + Diag(SEA->getLocation(), diag::err_attribute_missing_on_first_decl) << SEA; Diag(Old->getLocation(), diag::note_previous_declaration); + New->dropAttr(); } } diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp index 903e108ecd6d3..6b308966d8b6a 100644 --- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp +++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp @@ -100,8 +100,7 @@ int main() { // Verify that SYCL kernel caller functions are emitted for each device target. // -// main() shouldn't be emitted in device code. It is not annotated with -// sycl_kernel_entry_point or sycl_external attributes. +// main() shouldn't be emitted in device code. // CHECK-NOT: define {{[a-z_ ]*}}noundef i32 @main() #0 // IR for the SYCL kernel caller function generated for diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 8aefb82eca92d..00815e6b38372 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -18,7 +18,7 @@ namespace { // The first declaration of a SYCL external function is required to have this attribute. int foo(); // expected-note {{previous declaration is here}} -[[clang::sycl_external]] int foo(); // expected-error {{'sycl_external' must be applied to the first declaration}} +[[clang::sycl_external]] int foo(); // expected-error {{'clang::sycl_external' attribute does not appear on the first declaration}} // Subsequent declrations of a SYCL external function may optionally specify this attribute. [[clang::sycl_external]] int boo(); From 931fd76deaae79cc5cd74427df77d7c7a2eb169e Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 14 Jul 2025 10:34:24 -0700 Subject: [PATCH 31/33] Add additional test cases and address review comments --- .../clang/Basic/DiagnosticSemaKinds.td | 7 +- clang/lib/Sema/SemaDecl.cpp | 6 +- clang/lib/Sema/SemaSYCL.cpp | 2 +- .../SemaSYCL/sycl-external-attr-grammar.cpp | 21 +++--- .../SemaSYCL/sycl-external-attr-ignored.cpp | 9 ++- clang/test/SemaSYCL/sycl-external-attr.cpp | 75 ++++++++++++++----- 6 files changed, 77 insertions(+), 43 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 4c724bf3a5796..6c371740c34f9 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12882,11 +12882,10 @@ def err_sycl_special_type_num_init_method : Error< // SYCL external attribute diagnostics def err_sycl_attribute_invalid_linkage : Error< "'sycl_external' can only be applied to functions with external linkage">; -def err_sycl_attribute_avoid_main : Error< +def err_sycl_attribute_invalid_main : Error< "'sycl_external' cannot be applied to the 'main' function">; -def err_sycl_attribute_avoid_deleted_function - : Error<"'sycl_external' cannot be applied to an explicitly deleted " - "function">; +def err_sycl_attribute_invalid_deleted_function : Error< + "'sycl_external' cannot be applied to an explicitly deleted function">; // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 6e5041927344e..ef74d03a08b4e 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -4085,10 +4085,10 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S, } // SYCL 2020 section 5.10.1, "SYCL functions and member functions linkage": - // When a function is declared with sycl_external, that attribute must be + // When a function is declared with SYCL_EXTERNAL, that macro must be // used on the first declaration of that function in the translation unit. // Redeclarations of the function in the same translation unit may - // optionally use sycl_external, but this is not required. + // optionally use SYCL_EXTERNAL, but this is not required. if (LangOpts.SYCLIsDevice) { const SYCLExternalAttr *SEA = New->getAttr(); if (SEA && !Old->hasAttr()) { @@ -12459,7 +12459,7 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) { if (getLangOpts().SYCLIsDevice) { if (FD->hasAttr()) { - Diag(FD->getLocation(), diag::err_sycl_attribute_avoid_main); + Diag(FD->getLocation(), diag::err_sycl_attribute_invalid_main); FD->setInvalidDecl(); return; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7f4089364d8b2..ca690957ebcbc 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -257,7 +257,7 @@ void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) { } if (FD->isDeletedAsWritten()) { Diag(SEAttr->getLocation(), - diag::err_sycl_attribute_avoid_deleted_function); + diag::err_sycl_attribute_invalid_deleted_function); return; } } diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp index 4a6a028dbd988..c37165498577a 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -1,19 +1,16 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s -// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} -[[clang::sycl_external]] int a; +// expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} +[[clang::sycl_external(3)]] void bar() {} +// FIXME: this case should be diagnosed too +[[clang::sycl_external()]] void bad1(); -// expected-error@+2{{'clang::sycl_external' attribute only applies to functions}} -struct s { -[[clang::sycl_external]] int b; -}; +// expected-error@+1{{expected expression}} +[[clang::sycl_external(,)]] void bad2(); // expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} -[[clang::sycl_external(3)]] void bar() {} - -// FIXME: The first declaration of a function is required to have the attribute. -// The attribute may be optionally present on subsequent declarations -int foo(int c); +[[clang::sycl_external(3)]] void bad3(); -[[clang::sycl_external]] void foo(); +// expected-error@+1{{expected expression}} +[[clang::sycl_external(4,)]] void bad4(); diff --git a/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp index 21d44bfe73191..211b3f1989f11 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp @@ -1,5 +1,8 @@ // RUN: %clang_cc1 -fsyntax-only -verify %s +// These tests validate that the sycl_external attribute is ignored when SYCL +// support is not enabled. + // expected-warning@+1{{'clang::sycl_external' attribute ignored}} [[clang::sycl_external]] void bar() {} @@ -7,6 +10,6 @@ [[clang::sycl_external]] int a; // expected-warning@+2{{'clang::sycl_external' attribute ignored}} -namespace not_sycl { -[[clang::sycl_external]] void foo() {} -} +template +[[clang::sycl_external]] void ft(T) {} +template void ft(int); diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 00815e6b38372..508fcdb826848 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -1,45 +1,73 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -fsyntax-only -verify -DCPP17 %s // RUN: %clang_cc1 -fsycl-is-device -std=c++20 -fsyntax-only -verify -DCPP20 %s -// Semantic tests for sycl_external attribute -[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} +// Semantic tests for the sycl_external attribute. + +// expected-error@+1{{'sycl_external' can only be applied to functions with external linkage}} +[[clang::sycl_external]] static void func1() {} +// expected-error@+2{{'sycl_external' can only be applied to functions with external linkage}} namespace { - [[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} + [[clang::sycl_external]] void func2() {} struct UnnX {}; } -[[clang::sycl_external]] // expected-error {{'sycl_external' can only be applied to functions with external linkage}} - void func4(UnnX) {} +// expected-error@+2{{'sycl_external' can only be applied to functions with external linkage}} +namespace { struct S4 {}; } +[[clang::sycl_external]] void func4(S4) {} + +// FIXME: This case is currently being diagnosed as an error because clang implements +// default inheritance of attribute and explicit instantiation declaration names the +// symbol that causes the instantiated specialization to have internal linkage. +// expected-error@+3{{'sycl_external' can only be applied to functions with external linkage}} +namespace { struct S6 {}; } +template +[[clang::sycl_external]] void func6() {} +template void func6(); +// expected-note@-1{{in instantiation of function template specialization 'func6<(anonymous namespace)::S6>' requested here}} + +// expected-error@+3 2{{'sycl_external' can only be applied to functions with external linkage}} +namespace { struct S7 {}; } +template +[[clang::sycl_external]] void func7(); +template<> void func7() {} +// expected-note@-1{{in instantiation of function template specialization 'func7<(anonymous namespace)::S7>' requested here}} + +namespace { struct S8 {}; } +template +void func8(); +template<> [[clang::sycl_external]] void func8() {} +// expected-error@-1{{'clang::sycl_external' attribute does not appear on the first declaration}} +// expected-note@-2{{previous declaration is here}} // The first declaration of a SYCL external function is required to have this attribute. -int foo(); // expected-note {{previous declaration is here}} - -[[clang::sycl_external]] int foo(); // expected-error {{'clang::sycl_external' attribute does not appear on the first declaration}} +// expected-note@+1{{previous declaration is here}} +int foo(); +// expected-error@+1{{'clang::sycl_external' attribute does not appear on the first declaration}} +[[clang::sycl_external]] int foo(); // Subsequent declrations of a SYCL external function may optionally specify this attribute. [[clang::sycl_external]] int boo(); - [[clang::sycl_external]] int boo(); // OK - int boo(); // OK class C { [[clang::sycl_external]] void member(); }; -[[clang::sycl_external]] int main() // expected-error {{'sycl_external' cannot be applied to the 'main' function}} +// expected-error@+1{{'sycl_external' cannot be applied to the 'main' function}} +[[clang::sycl_external]] int main() { return 0; } +// expected-error@+2{{'sycl_external' cannot be applied to an explicitly deleted function}} class D { - [[clang::sycl_external]] void del() = delete; // expected-error {{'sycl_external' cannot be applied to an explicitly deleted function}} + [[clang::sycl_external]] void del() = delete; }; - struct NonCopyable { ~NonCopyable() = delete; [[clang::sycl_external]] NonCopyable(const NonCopyable&) = default; @@ -49,7 +77,8 @@ class A { [[clang::sycl_external]] A() {} - [[clang::sycl_external]] void func3() {} + [[clang::sycl_external]] void mf() {} + [[clang::sycl_external]] static void smf(); }; class B { @@ -59,12 +88,18 @@ class B { [[clang::sycl_external]] virtual void bar() = 0; }; -[[clang::sycl_external]] int *func0() { return nullptr; } - -[[clang::sycl_external]] void func2(int *) {} - [[clang::sycl_external]] constexpr int square(int x); -#ifdef CPP20 +// Devices that do not support the generic address space shall not specify +// a raw pointer or reference type as the return type or as a parameter type. +[[clang::sycl_external]] int *fun0(); +[[clang::sycl_external]] int &fun1(); +[[clang::sycl_external]] int &&fun2(); +[[clang::sycl_external]] void fun3(int *); +[[clang::sycl_external]] void fun4(int &); +[[clang::sycl_external]] void fun5(int &&); + +#if CPP20 [[clang::sycl_external]] consteval int func(); #endif + From e34f2a6a01eb5cdb62a05bbb54bd55bbbf5980e2 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 16 Jul 2025 11:27:16 -0700 Subject: [PATCH 32/33] Add test cases --- clang/include/clang/Basic/AttrDocs.td | 3 +- clang/lib/Sema/SemaDecl.cpp | 7 ++++ .../sycl-external-attr-appertainment.cpp | 33 +++++++++++++++++++ .../SemaSYCL/sycl-external-attr-grammar.cpp | 5 ++- clang/test/SemaSYCL/sycl-external-attr.cpp | 11 +++++-- 5 files changed, 52 insertions(+), 7 deletions(-) create mode 100644 clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 467ff504885da..7785018f9a344 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -501,8 +501,7 @@ may optionally be present on subsequent declarations. When compiling for a SYCL device target that does not support the generic address space, the function shall not specify a raw pointer or reference type as the return type or as a parameter type. -See section 5.9, "Address-space deduction", of the SYCL 2020 specification. - +See section 5.10, "SYCL offline linking", of the SYCL 2020 specification. The following examples demonstrate the use of this attribute: .. code-block:: c++ diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index ef74d03a08b4e..ef105b98c1370 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -16301,6 +16301,13 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, } } + if (FD && !FD->isInvalidDecl() && FD->hasAttr()) { + SYCLExternalAttr *SEAttr = FD->getAttr(); + if (FD->isDeletedAsWritten()) + Diag(SEAttr->getLocation(), + diag::err_sycl_attribute_invalid_deleted_function); + } + { // Do not call PopExpressionEvaluationContext() if it is a lambda because // one is already popped when finishing the lambda in BuildLambdaExpr(). diff --git a/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp b/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp new file mode 100644 index 0000000000000..45a1c90ccfe64 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++17 -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++20 -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++23 -verify %s + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +[[clang::sycl_external]] int bad1; + + +// expected-error@+2{{'clang::sycl_external' attribute only applies to functions}} +struct s { +[[clang::sycl_external]] int bad2; +}; + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +namespace [[clang::sycl_external]] bad3 {} + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +struct [[clang::sycl_external]] bad4; + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +enum [[clang::sycl_external]] bad5 {}; + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +int bad6(void (fp [[clang::sycl_external]])()); + +// expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} +[[clang::sycl_external]]; + +#if __cplusplus >= 202002L +// expected-error@+2{{'clang::sycl_external' attribute only applies to functions}} +template +concept bad8 [[clang::sycl_external]] = true; +#endif diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp index c37165498577a..3c0037dcdc644 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -1,9 +1,8 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s -// expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} -[[clang::sycl_external(3)]] void bar() {} -// FIXME: this case should be diagnosed too +// FIXME: this case should be diagnosed. +// This attribute takes no arguments. [[clang::sycl_external()]] void bad1(); // expected-error@+1{{expected expression}} diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 508fcdb826848..0fe5685481fc3 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -66,8 +66,12 @@ class C { // expected-error@+2{{'sycl_external' cannot be applied to an explicitly deleted function}} class D { - [[clang::sycl_external]] void del() = delete; + [[clang::sycl_external]] void mdel() = delete; }; + +// expected-error@+1{{'sycl_external' cannot be applied to an explicitly deleted function}} +[[clang::sycl_external]] void del() = delete; + struct NonCopyable { ~NonCopyable() = delete; [[clang::sycl_external]] NonCopyable(const NonCopyable&) = default; @@ -98,8 +102,11 @@ class B { [[clang::sycl_external]] void fun3(int *); [[clang::sycl_external]] void fun4(int &); [[clang::sycl_external]] void fun5(int &&); +template +[[clang::sycl_external]] void fun6(T) {} +template void fun6(int *); +template<> [[clang::sycl_external]] void fun6(long *) {} #if CPP20 [[clang::sycl_external]] consteval int func(); #endif - From 13a68d58a5e0a1ad681480a2f969cf0b6ca2c550 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Fri, 18 Jul 2025 08:33:30 -0700 Subject: [PATCH 33/33] Add FIXME comments, enable diagnostics for host, remove a needless decl --- clang/include/clang/Sema/SemaSYCL.h | 1 - clang/lib/Sema/SemaDecl.cpp | 24 ++++++++----------- .../sycl-external-attr-appertainment.cpp | 3 +++ .../SemaSYCL/sycl-external-attr-grammar.cpp | 5 ++-- clang/test/SemaSYCL/sycl-external-attr.cpp | 17 +++++++++---- 5 files changed, 27 insertions(+), 23 deletions(-) diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index 5b1e17b65b0bc..7ae556da2bec1 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -62,7 +62,6 @@ class SemaSYCL : public SemaBase { ParsedType ParsedTy); void handleKernelAttr(Decl *D, const ParsedAttr &AL); - void handleExternalAttr(Decl *D, const ParsedAttr &AL); void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL); void CheckSYCLExternalFunctionDecl(FunctionDecl *FD); diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index ef105b98c1370..cf95907ca7257 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -4089,14 +4089,12 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S, // used on the first declaration of that function in the translation unit. // Redeclarations of the function in the same translation unit may // optionally use SYCL_EXTERNAL, but this is not required. - if (LangOpts.SYCLIsDevice) { - const SYCLExternalAttr *SEA = New->getAttr(); - if (SEA && !Old->hasAttr()) { - Diag(SEA->getLocation(), diag::err_attribute_missing_on_first_decl) - << SEA; - Diag(Old->getLocation(), diag::note_previous_declaration); - New->dropAttr(); - } + const SYCLExternalAttr *SEA = New->getAttr(); + if (SEA && !Old->hasAttr()) { + Diag(SEA->getLocation(), diag::err_attribute_missing_on_first_decl) + << SEA; + Diag(Old->getLocation(), diag::note_previous_declaration); + New->dropAttr(); } // (C++98 8.3.5p3): @@ -12457,12 +12455,10 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) { return; } - if (getLangOpts().SYCLIsDevice) { - if (FD->hasAttr()) { - Diag(FD->getLocation(), diag::err_sycl_attribute_invalid_main); - FD->setInvalidDecl(); - return; - } + if (FD->hasAttr()) { + Diag(FD->getLocation(), diag::err_sycl_attribute_invalid_main); + FD->setInvalidDecl(); + return; } // Functions named main in hlsl are default entries, but don't have specific diff --git a/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp b/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp index 45a1c90ccfe64..d06c9c9d53580 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp @@ -1,5 +1,8 @@ +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -std=c++17 -verify %s // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++17 -verify %s +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -std=c++20 -verify %s // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++20 -verify %s +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -std=c++23 -verify %s // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -std=c++23 -verify %s // expected-error@+1{{'clang::sycl_external' attribute only applies to functions}} diff --git a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp index 3c0037dcdc644..a0169851cdaf0 100644 --- a/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -1,8 +1,7 @@ +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s - -// FIXME: this case should be diagnosed. -// This attribute takes no arguments. +// FIXME-expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} [[clang::sycl_external()]] void bad1(); // expected-error@+1{{expected expression}} diff --git a/clang/test/SemaSYCL/sycl-external-attr.cpp b/clang/test/SemaSYCL/sycl-external-attr.cpp index 0fe5685481fc3..a42af8c7e4a52 100644 --- a/clang/test/SemaSYCL/sycl-external-attr.cpp +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -1,4 +1,6 @@ +// RUN: %clang_cc1 -fsycl-is-host -std=c++17 -fsyntax-only -verify -DCPP17 %s // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -fsyntax-only -verify -DCPP17 %s +// RUN: %clang_cc1 -fsycl-is-host -std=c++20 -fsyntax-only -verify -DCPP20 %s // RUN: %clang_cc1 -fsycl-is-device -std=c++20 -fsyntax-only -verify -DCPP20 %s // Semantic tests for the sycl_external attribute. @@ -11,17 +13,12 @@ static void func1() {} namespace { [[clang::sycl_external]] void func2() {} - - struct UnnX {}; } // expected-error@+2{{'sycl_external' can only be applied to functions with external linkage}} namespace { struct S4 {}; } [[clang::sycl_external]] void func4(S4) {} -// FIXME: This case is currently being diagnosed as an error because clang implements -// default inheritance of attribute and explicit instantiation declaration names the -// symbol that causes the instantiated specialization to have internal linkage. // expected-error@+3{{'sycl_external' can only be applied to functions with external linkage}} namespace { struct S6 {}; } template @@ -29,6 +26,12 @@ template template void func6(); // expected-note@-1{{in instantiation of function template specialization 'func6<(anonymous namespace)::S6>' requested here}} +// FIXME: C++23 [temp.expl.spec]p12 states: +// ... Similarly, attributes appearing in the declaration of a template +// have no effect on an explicit specialization of that template. +// Clang currently instantiates and propagates attributes from a function +// template to its explicit specializations resulting in the following +// spurious error. // expected-error@+3 2{{'sycl_external' can only be applied to functions with external linkage}} namespace { struct S7 {}; } template @@ -36,6 +39,10 @@ template template<> void func7() {} // expected-note@-1{{in instantiation of function template specialization 'func7<(anonymous namespace)::S7>' requested here}} +// FIXME: The explicit function template specialization appears to trigger +// instantiation of a declaration from the primary template without the +// attribute leading to a spurious diagnostic that the sycl_external +// attribute is not present on the first declaration. namespace { struct S8 {}; } template void func8();