diff options
-rw-r--r-- | include/clang/AST/Decl.h | 2 | ||||
-rw-r--r-- | include/clang/Basic/DiagnosticSemaKinds.td | 2 | ||||
-rw-r--r-- | include/clang/Basic/Specifiers.h | 1 | ||||
-rw-r--r-- | lib/AST/Decl.cpp | 13 | ||||
-rw-r--r-- | lib/AST/DeclPrinter.cpp | 3 | ||||
-rw-r--r-- | lib/CodeGen/CGDecl.cpp | 3 | ||||
-rw-r--r-- | lib/CodeGen/CGOpenCLRuntime.cpp | 28 | ||||
-rw-r--r-- | lib/CodeGen/CGOpenCLRuntime.h | 46 | ||||
-rw-r--r-- | lib/CodeGen/CMakeLists.txt | 1 | ||||
-rw-r--r-- | lib/CodeGen/CodeGenModule.cpp | 14 | ||||
-rw-r--r-- | lib/CodeGen/CodeGenModule.h | 10 | ||||
-rw-r--r-- | lib/Sema/SemaDecl.cpp | 19 | ||||
-rw-r--r-- | test/CodeGenOpenCL/local.cl | 7 | ||||
-rw-r--r-- | test/SemaOpenCL/local.cl | 6 |
14 files changed, 144 insertions, 11 deletions
diff --git a/include/clang/AST/Decl.h b/include/clang/AST/Decl.h index 786cc24760..49a118cef7 100644 --- a/include/clang/AST/Decl.h +++ b/include/clang/AST/Decl.h @@ -815,7 +815,7 @@ public: return !isFileVarDecl(); // Return true for: Auto, Register. - // Return false for: Extern, Static, PrivateExtern. + // Return false for: Extern, Static, PrivateExtern, OpenCLWorkGroupLocal. return getStorageClass() >= SC_Auto; } diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index 9d8a9e6618..4dca86f89b 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -2589,6 +2589,8 @@ def err_at_least_one_initializer_needed_to_size_array : Error< def err_array_size_non_int : Error<"size of array has non-integer type %0">; def err_init_element_not_constant : Error< "initializer element is not a compile-time constant">; +def err_local_cant_init : Error< + "'__local' variable cannot have an initializer">; def err_block_extern_cant_init : Error< "'extern' variable cannot have an initializer">; def warn_extern_init : Warning<"'extern' variable has an initializer">; diff --git a/include/clang/Basic/Specifiers.h b/include/clang/Basic/Specifiers.h index cfce0ccbc9..be59ec5a56 100644 --- a/include/clang/Basic/Specifiers.h +++ b/include/clang/Basic/Specifiers.h @@ -146,6 +146,7 @@ namespace clang { SC_PrivateExtern, // These are only legal on variables. + SC_OpenCLWorkGroupLocal, SC_Auto, SC_Register }; diff --git a/lib/AST/Decl.cpp b/lib/AST/Decl.cpp index 252131c95d..4aa1d22be0 100644 --- a/lib/AST/Decl.cpp +++ b/lib/AST/Decl.cpp @@ -1119,12 +1119,13 @@ QualifierInfo::setTemplateParameterListsInfo(ASTContext &Context, const char *VarDecl::getStorageClassSpecifierString(StorageClass SC) { switch (SC) { - case SC_None: break; - case SC_Auto: return "auto"; break; - case SC_Extern: return "extern"; break; - case SC_PrivateExtern: return "__private_extern__"; break; - case SC_Register: return "register"; break; - case SC_Static: return "static"; break; + case SC_None: break; + case SC_Auto: return "auto"; break; + case SC_Extern: return "extern"; break; + case SC_OpenCLWorkGroupLocal: return "<<work-group-local>>"; break; + case SC_PrivateExtern: return "__private_extern__"; break; + case SC_Register: return "register"; break; + case SC_Static: return "static"; break; } assert(0 && "Invalid storage class"); diff --git a/lib/AST/DeclPrinter.cpp b/lib/AST/DeclPrinter.cpp index 5f9c1910e2..866153db11 100644 --- a/lib/AST/DeclPrinter.cpp +++ b/lib/AST/DeclPrinter.cpp @@ -381,7 +381,8 @@ void DeclPrinter::VisitFunctionDecl(FunctionDecl *D) { case SC_Extern: Out << "extern "; break; case SC_Static: Out << "static "; break; case SC_PrivateExtern: Out << "__private_extern__ "; break; - case SC_Auto: case SC_Register: llvm_unreachable("invalid for functions"); + case SC_Auto: case SC_Register: case SC_OpenCLWorkGroupLocal: + llvm_unreachable("invalid for functions"); } if (D->isInlineSpecified()) Out << "inline "; diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index 78a1101ce1..46f3f6b177 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -14,6 +14,7 @@ #include "CGDebugInfo.h" #include "CodeGenFunction.h" #include "CodeGenModule.h" +#include "CGOpenCLRuntime.h" #include "clang/AST/ASTContext.h" #include "clang/AST/CharUnits.h" #include "clang/AST/Decl.h" @@ -131,6 +132,8 @@ void CodeGenFunction::EmitVarDecl(const VarDecl &D) { case SC_PrivateExtern: // Don't emit it now, allow it to be emitted lazily on its first use. return; + case SC_OpenCLWorkGroupLocal: + return CGM.getOpenCLRuntime().EmitWorkGroupLocalVarDecl(*this, D); } assert(0 && "Unknown storage class"); diff --git a/lib/CodeGen/CGOpenCLRuntime.cpp b/lib/CodeGen/CGOpenCLRuntime.cpp new file mode 100644 index 0000000000..3a0e116e5a --- /dev/null +++ b/lib/CodeGen/CGOpenCLRuntime.cpp @@ -0,0 +1,28 @@ +//===----- CGOpenCLRuntime.cpp - Interface to OpenCL Runtimes -------------===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +// +// This provides an abstract class for OpenCL code generation. Concrete +// subclasses of this implement code generation for specific OpenCL +// runtime libraries. +// +//===----------------------------------------------------------------------===// + +#include "CGOpenCLRuntime.h" +#include "CodeGenFunction.h" +#include "llvm/GlobalValue.h" + +using namespace clang; +using namespace CodeGen; + +CGOpenCLRuntime::~CGOpenCLRuntime() {} + +void CGOpenCLRuntime::EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF, + const VarDecl &D) { + return CGF.EmitStaticVarDecl(D, llvm::GlobalValue::InternalLinkage); +} diff --git a/lib/CodeGen/CGOpenCLRuntime.h b/lib/CodeGen/CGOpenCLRuntime.h new file mode 100644 index 0000000000..9a8430fb75 --- /dev/null +++ b/lib/CodeGen/CGOpenCLRuntime.h @@ -0,0 +1,46 @@ +//===----- CGOpenCLRuntime.h - Interface to OpenCL Runtimes -----*- C++ -*-===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// +// +// This provides an abstract class for OpenCL code generation. Concrete +// subclasses of this implement code generation for specific OpenCL +// runtime libraries. +// +//===----------------------------------------------------------------------===// + +#ifndef CLANG_CODEGEN_OPENCLRUNTIME_H +#define CLANG_CODEGEN_OPENCLRUNTIME_H + +namespace clang { + +class VarDecl; + +namespace CodeGen { + +class CodeGenFunction; +class CodeGenModule; + +class CGOpenCLRuntime { +protected: + CodeGenModule &CGM; + +public: + CGOpenCLRuntime(CodeGenModule &CGM) : CGM(CGM) {} + virtual ~CGOpenCLRuntime(); + + /// Emit the IR required for a work-group-local variable declaration, and add + /// an entry to CGF's LocalDeclMap for D. The base class does this using + /// CodeGenFunction::EmitStaticVarDecl to emit an internal global for D. + virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF, + const VarDecl &D); +}; + +} +} + +#endif diff --git a/lib/CodeGen/CMakeLists.txt b/lib/CodeGen/CMakeLists.txt index 80e46d2be7..c080dde467 100644 --- a/lib/CodeGen/CMakeLists.txt +++ b/lib/CodeGen/CMakeLists.txt @@ -31,6 +31,7 @@ add_clang_library(clangCodeGen CGObjCGNU.cpp CGObjCMac.cpp CGObjCRuntime.cpp + CGOpenCLRuntime.cpp CGRecordLayoutBuilder.cpp CGRTTI.cpp CGStmt.cpp diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp index 917f4b7545..1943a744c9 100644 --- a/lib/CodeGen/CodeGenModule.cpp +++ b/lib/CodeGen/CodeGenModule.cpp @@ -18,6 +18,7 @@ #include "CGCall.h" #include "CGCXXABI.h" #include "CGObjCRuntime.h" +#include "CGOpenCLRuntime.h" #include "TargetInfo.h" #include "clang/Frontend/CodeGenOptions.h" #include "clang/AST/ASTContext.h" @@ -65,15 +66,17 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO, ABI(createCXXABI(*this)), Types(C, M, TD, getTargetCodeGenInfo().getABIInfo(), ABI, CGO), TBAA(0), - VTables(*this), ObjCRuntime(0), DebugInfo(0), ARCData(0), RRData(0), - CFConstantStringClassRef(0), ConstantStringClassRef(0), + VTables(*this), ObjCRuntime(0), OpenCLRuntime(0), DebugInfo(0), ARCData(0), + RRData(0), CFConstantStringClassRef(0), ConstantStringClassRef(0), NSConstantStringType(0), VMContext(M.getContext()), NSConcreteGlobalBlock(0), NSConcreteStackBlock(0), BlockObjectAssign(0), BlockObjectDispose(0), BlockDescriptorType(0), GenericBlockLiteralType(0) { if (Features.ObjC1) - createObjCRuntime(); + createObjCRuntime(); + if (Features.OpenCL) + createOpenCLRuntime(); // Enable TBAA unless it's suppressed. if (!CodeGenOpts.RelaxedAliasing && CodeGenOpts.OptimizationLevel > 0) @@ -109,6 +112,7 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO, CodeGenModule::~CodeGenModule() { delete ObjCRuntime; + delete OpenCLRuntime; delete &ABI; delete TBAA; delete DebugInfo; @@ -123,6 +127,10 @@ void CodeGenModule::createObjCRuntime() { ObjCRuntime = CreateMacObjCRuntime(*this); } +void CodeGenModule::createOpenCLRuntime() { + OpenCLRuntime = new CGOpenCLRuntime(*this); +} + void CodeGenModule::Release() { EmitDeferred(); EmitCXXGlobalInitFunc(); diff --git a/lib/CodeGen/CodeGenModule.h b/lib/CodeGen/CodeGenModule.h index aabd7703d0..a5938d9032 100644 --- a/lib/CodeGen/CodeGenModule.h +++ b/lib/CodeGen/CodeGenModule.h @@ -75,6 +75,7 @@ namespace CodeGen { class CGCXXABI; class CGDebugInfo; class CGObjCRuntime; + class CGOpenCLRuntime; class BlockFieldFlags; class FunctionArgList; @@ -226,6 +227,7 @@ class CodeGenModule : public CodeGenTypeCache { friend class CodeGenVTables; CGObjCRuntime* ObjCRuntime; + CGOpenCLRuntime* OpenCLRuntime; CGDebugInfo* DebugInfo; ARCEntrypoints *ARCData; RREntrypoints *RRData; @@ -317,6 +319,8 @@ class CodeGenModule : public CodeGenTypeCache { /// Lazily create the Objective-C runtime void createObjCRuntime(); + void createOpenCLRuntime(); + llvm::LLVMContext &VMContext; /// @name Cache for Blocks Runtime Globals @@ -356,6 +360,12 @@ public: /// been configured. bool hasObjCRuntime() { return !!ObjCRuntime; } + /// getObjCRuntime() - Return a reference to the configured OpenCL runtime. + CGOpenCLRuntime &getOpenCLRuntime() { + assert(OpenCLRuntime != 0); + return *OpenCLRuntime; + } + /// getCXXABI() - Return a reference to the configured C++ ABI. CGCXXABI &getCXXABI() { return ABI; } diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp index 497aa6c9ad..52a05db7cf 100644 --- a/lib/Sema/SemaDecl.cpp +++ b/lib/Sema/SemaDecl.cpp @@ -3734,6 +3734,13 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, } } + if (getLangOptions().OpenCL) { + // Set up the special work-group-local storage class for variables in the + // OpenCL __local address space. + if (R.getAddressSpace() == LangAS::opencl_local) + SC = SC_OpenCLWorkGroupLocal; + } + bool isExplicitSpecialization = false; VarDecl *NewVD; if (!getLangOptions().CPlusPlus) { @@ -3883,6 +3890,7 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, case SC_Static: case SC_Extern: case SC_PrivateExtern: + case SC_OpenCLWorkGroupLocal: break; } } @@ -5715,6 +5723,14 @@ void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, } } + // OpenCL 1.1 6.5.2: "Variables allocated in the __local address space inside + // a kernel function cannot be initialized." + if (VDecl->getStorageClass() == SC_OpenCLWorkGroupLocal) { + Diag(VDecl->getLocation(), diag::err_local_cant_init); + VDecl->setInvalidDecl(); + return; + } + // Capture the variable that is being initialized and the style of // initialization. InitializedEntity Entity = InitializedEntity::InitializeVariable(VDecl); @@ -6131,6 +6147,9 @@ void Sema::ActOnCXXForRangeDecl(Decl *D) { case SC_Register: Error = 4; break; + case SC_OpenCLWorkGroupLocal: + assert(0 && "Unexpected storage class"); + break; } // FIXME: constexpr isn't allowed here. //if (DS.isConstexprSpecified()) diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl new file mode 100644 index 0000000000..32fa7be0f7 --- /dev/null +++ b/test/CodeGenOpenCL/local.cl @@ -0,0 +1,7 @@ +// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s + +__kernel void foo(void) { + // CHECK: @foo.i = internal addrspace(2) + __local int i; + ++i; +} diff --git a/test/SemaOpenCL/local.cl b/test/SemaOpenCL/local.cl new file mode 100644 index 0000000000..8637cfff30 --- /dev/null +++ b/test/SemaOpenCL/local.cl @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only + +__kernel void foo(void) { + __local int i; + __local int j = 2; // expected-error {{'__local' variable cannot have an initializer}} +} |