[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.podchishchaeva@intel.com>
This commit is contained in:
schittir
GitHub
parent
85043c1c14
commit
fdfcebb38d
@ -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 = [
|
||||
|
||||
@ -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 = [{
|
||||
|
||||
@ -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
|
||||
|
||||
@ -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<
|
||||
|
||||
@ -64,6 +64,7 @@ public:
|
||||
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);
|
||||
};
|
||||
|
||||
@ -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;
|
||||
|
||||
@ -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().
|
||||
|
||||
@ -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;
|
||||
|
||||
@ -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.
|
||||
|
||||
@ -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);
|
||||
}
|
||||
|
||||
@ -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();
|
||||
}
|
||||
|
||||
@ -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 different template instantiations.
|
||||
tmpl(GLOB);
|
||||
// CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast
|
||||
// CHECK: call spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]])
|
||||
// CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr addrspace(4) [[GLOB]].ascast
|
||||
// CHECK-DAG: call spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]])
|
||||
tmpl(LOC);
|
||||
// CHECK: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast
|
||||
// CHECK: call spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]])
|
||||
// CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr addrspace(4) [[LOC]].ascast
|
||||
// CHECK-DAG: call spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]])
|
||||
tmpl(PRIV);
|
||||
// CHECK: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast
|
||||
// CHECK: call spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef [[PRIV_LOAD5]])
|
||||
// CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr addrspace(4) [[PRIV]].ascast
|
||||
// CHECK-DAG: call spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef [[PRIV_LOAD5]])
|
||||
tmpl(NoAS);
|
||||
// CHECK: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
|
||||
// CHECK: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef [[NoAS_LOAD5]])
|
||||
// CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(4), ptr addrspace(4) [[NoAS]].ascast
|
||||
// CHECK-DAG: call spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef [[NoAS_LOAD5]])
|
||||
}
|
||||
|
||||
// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %
|
||||
// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %
|
||||
// CHECK: define linkonce_odr spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef %
|
||||
// CHECK: define linkonce_odr spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %
|
||||
// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %
|
||||
// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %
|
||||
// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS0iEvT_(ptr noundef %
|
||||
// CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPiEvT_(ptr addrspace(4) noundef %
|
||||
|
||||
@ -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;
|
||||
|
||||
|
||||
|
||||
@ -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;
|
||||
|
||||
@ -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 different template instantiations.
|
||||
tmpl(GLOB);
|
||||
// CHECK: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8
|
||||
// CHECK: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]])
|
||||
// CHECK-DAG: [[GLOB_LOAD4:%[a-zA-Z0-9]+]] = load ptr addrspace(1), ptr [[GLOB]].ascast, align 8
|
||||
// CHECK-DAG: call void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef [[GLOB_LOAD4]])
|
||||
tmpl(LOC);
|
||||
// CHECK: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4
|
||||
// CHECK: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]])
|
||||
// CHECK-DAG: [[LOC_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(3), ptr [[LOC]].ascast, align 4
|
||||
// CHECK-DAG: call void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef [[LOC_LOAD5]])
|
||||
tmpl(PRIV);
|
||||
// CHECK: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4
|
||||
// CHECK: call void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef [[PRIV_LOAD5]])
|
||||
// CHECK-DAG: [[PRIV_LOAD5:%[a-zA-Z0-9]+]] = load ptr addrspace(5), ptr [[PRIV]].ascast, align 4
|
||||
// CHECK-DAG: call void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef [[PRIV_LOAD5]])
|
||||
tmpl(NoAS);
|
||||
// CHECK: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
|
||||
// CHECK: call void @_Z4tmplIPiEvT_(ptr noundef [[NoAS_LOAD5]])
|
||||
// CHECK-DAG: [[NoAS_LOAD5:%[a-zA-Z0-9]+]] = load ptr, ptr [[NoAS]].ascast, align 8
|
||||
// CHECK-DAG: call void @_Z4tmplIPiEvT_(ptr noundef [[NoAS_LOAD5]])
|
||||
}
|
||||
|
||||
// CHECK: declare void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef)
|
||||
// CHECK: declare void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef)
|
||||
// CHECK: declare void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef)
|
||||
// CHECK: declare void @_Z4tmplIPiEvT_(ptr noundef)
|
||||
// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS1iEvT_(ptr addrspace(1) noundef %
|
||||
// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS3iEvT_(ptr addrspace(3) noundef %
|
||||
// CHECK-DAG: define linkonce_odr void @_Z4tmplIPU3AS5iEvT_(ptr addrspace(5) noundef %
|
||||
// CHECK-DAG: define linkonce_odr void @_Z4tmplIPiEvT_(ptr noundef %
|
||||
|
||||
|
||||
@ -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
|
||||
|
||||
@ -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;
|
||||
}
|
||||
|
||||
@ -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]]
|
||||
|
||||
@ -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}
|
||||
//.
|
||||
|
||||
@ -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; }
|
||||
|
||||
|
||||
@ -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
|
||||
|
||||
85
clang/test/CodeGenSYCL/sycl-external-attr.cpp
Normal file
85
clang/test/CodeGenSYCL/sycl-external-attr.cpp
Normal file
@ -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;
|
||||
}
|
||||
@ -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]])
|
||||
|
||||
@ -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);
|
||||
|
||||
@ -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);
|
||||
|
||||
@ -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)
|
||||
|
||||
36
clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp
Normal file
36
clang/test/SemaSYCL/sycl-external-attr-appertainment.cpp
Normal file
@ -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
|
||||
14
clang/test/SemaSYCL/sycl-external-attr-grammar.cpp
Normal file
14
clang/test/SemaSYCL/sycl-external-attr-grammar.cpp
Normal file
@ -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();
|
||||
15
clang/test/SemaSYCL/sycl-external-attr-ignored.cpp
Normal file
15
clang/test/SemaSYCL/sycl-external-attr-ignored.cpp
Normal file
@ -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);
|
||||
154
clang/test/SemaSYCL/sycl-external-attr.cpp
Normal file
154
clang/test/SemaSYCL/sycl-external-attr.cpp
Normal file
@ -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
|
||||
Loading…
x
Reference in New Issue
Block a user