Author: schittir Date: 2025-08-20T12:37:37-04:00 New Revision: fdfcebb38d18fa328d51794e4f8b6914dc87759b
URL: https://github.com/llvm/llvm-project/commit/fdfcebb38d18fa328d51794e4f8b6914dc87759b DIFF: https://github.com/llvm/llvm-project/commit/fdfcebb38d18fa328d51794e4f8b6914dc87759b.diff LOG: [clang][SYCL] Add sycl_external attribute and restrict emitting device code (#140282) This patch is part of the upstreaming effort for supporting SYCL language front end. It makes the following changes: 1. Adds sycl_external attribute for functions with external linkage, which is intended for use to implement the SYCL_EXTERNAL macro as specified by the SYCL 2020 specification 2. Adds checks to avoid emitting device code when sycl_external and sycl_kernel_entry_point attributes are not enabled 3. Fixes test failures caused by the above changes This patch is missing diagnostics for the following diagnostics listed in the SYCL 2020 specification's section 5.10.1, which will be addressed in a subsequent PR: Functions that are declared using SYCL_EXTERNAL have the following additional restrictions beyond those imposed on other device functions: 1. If the SYCL backend does not support the generic address space then the function cannot use raw pointers as parameter or return types. Explicit pointer classes must be used instead; 2. The function cannot call group::parallel_for_work_item; 3. The function cannot be called from a parallel_for_work_group scope. In addition to that, the subsequent PR will also implement diagnostics for inline functions including virtual functions defined as inline. --------- Co-authored-by: Mariya Podchishchaeva <mariya.podchishcha...@intel.com> Added: clang/test/CodeGenSYCL/sycl-external-attr.cpp clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp clang/test/SemaSYCL/sycl-external-attr-grammar.cpp clang/test/SemaSYCL/sycl-external-attr-ignored.cpp clang/test/SemaSYCL/sycl-external-attr.cpp Modified: clang/include/clang/Basic/Attr.td clang/include/clang/Basic/AttrDocs.td clang/include/clang/Basic/DiagnosticGroups.td clang/include/clang/Basic/DiagnosticSemaKinds.td clang/include/clang/Sema/SemaSYCL.h clang/lib/AST/ASTContext.cpp clang/lib/Sema/SemaDecl.cpp clang/lib/Sema/SemaDeclAttr.cpp clang/lib/Sema/SemaSYCL.cpp clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c clang/test/CodeGenSYCL/address-space-conversions.cpp clang/test/CodeGenSYCL/address-space-deduction.cpp clang/test/CodeGenSYCL/address-space-mangling.cpp clang/test/CodeGenSYCL/amd-address-space-conversions.cpp clang/test/CodeGenSYCL/cuda-address-space-conversions.cpp clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp clang/test/CodeGenSYCL/field-annotate-addr-space.cpp clang/test/CodeGenSYCL/function-attrs.cpp clang/test/CodeGenSYCL/functionptr-addrspace.cpp clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp clang/test/CodeGenSYCL/unique_stable_name.cpp clang/test/Headers/spirv_functions.cpp clang/test/Headers/spirv_ids.cpp clang/test/Misc/pragma-attribute-supported-attributes-list.test Removed: ################################################################################ diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 8c8e0b3bca46c..29364c5903d31 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1632,6 +1632,13 @@ def DeviceKernel : DeclOrTypeAttr { }]; } +def SYCLExternal : InheritableAttr { + let Spellings = [CXX11<"clang", "sycl_external">]; + let Subjects = SubjectList<[Function], ErrorDiag>; + let LangOpts = [SYCLHost, SYCLDevice]; + let Documentation = [SYCLExternalDocs]; +} + def SYCLKernelEntryPoint : InheritableAttr { let Spellings = [CXX11<"clang", "sycl_kernel_entry_point">]; let Args = [ diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 00e8fc0787884..b9405c5fc86c1 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -476,6 +476,47 @@ 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 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.10, "SYCL offline linking", of the SYCL 2020 specification. +The following examples demonstrate the use of this attribute: + +.. code-block:: c++ + + [[clang::sycl_external]] void Foo(); // Ok. + + [[clang::sycl_external]] void Bar() { /* ... */ } // Ok. + + [[clang::sycl_external]] extern void Baz(); // Ok. + + [[clang::sycl_external]] static void Quux() { /* ... */ } // error: Quux() has internal linkage. + + }]; +} + def SYCLKernelEntryPointDocs : Documentation { let Category = DocCatFunction; let Content = [{ diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td index e29c4694fa5ea..0c994e0b5ca4d 100644 --- a/clang/include/clang/Basic/DiagnosticGroups.td +++ b/clang/include/clang/Basic/DiagnosticGroups.td @@ -652,6 +652,7 @@ def NonNull : DiagGroup<"nonnull">; def NonPODVarargs : DiagGroup<"non-pod-varargs">; def ClassVarargs : DiagGroup<"class-varargs", [NonPODVarargs]>; def : DiagGroup<"nonportable-cfstrings">; +def NonPortableSYCL : DiagGroup<"nonportable-sycl">; def NonVirtualDtor : DiagGroup<"non-virtual-dtor">; def GNUNullPointerArithmetic : DiagGroup<"gnu-null-pointer-arithmetic">; def NullPointerArithmetic diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index d8200516c7a33..51005bf09fda5 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -12958,6 +12958,17 @@ 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_external_invalid_linkage : Error< + "%0 can only be applied to functions with external linkage">; +def err_sycl_external_invalid_main : Error< + "%0 cannot be applied to the 'main' function">; +def err_sycl_external_invalid_deleted_function : Error< + "%0 cannot be applied to an explicitly deleted function">; +def warn_sycl_external_missing_on_first_decl : Warning< + "%0 attribute does not appear on the first declaration">, + InGroup<NonPortableSYCL>; + // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< "the %0 attribute cannot be applied to a" @@ -12972,7 +12983,7 @@ def err_sycl_kernel_name_conflict : Error< "the %0 kernel name argument conflicts with a previous declaration">; def warn_sycl_kernel_name_not_a_class_type : Warning< "%0 is not a valid SYCL kernel name type; a non-union class type is required">, - InGroup<DiagGroup<"nonportable-sycl">>, DefaultError; + InGroup<NonPortableSYCL>, DefaultError; def warn_sycl_entry_point_redundant_declaration : Warning< "redundant %0 attribute">, InGroup<RedundantAttribute>; def err_sycl_entry_point_after_definition : Error< diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h index b47b2f155ef93..7ae556da2bec1 100644 --- a/clang/include/clang/Sema/SemaSYCL.h +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -64,6 +64,7 @@ class SemaSYCL : public SemaBase { void handleKernelAttr(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/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 059d2a6b1da7d..bbb957067c4c8 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -13127,6 +13127,14 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (D->hasAttr<WeakRefAttr>()) return false; + // SYCL device compilation requires that functions defined with the + // sycl_kernel_entry_point or sycl_external attributes be emitted. All + // other entities are emitted only if they are used by a function + // defined with one of those attributes. + if (LangOpts.SYCLIsDevice) + return isa<FunctionDecl>(D) && (D->hasAttr<SYCLKernelEntryPointAttr>() || + D->hasAttr<SYCLExternalAttr>()); + // Aliases and used decls are required. if (D->hasAttr<AliasAttr>() || D->hasAttr<UsedAttr>()) return true; @@ -13136,15 +13144,6 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (!FD->doesThisDeclarationHaveABody()) return FD->doesDeclarationForceExternallyVisibleDefinition(); - // Function definitions with the sycl_kernel_entry_point attribute are - // required during device compilation so that SYCL kernel caller offload - // entry points are emitted. - if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelEntryPointAttr>()) - return true; - - // FIXME: Functions declared with SYCL_EXTERNAL are required during - // device compilation. - // Constructors and destructors are required. if (FD->hasAttr<ConstructorAttr>() || FD->hasAttr<DestructorAttr>()) return true; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 64352ab29be5d..160d7353cacd9 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -3115,6 +3115,10 @@ static void checkNewAttributesAfterDef(Sema &S, Decl *New, const Decl *Old) { cast<SYCLKernelEntryPointAttr>(NewAttribute)->setInvalidAttr(); ++I; continue; + } else if (isa<SYCLExternalAttr>(NewAttribute)) { + // SYCLExternalAttr may be added after a definition. + ++I; + continue; } S.Diag(NewAttribute->getLocation(), @@ -4140,6 +4144,18 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD, Scope *S, diag::note_carries_dependency_missing_first_decl) << 0/*Function*/; } + // SYCL 2020 section 5.10.1, "SYCL functions and member functions linkage": + // 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. + const SYCLExternalAttr *SEA = New->getAttr<SYCLExternalAttr>(); + if (SEA && !Old->hasAttr<SYCLExternalAttr>()) { + Diag(SEA->getLocation(), diag::warn_sycl_external_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. @@ -12325,6 +12341,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD, if (NewFD->hasAttr<SYCLKernelEntryPointAttr>()) SYCL().CheckSYCLEntryPointFunctionDecl(NewFD); + if (NewFD->hasAttr<SYCLExternalAttr>()) + SYCL().CheckSYCLExternalFunctionDecl(NewFD); + // Semantic checking for this function declaration (in isolation). if (getLangOpts().CPlusPlus) { @@ -12513,6 +12532,13 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) { return; } + if (FD->hasAttr<SYCLExternalAttr>()) { + Diag(FD->getLocation(), diag::err_sycl_external_invalid_main) + << FD->getAttr<SYCLExternalAttr>(); + 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) @@ -16351,6 +16377,14 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation, } } + if (FD && !FD->isInvalidDecl() && FD->hasAttr<SYCLExternalAttr>()) { + SYCLExternalAttr *SEAttr = FD->getAttr<SYCLExternalAttr>(); + if (FD->isDeletedAsWritten()) + Diag(SEAttr->getLocation(), + diag::err_sycl_external_invalid_deleted_function) + << SEAttr; + } + { // Do not call PopExpressionEvaluationContext() if it is a lambda because // one is already popped when finishing the lambda in BuildLambdaExpr(). diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 7a185106e4c6e..625e380f561e2 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7061,6 +7061,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL, case ParsedAttr::AT_EnumExtensibility: handleEnumExtensibilityAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLExternal: + handleSimpleAttribute<SYCLExternalAttr>(S, 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 dd25336539816..2f97f6290f0e8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -250,6 +250,23 @@ static bool CheckSYCLKernelName(Sema &S, SourceLocation Loc, return false; } +void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) { + const auto *SEAttr = FD->getAttr<SYCLExternalAttr>(); + assert(SEAttr && "Missing sycl_external attribute"); + if (!FD->isInvalidDecl() && !FD->isTemplated()) { + if (!FD->isExternallyVisible()) + if (!FD->isFunctionTemplateSpecialization() || + FD->getTemplateSpecializationInfo()->isExplicitSpecialization()) + Diag(SEAttr->getLocation(), diag::err_sycl_external_invalid_linkage) + << SEAttr; + } + if (FD->isDeletedAsWritten()) { + Diag(SEAttr->getLocation(), + diag::err_sycl_external_invalid_deleted_function) + << SEAttr; + } +} + void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) { // Ensure that all attributes present on the declaration are consistent // and warn about any redundant ones. 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..30f4ecb589a5c 100644 --- a/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c +++ b/clang/test/CodeGenSPIRV/Builtins/generic_cast_to_ptr_explicit.c @@ -1,33 +1,36 @@ -// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device %s -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device -x c++ %s -emit-llvm -o - | FileCheck %s // RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s // RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s -// CHECK: spir_func noundef ptr @test_cast_to_private( -// CHECK-SAME: ptr addrspace(4) noundef readnone [[P:%.*]] +#ifdef __SYCL_DEVICE_ONLY__ +#define SYCL_EXTERNAL [[clang::sycl_external]] +#else +#define SYCL_EXTERNAL +#endif + +// CHECK: spir_func noundef ptr @{{.*}}test_cast_to_private{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]] // CHECK-NEXT: [[ENTRY:.*:]] // 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) { +SYCL_EXTERNAL __attribute__((opencl_private)) int* test_cast_to_private(int* p) { return __builtin_spirv_generic_cast_to_ptr_explicit(p, 7); } -// CHECK: spir_func noundef ptr addrspace(1) @test_cast_to_global( -// CHECK-SAME: ptr addrspace(4) noundef readnone [[P:%.*]] +// CHECK: spir_func noundef ptr addrspace(1) @{{.*}}test_cast_to_global{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]] // CHECK-NEXT: [[ENTRY:.*:]] // 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) { +SYCL_EXTERNAL __attribute__((opencl_global)) int* test_cast_to_global(int* p) { return __builtin_spirv_generic_cast_to_ptr_explicit(p, 5); } -// CHECK: spir_func noundef ptr addrspace(3) @test_cast_to_local( -// CHECK-SAME: ptr addrspace(4) noundef readnone [[P:%.*]] +// CHECK: spir_func noundef ptr addrspace(3) @{{.*}}test_cast_to_local{{.*}}(ptr addrspace(4) noundef readnone [[P:%.*]] // CHECK-NEXT: [[ENTRY:.*:]] // 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) { +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/CodeGenSPIRV/Builtins/ids_and_ranges.c b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c index f71af779ec358..bff850b3622b9 100644 --- a/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c +++ b/clang/test/CodeGenSPIRV/Builtins/ids_and_ranges.c @@ -1,106 +1,106 @@ -// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64 +// RUN: %clang_cc1 -O1 -triple spirv64 -fsycl-is-device -x c++ %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64 // RUN: %clang_cc1 -O1 -triple spirv64 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK64 // RUN: %clang_cc1 -O1 -triple spirv32 -cl-std=CL3.0 -x cl %s -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK32 -// CHECK: @test_num_workgroups( +// CHECK: @{{.*}}test_num_workgroups{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_workgroup_size( +// CHECK: @{{.*}}test_workgroup_size{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_workgroup_id( +// CHECK: @{{.*}}test_workgroup_id{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_local_invocation_id( +// CHECK: @{{.*}}test_local_invocation_id{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_global_invocation_id( +// CHECK: @{{.*}}test_global_invocation_id{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_global_size( +// CHECK: @{{.*}}test_global_size{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_global_offset( +// CHECK: @{{.*}}test_global_offset{{.*}}( // CHECK-NEXT: [[ENTRY:.*:]] // 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); } -// CHECK: @test_subgroup_size( +// CHECK: @{{.*}}test_subgroup_size{{.*}}( // 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(); } -// CHECK: @test_subgroup_max_size( +// CHECK: @{{.*}}test_subgroup_max_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(); } -// CHECK: @test_num_subgroups( +// CHECK: @{{.*}}test_num_subgroups{{.*}}( // 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(); } -// CHECK: @test_subgroup_id( +// CHECK: @{{.*}}test_subgroup_id{{.*}}( // 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(); } -// CHECK: @test_subgroup_local_invocation_id( +// CHECK: @{{.*}}test_subgroup_local_invocation_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/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 <typename 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 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 diff erent 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/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/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 <typename T> -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 diff erent 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 <typename 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 + // 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/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/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp index 83a77a617240a..81f893644bc7c 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 <typename Name, typename Func> -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { +[[clang::sycl_kernel_entry_point(Name)]] 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<class fake_kernel>([] { 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/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/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp index b5687523aee36..cd1d4d801951d 100644 --- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp +++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp @@ -100,11 +100,8 @@ 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. +// CHECK-NOT: @main() // IR for the SYCL kernel caller function generated for // single_purpose_kernel_task with single_purpose_kernel_name as the SYCL kernel diff --git a/clang/test/CodeGenSYCL/sycl-external-attr.cpp b/clang/test/CodeGenSYCL/sycl-external-attr.cpp new file mode 100644 index 0000000000000..2c22e65874713 --- /dev/null +++ b/clang/test/CodeGenSYCL/sycl-external-attr.cpp @@ -0,0 +1,85 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s + +// This test code generation when sycl_external attribute is used + +// Function defined and not used - symbols emitted +[[clang::sycl_external]] int square(int x) { return x*x; } +// CHECK: define dso_local spir_func noundef i32 @_Z6squarei + +// Function defined and used - symbols emitted +[[clang::sycl_external]] int squareUsed(int x) { return x*x; } +// CHECK: define dso_local spir_func noundef i32 @_Z10squareUsedi + +// FIXME: Constexpr function defined and not used - symbols emitted +[[clang::sycl_external]] constexpr int squareInlined(int x) { return x*x; } +// CHECK: define linkonce_odr spir_func noundef i32 @_Z13squareInlinedi + +// Function declared but not defined or used - no symbols emitted +[[clang::sycl_external]] int declOnly(); +// CHECK-NOT: define {{.*}} i32 @_Z8declOnlyv +// CHECK-NOT: declare {{.*}} i32 @_Z8declOnlyv + +// Function declared and used in host but not defined - no symbols emitted +[[clang::sycl_external]] void declUsedInHost(int y); + +// Function declared and used in device but not defined - emit external reference +[[clang::sycl_external]] void declUsedInDevice(int y); +// CHECK: define dso_local spir_func void @_Z9deviceUsev +[[clang::sycl_external]] void deviceUse() { declUsedInDevice(3); } +// CHECK: declare spir_func void @_Z16declUsedInDevicei + +// Function declared with the attribute and later defined - definition emitted +[[clang::sycl_external]] int func1(int arg); +int func1(int arg) { return arg; } +// CHECK: define dso_local spir_func noundef i32 @_Z5func1i + +class A { +// Unused defaulted special member functions - no symbols emitted + [[clang::sycl_external]] A& operator=(A& a) = default; +}; + +class B { + [[clang::sycl_external]] virtual void BFunc1WithAttr() { int i = 1; } +// CHECK: define linkonce_odr spir_func void @_ZN1B14BFunc1WithAttrEv + virtual void BFunc2NoAttr() { int i = 2; } +}; + +class C { +// Special member function defined - definition emitted + [[clang::sycl_external]] ~C() {} +// CHECK: define linkonce_odr spir_func void @_ZN1CD1Ev +}; + +// Function reachable from an unused function - definition emitted +int ret1() { return 1; } +[[clang::sycl_external]] int withAttr() { return ret1(); } +// CHECK: define dso_local spir_func noundef i32 @_Z8withAttrv +// CHECK: define dso_local spir_func noundef i32 @_Z4ret1v + +template <typename T> +[[clang::sycl_external]] void tFunc1(T arg) {} +// Explicit specialization defined - symbols emitted +template<> +[[clang::sycl_external]] void tFunc1<int>(int arg) {} +// CHECK: define dso_local spir_func void @_Z6tFunc1IiEvT_ + +template <typename T> +[[clang::sycl_external]] void tFunc2(T arg) {} +template void tFunc2<int>(int arg); +// CHECK: define weak_odr spir_func void @_Z6tFunc2IiEvT_ +template<> void tFunc2<char>(char arg) {} +// CHECK: define dso_local spir_func void @_Z6tFunc2IcEvT_ +template<> [[clang::sycl_external]] void tFunc2<long>(long arg) {} +// CHECK: define dso_local spir_func void @_Z6tFunc2IlEvT_ + +// Functions defined without the sycl_external attribute that are used +// in host code, but not in device code are not emitted. +int squareNoAttr(int x) { return x*x; } +// CHECK-NOT: define {{.*}} i32 @_Z12squareNoAttri + +int main() { + declUsedInHost(4); + int i = squareUsed(5); + int j = squareNoAttr(6); + return 0; +} 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 <typename KernelName, typename KernelType> kernelFunc(); } +template<typename KernelType> +void unnamed_kernel_single_task(KernelType kernelFunc) { + kernel_single_task<KernelType>(kernelFunc); +} + +template <typename KernelName, typename KernelType> +void not_kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} + int main() { - kernel_single_task<class kernel2>(func<Derp>); - // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(ptr noundef @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) + not_kernel_single_task<class kernel2>(func<Derp>); + // 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<class kernel3>(l2); + kernel_single_task<decltype(l2)>(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<class kernel>( + 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<int>(); - // CHECK: call spir_func void @_Z14template_paramIiEvv + // CHECK: call void @_Z14template_paramIiEvv template_param<decltype(x)>(); - // CHECK: call spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_in_dependent_function<int>(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIiEvv + // CHECK: call void @_Z28lambda_in_dependent_functionIiEvv lambda_in_dependent_function<decltype(x)>(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_no_dep<int, double>(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<decltype(y), decltype(z)>(); - // 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<decltype(z), decltype(y)>(); - // 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/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/Headers/spirv_ids.cpp b/clang/test/Headers/spirv_ids.cpp index 466be5deee87a..a1c953eff7771 100644 --- a/clang/test/Headers/spirv_ids.cpp +++ b/clang/test/Headers/spirv_ids.cpp @@ -80,7 +80,7 @@ // NV: call noundef i32 @_Z25__spirv_BuiltInSubgroupIdv() #2 // NV: call noundef i32 @_Z40__spirv_BuiltInSubgroupLocalInvocationIdv() #2 -void test_id_and_range() { +[[clang::sycl_external]] void test_id_and_range() { __spirv_BuiltInNumWorkgroups(0); __spirv_BuiltInNumWorkgroups(1); __spirv_BuiltInNumWorkgroups(2); diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test index b9cf7cf9462fe..37ff33e5a1523 100644 --- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test +++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test @@ -182,6 +182,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) 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..d06c9c9d53580 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp @@ -0,0 +1,36 @@ +// 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}} +[[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<typename> +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 new file mode 100644 index 0000000000000..a0169851cdaf0 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-grammar.cpp @@ -0,0 +1,14 @@ +// RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s + +// FIXME-expected-error@+1{{'clang::sycl_external' attribute takes no arguments}} +[[clang::sycl_external()]] void bad1(); + +// 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 bad3(); + +// 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 new file mode 100644 index 0000000000000..211b3f1989f11 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr-ignored.cpp @@ -0,0 +1,15 @@ +// 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() {} + +// expected-warning@+1{{'clang::sycl_external' attribute ignored}} +[[clang::sycl_external]] int a; + +// expected-warning@+2{{'clang::sycl_external' attribute ignored}} +template<typename T> +[[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 new file mode 100644 index 0000000000000..ebda94e7d5030 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-external-attr.cpp @@ -0,0 +1,154 @@ +// RUN: %clang_cc1 -fsycl-is-host -std=c++17 -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -fsyntax-only -verify %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. + +// expected-error@+1{{'clang::sycl_external' can only be applied to functions with external linkage}} +[[clang::sycl_external]] +static void func1() {} + +// expected-error@+2{{'clang::sycl_external' can only be applied to functions with external linkage}} +namespace { + [[clang::sycl_external]] + void func2() {} +} + +// expected-error@+2{{'clang::sycl_external' can only be applied to functions with external linkage}} +namespace { struct S4 {}; } +[[clang::sycl_external]] void func4(S4) {} + +// expected-error@+3{{'clang::sycl_external' can only be applied to functions with external linkage}} +namespace { struct S5 {}; } +template<typename> [[clang::sycl_external]] void func5(); +template<> [[clang::sycl_external]] void func5<S5>() {} + +namespace { struct S6 {}; } +template<typename> +[[clang::sycl_external]] void func6() {} +template void func6<S6>(); + +// 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{{'clang::sycl_external' can only be applied to functions with external linkage}} +namespace { struct S7 {}; } +template<typename> +[[clang::sycl_external]] void func7(); +template<> void func7<S7>() {} + +// 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<typename> +void func8(); +template<> [[clang::sycl_external]] void func8<S8>() {} +// expected-warning@-1{{'clang::sycl_external' attribute does not appear on the first declaration}} +// expected-error@-2{{'clang::sycl_external' can only be applied to functions with external linkage}} +// expected-note@-3{{previous declaration is here}} + +namespace { struct S9 {}; } +struct T9 { + using type = S9; +}; +template<typename> +[[clang::sycl_external]] void func9() {} +template<typename T> +[[clang::sycl_external]] void test_func9() { + func9<typename T::type>(); +} +template void test_func9<T9>(); + +// The first declaration of a SYCL external function is required to have this attribute. +// expected-note@+1{{previous declaration is here}} +int foo(); +// expected-warning@+1{{'clang::sycl_external' attribute does not appear on the first declaration}} +[[clang::sycl_external]] int foo(); + +// expected-note@+1{{previous declaration is here}} +void goo(); +// expected-warning@+1{{'clang::sycl_external' attribute does not appear on the first declaration}} +[[clang::sycl_external]] void goo(); +void goo() {} + +// expected-note@+1{{previous declaration is here}} +void hoo() {} +// expected-warning@+1{{'clang::sycl_external' attribute does not appear on the first declaration}} +[[clang::sycl_external]] void hoo(); + +// expected-note@+1{{previous declaration is here}} +void joo(); +void use_joo() { + joo(); +} +// expected-warning@+1{{'clang::sycl_external' attribute does not appear on the first declaration}} +[[clang::sycl_external]] void joo(); + +// Subsequent declarations 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(); +}; + +// expected-error@+1{{'clang::sycl_external' cannot be applied to the 'main' function}} +[[clang::sycl_external]] int main() +{ + return 0; +} + +// expected-error@+2{{'clang::sycl_external' cannot be applied to an explicitly deleted function}} +class D { + [[clang::sycl_external]] void mdel() = delete; +}; + +// expected-error@+1{{'clang::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; +}; + +class A { + [[clang::sycl_external]] + A() {} + + [[clang::sycl_external]] void mf() {} + [[clang::sycl_external]] static void smf(); +}; + +class B { +public: + [[clang::sycl_external]] virtual void foo() {} + + [[clang::sycl_external]] virtual void bar() = 0; +}; +[[clang::sycl_external]] void B::bar() {} + +[[clang::sycl_external]] constexpr int square(int x); + +// 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 &&); +template<typename T> +[[clang::sycl_external]] void fun6(T) {} +template void fun6(int *); +template<> [[clang::sycl_external]] void fun6<long*>(long *) {} + +#if CPP20 +[[clang::sycl_external]] consteval int func(); +#endif _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
