Hi, This patch introduces support for function scope __local variables. It does so by introducing a new storage class, SC_WorkGroupLocal, for __local variables, and an OpenCL runtime class which controls how __local variable declarations are emitted. The default behaviour is simply to emit an internal global.
Reviews appreciated. Thanks, -- Peter
>From 033c98f53aaa88daafa59798fca2c686d6354245 Mon Sep 17 00:00:00 2001 From: Peter Collingbourne <[email protected]> Date: Wed, 27 Jul 2011 18:43:39 +0100 Subject: [PATCH] OpenCL: introduce support for function scope __local variables --- include/clang/AST/Decl.h | 2 +- include/clang/Basic/DiagnosticSemaKinds.td | 2 + include/clang/Basic/Specifiers.h | 1 + lib/AST/Decl.cpp | 1 + lib/AST/DeclPrinter.cpp | 3 +- lib/CodeGen/CGDecl.cpp | 3 ++ lib/CodeGen/CGOpenCLRuntime.cpp | 28 ++++++++++++++++++ lib/CodeGen/CGOpenCLRuntime.h | 43 ++++++++++++++++++++++++++++ lib/CodeGen/CMakeLists.txt | 1 + lib/CodeGen/CodeGenModule.cpp | 14 +++++++-- lib/CodeGen/CodeGenModule.h | 10 ++++++ lib/Sema/SemaDecl.cpp | 19 ++++++++++++ test/CodeGenOpenCL/local.cl | 7 ++++ test/SemaOpenCL/local.cl | 6 ++++ 14 files changed, 135 insertions(+), 5 deletions(-) create mode 100644 lib/CodeGen/CGOpenCLRuntime.cpp create mode 100644 lib/CodeGen/CGOpenCLRuntime.h create mode 100644 test/CodeGenOpenCL/local.cl create mode 100644 test/SemaOpenCL/local.cl diff --git a/include/clang/AST/Decl.h b/include/clang/AST/Decl.h index 938de81..519bad5 100644 --- a/include/clang/AST/Decl.h +++ b/include/clang/AST/Decl.h @@ -802,7 +802,7 @@ public: return !isFileVarDecl(); // Return true for: Auto, Register. - // Return false for: Extern, Static, PrivateExtern. + // Return false for: Extern, Static, PrivateExtern, WorkGroupLocal. return getStorageClass() >= SC_Auto; } diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index 239bd45..b16261b 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -2436,6 +2436,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 cfce0cc..ee09587 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_WorkGroupLocal, SC_Auto, SC_Register }; diff --git a/lib/AST/Decl.cpp b/lib/AST/Decl.cpp index ed8171d..0c9cf8a 100644 --- a/lib/AST/Decl.cpp +++ b/lib/AST/Decl.cpp @@ -1126,6 +1126,7 @@ const char *VarDecl::getStorageClassSpecifierString(StorageClass SC) { case SC_PrivateExtern: return "__private_extern__"; break; case SC_Register: return "register"; break; case SC_Static: return "static"; break; + case SC_WorkGroupLocal:return "work-group-local"; break; } assert(0 && "Invalid storage class"); diff --git a/lib/AST/DeclPrinter.cpp b/lib/AST/DeclPrinter.cpp index 08112cb..d21cd4f 100644 --- a/lib/AST/DeclPrinter.cpp +++ b/lib/AST/DeclPrinter.cpp @@ -364,7 +364,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_WorkGroupLocal: + llvm_unreachable("invalid for functions"); } if (D->isInlineSpecified()) Out << "inline "; diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp index a372d7d..6895e62 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" @@ -130,6 +131,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_WorkGroupLocal: + 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 0000000..3a0e116 --- /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 0000000..252281d --- /dev/null +++ b/lib/CodeGen/CGOpenCLRuntime.h @@ -0,0 +1,43 @@ +//===----- 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(); + + virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF, + const VarDecl &D); +}; + +} +} + +#endif diff --git a/lib/CodeGen/CMakeLists.txt b/lib/CodeGen/CMakeLists.txt index 80e46d2..c080dde 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 12f09b1..d2baa86 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" @@ -64,8 +65,8 @@ 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), VMContext(M.getContext()), NSConcreteGlobalBlockDecl(0), NSConcreteStackBlockDecl(0), NSConcreteGlobalBlock(0), NSConcreteStackBlock(0), @@ -73,7 +74,9 @@ CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO, 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 318c3ea..a38afa1 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; @@ -222,6 +223,7 @@ class CodeGenModule : public CodeGenTypeCache { friend class CodeGenVTables; CGObjCRuntime* ObjCRuntime; + CGOpenCLRuntime* OpenCLRuntime; CGDebugInfo* DebugInfo; ARCEntrypoints *ARCData; RREntrypoints *RRData; @@ -297,6 +299,8 @@ class CodeGenModule : public CodeGenTypeCache { /// Lazily create the Objective-C runtime void createObjCRuntime(); + void createOpenCLRuntime(); + llvm::LLVMContext &VMContext; /// @name Cache for Blocks Runtime Globals @@ -340,6 +344,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 cb82b0c..80a4eb4 100644 --- a/lib/Sema/SemaDecl.cpp +++ b/lib/Sema/SemaDecl.cpp @@ -3615,6 +3615,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_WorkGroupLocal; + } + bool isExplicitSpecialization = false; VarDecl *NewVD; if (!getLangOptions().CPlusPlus) { @@ -3745,6 +3752,7 @@ Sema::ActOnVariableDeclarator(Scope *S, Declarator &D, DeclContext *DC, case SC_Static: case SC_Extern: case SC_PrivateExtern: + case SC_WorkGroupLocal: break; } } @@ -5413,6 +5421,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_WorkGroupLocal) { + 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); @@ -5850,6 +5866,9 @@ void Sema::ActOnCXXForRangeDecl(Decl *D) { case SC_Register: Error = 4; break; + case SC_WorkGroupLocal: + 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 0000000..32fa7be --- /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 0000000..8637cff --- /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}} +} -- 1.7.5.3
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
