intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-02-07 07:39:11 +08:00
Code Issues Packages Projects Releases Wiki Activity

[OpenCL] Add LangAS::opencl_private to represent private address space in AST

Browse Source 
Currently Clang uses default address space (0) to represent private address space for OpenCL
in AST. There are two issues with this:

Multiple address spaces including private address space cannot be diagnosed.
There is no mangling for default address space. For example, if private int* is emitted as
i32 addrspace(5)* in IR. It is supposed to be mangled as PUAS5i but it is mangled as
Pi instead.

This patch attempts to represent OpenCL private address space explicitly in AST. It adds
a new enum LangAS::opencl_private and adds it to the variable types which are implicitly
private:

automatic variables without address space qualifier

function parameter

pointee type without address space qualifier (OpenCL 1.2 and below)

Differential Revision: https://reviews.llvm.org/D35082

llvm-svn: 315668
This commit is contained in:
Yaxun Liu
2017-10-13 03:37:48 +00:00
parent 662bb00028
commit b7318e02c1
21 changed files with 318 additions and 106 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
clang/include/clang/Basic/AddressSpaces.h
View File

@@ -25,16 +25,17 @@ namespace LangAS {
///
enum ID {
// The default value 0 is the value used in QualType for the the situation
// where there is no address space qualifier. For most languages, this also
// corresponds to the situation where there is no address space qualifier in
// the source code, except for OpenCL, where the address space value 0 in
// QualType represents private address space in OpenCL source code.
// where there is no address space qualifier.
Default = 0,
// OpenCL specific address spaces.
// In OpenCL each l-value must have certain non-default address space, each
// r-value must have no address space (i.e. the default address space). The
// pointee of a pointer must have non-default address space.
opencl_global,
opencl_local,
opencl_constant,
opencl_private,
opencl_generic,
// CUDA specific address spaces.

1
clang/lib/AST/ASTContext.cpp
View File

@@ -707,6 +707,7 @@ static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
1, // opencl_global
3, // opencl_local
2, // opencl_constant
0, // opencl_private
4, // opencl_generic
5, // cuda_device
6, // cuda_constant

26
clang/lib/AST/Expr.cpp
View File

@@ -3293,20 +3293,20 @@ Expr::isNullPointerConstant(ASTContext &Ctx,
// Check that it is a cast to void*.
if (const PointerType *PT = CE->getType()->getAs<PointerType>()) {
QualType Pointee = PT->getPointeeType();
Qualifiers Q = Pointee.getQualifiers();
// In OpenCL v2.0 generic address space acts as a placeholder
// and should be ignored.
bool IsASValid = true;
if (Ctx.getLangOpts().OpenCLVersion >= 200) {
if (Pointee.getAddressSpace() == LangAS::opencl_generic)
Q.removeAddressSpace();
else
IsASValid = false;
}
// Only (void*)0 or equivalent are treated as nullptr. If pointee type
// has non-default address space it is not treated as nullptr.
// (__generic void*)0 in OpenCL 2.0 should not be treated as nullptr
// since it cannot be assigned to a pointer to constant address space.
bool PointeeHasDefaultAS =
Pointee.getAddressSpace() == LangAS::Default ||
(Ctx.getLangOpts().OpenCLVersion >= 200 &&
Pointee.getAddressSpace() == LangAS::opencl_generic) ||
(Ctx.getLangOpts().OpenCL &&
Ctx.getLangOpts().OpenCLVersion < 200 &&
Pointee.getAddressSpace() == LangAS::opencl_private);
if (IsASValid && !Q.hasQualifiers() &&
Pointee->isVoidType() && // to void*
CE->getSubExpr()->getType()->isIntegerType()) // from int.
if (PointeeHasDefaultAS && Pointee->isVoidType() && // to void*
CE->getSubExpr()->getType()->isIntegerType()) // from int.
return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
}
}

11
clang/lib/AST/ItaniumMangle.cpp
View File

@@ -2227,15 +2227,17 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
if (Context.getASTContext().addressSpaceMapManglingFor(AS)) {
// <target-addrspace> ::= "AS" <address-space-number>
unsigned TargetAS = Context.getASTContext().getTargetAddressSpace(AS);
ASString = "AS" + llvm::utostr(TargetAS);
if (TargetAS != 0)
ASString = "AS" + llvm::utostr(TargetAS);
} else {
switch (AS) {
default: llvm_unreachable("Not a language specific address space");
// <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant |
// "generic" ]
// <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" |
// "private"| "generic" ]
case LangAS::opencl_global: ASString = "CLglobal"; break;
case LangAS::opencl_local: ASString = "CLlocal"; break;
case LangAS::opencl_constant: ASString = "CLconstant"; break;
case LangAS::opencl_private: ASString = "CLprivate"; break;
case LangAS::opencl_generic: ASString = "CLgeneric"; break;
// <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
case LangAS::cuda_device: ASString = "CUdevice"; break;
@@ -2243,7 +2245,8 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
case LangAS::cuda_shared: ASString = "CUshared"; break;
}
}
mangleVendorQualifier(ASString);
if (!ASString.empty())
mangleVendorQualifier(ASString);
}
// The ARC ownership qualifiers start with underscores.

12
clang/lib/AST/TypePrinter.cpp
View File

@@ -1677,16 +1677,19 @@ void Qualifiers::print(raw_ostream &OS, const PrintingPolicy& Policy,
addSpace = true;
}
if (unsigned addrspace = getAddressSpace()) {
if (addSpace)
OS << ' ';
addSpace = true;
switch (addrspace) {
if (addrspace != LangAS::opencl_private) {
if (addSpace)
OS << ' ';
addSpace = true;
switch (addrspace) {
case LangAS::opencl_global:
OS << "__global";
break;
case LangAS::opencl_local:
OS << "__local";
break;
case LangAS::opencl_private:
break;
case LangAS::opencl_constant:
case LangAS::cuda_constant:
OS << "__constant";
@@ -1705,6 +1708,7 @@ void Qualifiers::print(raw_ostream &OS, const PrintingPolicy& Policy,
OS << "__attribute__((address_space(";
OS << addrspace - LangAS::FirstTargetAddressSpace;
OS << ")))";
}
}
}
if (Qualifiers::GC gc = getObjCGCAttr()) {

4
clang/lib/Basic/Targets/AMDGPU.cpp
View File

@@ -47,6 +47,7 @@ static const LangAS::Map AMDGPUPrivIsZeroDefIsGenMap = {
1, // opencl_global
3, // opencl_local
2, // opencl_constant
0, // opencl_private
4, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -58,6 +59,7 @@ static const LangAS::Map AMDGPUGenIsZeroDefIsGenMap = {
1, // opencl_global
3, // opencl_local
2, // opencl_constant
5, // opencl_private
0, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -69,6 +71,7 @@ static const LangAS::Map AMDGPUPrivIsZeroDefIsPrivMap = {
1, // opencl_global
3, // opencl_local
2, // opencl_constant
0, // opencl_private
4, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -80,6 +83,7 @@ static const LangAS::Map AMDGPUGenIsZeroDefIsPrivMap = {
1, // opencl_global
3, // opencl_local
2, // opencl_constant
5, // opencl_private
0, // opencl_generic
1, // cuda_device
2, // cuda_constant

1
clang/lib/Basic/Targets/NVPTX.h
View File

@@ -28,6 +28,7 @@ static const unsigned NVPTXAddrSpaceMap[] = {
1, // opencl_global
3, // opencl_local
4, // opencl_constant
0, // opencl_private
// FIXME: generic has to be added to the target
0, // opencl_generic
1, // cuda_device

1
clang/lib/Basic/Targets/SPIR.h
View File

@@ -27,6 +27,7 @@ static const unsigned SPIRAddrSpaceMap[] = {
1, // opencl_global
3, // opencl_local
2, // opencl_constant
0, // opencl_private
4, // opencl_generic
0, // cuda_device
0, // cuda_constant

1
clang/lib/Basic/Targets/TCE.h
View File

@@ -35,6 +35,7 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
3, // opencl_global
4, // opencl_local
5, // opencl_constant
0, // opencl_private
// FIXME: generic has to be added to the target
0, // opencl_generic
0, // cuda_device

4
clang/lib/CodeGen/CGDecl.cpp
View File

@@ -956,7 +956,9 @@ void CodeGenFunction::EmitLifetimeEnd(llvm::Value *Size, llvm::Value *Addr) {
CodeGenFunction::AutoVarEmission
CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
QualType Ty = D.getType();
assert(Ty.getAddressSpace() == LangAS::Default);
assert(
Ty.getAddressSpace() == LangAS::Default ||
(Ty.getAddressSpace() == LangAS::opencl_private && getLangOpts().OpenCL));
AutoVarEmission emission(D);

7
clang/lib/Sema/SemaChecking.cpp
View File

@@ -340,7 +340,7 @@ static bool SemaOpenCLBuiltinNDRangeAndBlock(Sema &S, CallExpr *TheCall) {
// First argument is an ndrange_t type.
Expr *NDRangeArg = TheCall->getArg(0);
if (NDRangeArg->getType().getAsString() != "ndrange_t") {
if (NDRangeArg->getType().getUnqualifiedType().getAsString() != "ndrange_t") {
S.Diag(NDRangeArg->getLocStart(),
diag::err_opencl_builtin_expected_type)
<< TheCall->getDirectCallee() << "'ndrange_t'";
@@ -784,8 +784,11 @@ static bool SemaOpenCLBuiltinToAddr(Sema &S, unsigned BuiltinID,
case Builtin::BIto_local:
Qual.setAddressSpace(LangAS::opencl_local);
break;
case Builtin::BIto_private:
Qual.setAddressSpace(LangAS::opencl_private);
break;
default:
Qual.removeAddressSpace();
llvm_unreachable("Invalid builtin function");
}
Call->setType(S.Context.getPointerType(S.Context.getQualifiedType(
RT.getUnqualifiedType(), Qual)));

25
clang/lib/Sema/SemaDecl.cpp
View File

@@ -6324,7 +6324,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
// The event type cannot be used with the __local, __constant and __global
// address space qualifiers.
if (R->isEventT()) {
if (R.getAddressSpace()) {
if (R.getAddressSpace() != LangAS::opencl_private) {
Diag(D.getLocStart(), diag::err_event_t_addr_space_qual);
D.setInvalidType();
}
@@ -7427,7 +7427,7 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
return;
}
}
} else if (T.getAddressSpace() != LangAS::Default) {
} else if (T.getAddressSpace() != LangAS::opencl_private) {
// Do not allow other address spaces on automatic variable.
Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1;
NewVD->setInvalidDecl();
@@ -8062,7 +8062,8 @@ static OpenCLParamType getOpenCLKernelParameterType(Sema &S, QualType PT) {
if (PointeeType->isPointerType())
return PtrPtrKernelParam;
if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
PointeeType.getAddressSpace() == 0)
PointeeType.getAddressSpace() == LangAS::opencl_private ||
PointeeType.getAddressSpace() == LangAS::Default)
return InvalidAddrSpacePtrKernelParam;
return PtrKernelParam;
}
@@ -8832,9 +8833,7 @@ Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
// OpenCL v1.1 s6.5: Using an address space qualifier in a function return
// type declaration will generate a compilation error.
unsigned AddressSpace = NewFD->getReturnType().getAddressSpace();
if (AddressSpace == LangAS::opencl_local ||
AddressSpace == LangAS::opencl_global ||
AddressSpace == LangAS::opencl_constant) {
if (AddressSpace != LangAS::Default) {
Diag(NewFD->getLocation(),
diag::err_opencl_return_value_with_address_space);
NewFD->setInvalidDecl();
@@ -11939,13 +11938,13 @@ ParmVarDecl *Sema::CheckParameter(DeclContext *DC, SourceLocation StartLoc,
// duration shall not be qualified by an address-space qualifier."
// Since all parameters have automatic store duration, they can not have
// an address space.
if (T.getAddressSpace() != 0) {
// OpenCL allows function arguments declared to be an array of a type
// to be qualified with an address space.
if (!(getLangOpts().OpenCL && T->isArrayType())) {
Diag(NameLoc, diag::err_arg_with_address_space);
New->setInvalidDecl();
}
if (T.getAddressSpace() != LangAS::Default &&
// OpenCL allows function arguments declared to be an array of a type
// to be qualified with an address space.
!(getLangOpts().OpenCL &&
(T->isArrayType() || T.getAddressSpace() == LangAS::opencl_private))) {
Diag(NameLoc, diag::err_arg_with_address_space);
New->setInvalidDecl();
}
return New;

130
clang/lib/Sema/SemaType.cpp
View File

@@ -4938,7 +4938,6 @@ TypeSourceInfo *Sema::GetTypeForDeclarator(Declarator &D, Scope *S) {
TypeSourceInfo *ReturnTypeInfo = nullptr;
QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo);
if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount)
inferARCWriteback(state, T);
@@ -5752,9 +5751,10 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
ASIdx = LangAS::opencl_constant; break;
case AttributeList::AT_OpenCLGenericAddressSpace:
ASIdx = LangAS::opencl_generic; break;
case AttributeList::AT_OpenCLPrivateAddressSpace:
ASIdx = LangAS::opencl_private; break;
default:
assert(Attr.getKind() == AttributeList::AT_OpenCLPrivateAddressSpace);
ASIdx = 0; break;
llvm_unreachable("Invalid address space");
}
Type = S.Context.getAddrSpaceQualType(Type, ASIdx);
@@ -6986,6 +6986,92 @@ static void HandleOpenCLAccessAttr(QualType &CurType, const AttributeList &Attr,
}
}
static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
QualType &T, TypeAttrLocation TAL) {
Declarator &D = State.getDeclarator();
// Handle the cases where address space should not be deduced.
//
// The pointee type of a pointer type is alwasy deduced since a pointer always
// points to some memory location which should has an address space.
//
// There are situations that at the point of certain declarations, the address
// space may be unknown and better to be left as default. For example, when
// definining a typedef or struct type, they are not associated with any
// specific address space. Later on, they may be used with any address space
// to declare a variable.
//
// The return value of a function is r-value, therefore should not have
// address space.
//
// The void type does not occupy memory, therefore should not have address
// space, except when it is used as a pointee type.
//
// Since LLVM assumes function type is in default address space, it should not
// have address space.
auto ChunkIndex = State.getCurrentChunkIndex();
bool IsPointee =
ChunkIndex > 0 &&
(D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
bool IsFuncReturnType =
ChunkIndex > 0 &&
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
bool IsFuncType =
ChunkIndex < D.getNumTypeObjects() &&
D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
if ( // Do not deduce addr space for function return type and function type,
// otherwise it will fail some sema check.
IsFuncReturnType || IsFuncType ||
// Do not deduce addr space for member types of struct, except the pointee
// type of a pointer member type.
(D.getContext() == Declarator::MemberContext && !IsPointee) ||
// Do not deduce addr space for types used to define a typedef and the
// typedef itself, except the pointee type of a pointer type which is used
// to define the typedef.
(D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
!IsPointee) ||
// Do not deduce addr space of the void type, e.g. in f(void), otherwise
// it will fail some sema check.
(T->isVoidType() && !IsPointee))
return;
unsigned ImpAddr;
// Put OpenCL automatic variable in private address space.
// OpenCL v1.2 s6.5:
// The default address space name for arguments to a function in a
// program, or local variables of a function is __private. All function
// arguments shall be in the __private address space.
if (State.getSema().getLangOpts().OpenCLVersion <= 120) {
ImpAddr = LangAS::opencl_private;
} else {
// If address space is not set, OpenCL 2.0 defines non private default
// address spaces for some cases:
// OpenCL 2.0, section 6.5:
// The address space for a variable at program scope or a static variable
// inside a function can either be __global or __constant, but defaults to
// __global if not specified.
// (...)
// Pointers that are declared without pointing to a named address space
// point to the generic address space.
if (IsPointee) {
ImpAddr = LangAS::opencl_generic;
} else {
if (D.getContext() == Declarator::FileContext) {
ImpAddr = LangAS::opencl_global;
} else {
if (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static ||
D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
ImpAddr = LangAS::opencl_global;
} else {
ImpAddr = LangAS::opencl_private;
}
}
}
}
T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr);
}
static void processTypeAttrs(TypeProcessingState &state, QualType &type,
TypeAttrLocation TAL, AttributeList *attrs) {
// Scan through and apply attributes to this type where it makes sense. Some
@@ -7157,39 +7243,11 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
}
}
// If address space is not set, OpenCL 2.0 defines non private default
// address spaces for some cases:
// OpenCL 2.0, section 6.5:
// The address space for a variable at program scope or a static variable
// inside a function can either be __global or __constant, but defaults to
// __global if not specified.
// (...)
// Pointers that are declared without pointing to a named address space point
// to the generic address space.
if (state.getSema().getLangOpts().OpenCLVersion >= 200 &&
!hasOpenCLAddressSpace && type.getAddressSpace() == 0 &&
(TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) {
Declarator &D = state.getDeclarator();
if (state.getCurrentChunkIndex() > 0 &&
(D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
DeclaratorChunk::Pointer ||
D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
DeclaratorChunk::BlockPointer)) {
type = state.getSema().Context.getAddrSpaceQualType(
type, LangAS::opencl_generic);
} else if (state.getCurrentChunkIndex() == 0 &&
D.getContext() == Declarator::FileContext &&
!D.isFunctionDeclarator() && !D.isFunctionDefinition() &&
D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
!type->isSamplerT())
type = state.getSema().Context.getAddrSpaceQualType(
type, LangAS::opencl_global);
else if (state.getCurrentChunkIndex() == 0 &&
D.getContext() == Declarator::BlockContext &&
D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static)
type = state.getSema().Context.getAddrSpaceQualType(
type, LangAS::opencl_global);
}
if (!state.getSema().getLangOpts().OpenCL ||
type.getAddressSpace() != LangAS::Default)
return;
deduceOpenCLImplicitAddrSpace(state, type, TAL);
}
void Sema::completeExprArrayBound(Expr *E) {

22
clang/test/CodeGenOpenCL/address-spaces-mangling.cl
View File

@@ -1,5 +1,7 @@
// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s
// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s
// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN10 %s
// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN20 %s
// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN10 %s
// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN20 %s
// We check that the address spaces are mangled the same in both version of OpenCL
// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm -o - | FileCheck -check-prefix=OCL-20 %s
@@ -10,15 +12,17 @@
// warnings, but we do want it for comparison purposes.
__attribute__((overloadable))
void ff(int *arg) { }
// ASMANG: @_Z2ffPi
// NOASMANG: @_Z2ffPi
// ASMANG10: @_Z2ffPi
// ASMANG20: @_Z2ffPU3AS4i
// NOASMANG10: @_Z2ffPi
// NOASMANG20: @_Z2ffPU9CLgenerici
// OCL-20-DAG: @_Z2ffPU3AS4i
// OCL-12-DAG: @_Z2ffPi
__attribute__((overloadable))
void f(private int *arg) { }
// ASMANG: @_Z1fPi
// NOASMANG: @_Z1fPi
// NOASMANG: @_Z1fPU9CLprivatei
// OCL-20-DAG: @_Z1fPi
// OCL-12-DAG: @_Z1fPi
@@ -42,3 +46,11 @@ void f(constant int *arg) { }
// NOASMANG: @_Z1fPU10CLconstanti
// OCL-20-DAG: @_Z1fPU3AS2i
// OCL-12-DAG: @_Z1fPU3AS2i
#if __OPENCL_C_VERSION__ >= 200
__attribute__((overloadable))
void f(generic int *arg) { }
// ASMANG20: @_Z1fPU3AS4i
// NOASMANG20: @_Z1fPU9CLgenerici
// OCL-20-DAG: @_Z1fPU3AS4i
#endif

67
clang/test/CodeGenOpenCL/address-spaces.cl
View File

@@ -7,6 +7,24 @@
// RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
// RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
// SPIR: %struct.S = type { i32, i32, i32* }
// CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* }
struct S {
int x;
int y;
int *z;
};
// CL20-DAG: @g_extern_var = external addrspace(1) global float
// CL20-DAG: @l_extern_var = external addrspace(1) global float
// CL20-DAG: @test_static.l_static_var = internal addrspace(1) global float 0.000000e+00
// CL20-DAG: @g_static_var = internal addrspace(1) global float 0.000000e+00
#ifdef CL20
// CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer
struct S g_s;
#endif
// SPIR: i32* %arg
// GIZ: i32 addrspace(5)* %arg
void f__p(__private int *arg) {}
@@ -58,3 +76,52 @@ void f(int *arg) {
// CL20-DAG: @f.ii = internal addrspace(1) global i32 0
#endif
}
typedef int int_td;
typedef int *intp_td;
// SPIR: define void @test_typedef(i32 addrspace(1)* %x, i32 addrspace(2)* %y, i32* %z)
void test_typedef(global int_td *x, constant int_td *y, intp_td z) {
*x = *y;
*z = 0;
}
// SPIR: define void @test_struct()
void test_struct() {
// SPIR: %ps = alloca %struct.S*
// CL20SPIR: %ps = alloca %struct.S addrspace(4)*
struct S *ps;
// SPIR: store i32 0, i32* %x
// CL20SPIR: store i32 0, i32 addrspace(4)* %x
ps->x = 0;
#ifdef CL20
// CL20SPIR: store i32 0, i32 addrspace(1)* getelementptr inbounds (%struct.S, %struct.S addrspace(1)* @g_s, i32 0, i32 0)
g_s.x = 0;
#endif
}
// SPIR-LABEL: define void @test_void_par()
void test_void_par(void) {}
// SPIR-LABEL: define i32 @test_func_return_type()
int test_func_return_type(void) {
return 0;
}
#ifdef CL20
extern float g_extern_var;
// CL20-LABEL: define {{.*}}void @test_extern(
kernel void test_extern(global float *buf) {
extern float l_extern_var;
buf[0] += g_extern_var + l_extern_var;
}
static float g_static_var;
// CL20-LABEL: define {{.*}}void @test_static(
kernel void test_static(global float *buf) {
static float l_static_var;
buf[0] += g_static_var + l_static_var;
}
#endif

23
clang/test/SemaOpenCL/address-spaces.cl
View File

@@ -1,4 +1,5 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
__constant int ci = 1;
@@ -7,9 +8,15 @@ __kernel void foo(__global int *gip) {
__local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
int *ip;
#if __OPENCL_C_VERSION__ < 200
ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
#else
ip = gip;
ip = &li;
ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
#endif
}
void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
@@ -40,3 +47,19 @@ void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l
l = (local int*) l2;
p = (private int*) p2;
}
__private int func_return_priv(void); //expected-error {{return value cannot be qualified with address space}}
__global int func_return_global(void); //expected-error {{return value cannot be qualified with address space}}
__local int func_return_local(void); //expected-error {{return value cannot be qualified with address space}}
__constant int func_return_constant(void); //expected-error {{return value cannot be qualified with address space}}
#if __OPENCL_C_VERSION__ >= 200
__generic int func_return_generic(void); //expected-error {{return value cannot be qualified with address space}}
#endif
void func_multiple_addr(void) {
typedef __private int private_int_t;
__local __private int var1; // expected-error {{multiple address spaces specified for type}}
__local __private int *var2; // expected-error {{multiple address spaces specified for type}}
__local private_int_t var3; // expected-error {{multiple address spaces specified for type}}
__local private_int_t *var4; // expected-error {{multiple address spaces specified for type}}
}

2
clang/test/SemaOpenCL/cl20-device-side-enqueue.cl
View File

@@ -222,7 +222,7 @@ kernel void foo(global int *buf)
kernel void bar(global int *buf)
{
ndrange_t n;
__private ndrange_t n;
buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected 'ndrange_t' argument type}}
buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected block argument type}}

9
clang/test/SemaOpenCL/extern.cl
View File

@@ -1,9 +0,0 @@
// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
// expected-no-diagnostics
// CHECK: @foo = external addrspace(2) constant float
extern constant float foo;
kernel void test(global float* buf) {
buf[0] += foo;
}

34
clang/test/SemaOpenCL/storageclass-cl20.cl
View File

@@ -1,21 +1,41 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
static constant int G1 = 0;
int G2 = 0;
global int G3 = 0;
local int G4 = 0; // expected-error{{program scope variable must reside in global or constant address space}}
void kernel foo() {
static int S1 = 5;
static global int S2 = 5;
static private int S3 = 5; // expected-error{{static local variable must reside in global or constant address space}}
static float g_implicit_static_var = 0;
static constant float g_constant_static_var = 0;
static global float g_global_static_var = 0;
static local float g_local_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
static generic float g_generic_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
extern float g_implicit_extern_var;
extern constant float g_constant_extern_var;
extern global float g_global_extern_var;
extern local float g_local_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
extern private float g_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
extern generic float g_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
void kernel foo() {
constant int L1 = 0;
local int L2;
global int L3; // expected-error{{function scope variable cannot be declared in global address space}}
generic int L4; // expected-error{{automatic variable qualified with an invalid address space}}
__attribute__((address_space(100))) int L5; // expected-error{{automatic variable qualified with an invalid address space}}
extern global int G5;
extern int G6; // expected-error{{extern variable must reside in global or constant address space}}
static float l_implicit_static_var = 0;
static constant float l_constant_static_var = 0;
static global float l_global_static_var = 0;
static local float l_local_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
static private float l_private_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
static generic float l_generic_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
extern float l_implicit_extern_var;
extern constant float l_constant_extern_var;
extern global float l_global_extern_var;
extern local float l_local_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
extern private float l_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
extern generic float l_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
}

31
clang/test/SemaOpenCL/storageclass.cl
View File

@@ -5,6 +5,20 @@ constant int G2 = 0;
int G3 = 0; // expected-error{{program scope variable must reside in constant address space}}
global int G4 = 0; // expected-error{{program scope variable must reside in constant address space}}
static float g_implicit_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
static constant float g_constant_static_var = 0;
static global float g_global_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
static local float g_local_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
static generic float g_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{program scope variable must reside in constant address space}}
extern float g_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern constant float g_constant_extern_var;
extern global float g_global_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern local float g_local_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern private float g_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern generic float g_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
void kernel foo(int x) {
// static is not allowed at local scope before CL2.0
static int S1 = 5; // expected-error{{variables in function scope cannot be declared static}}
@@ -45,10 +59,17 @@ void f() {
__attribute__((address_space(100))) int L4; // expected-error{{automatic variable qualified with an invalid address space}}
}
static float l_implicit_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
static constant float l_constant_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
static global float l_global_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
static local float l_local_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
static private float l_private_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
static generic float l_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{variables in function scope cannot be declared static}}
extern constant float L5;
extern local float L6; // expected-error{{extern variable must reside in constant address space}}
static int L7 = 0; // expected-error{{variables in function scope cannot be declared static}}
static int L8; // expected-error{{variables in function scope cannot be declared static}}
extern float l_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern constant float l_constant_extern_var;
extern global float l_global_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern local float l_local_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern private float l_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
extern generic float l_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
}

4
clang/test/SemaTemplate/address_space-dependent.cpp
View File

@@ -43,7 +43,7 @@ void neg() {
template <long int I>
void tooBig() {
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388599)}}
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388598)}}
}
template <long int I>
@@ -101,7 +101,7 @@ int main() {
car<1, 2, 3>(); // expected-note {{in instantiation of function template specialization 'car<1, 2, 3>' requested here}}
HasASTemplateFields<1> HASTF;
neg<-1>(); // expected-note {{in instantiation of function template specialization 'neg<-1>' requested here}}
correct<0x7FFFF7>();
correct<0x7FFFF6>();
tooBig<8388650>(); // expected-note {{in instantiation of function template specialization 'tooBig<8388650>' requested here}}
__attribute__((address_space(1))) char *x;