[cfe-commits] r127915 - in /cfe/trunk: include/clang/AST/ include/clang/Basic/ include/clang/Driver/ include/clang/Parse/ include/clang/Sema/ lib/AST/ lib/Basic/ lib/CodeGen/ lib/Frontend/ lib/Parse/ lib/Sema/ test/CodeGenOpenCL/ test/Parser/

Peter Collingbourne peter at pcc.me.uk
Fri Mar 18 15:38:30 PDT 2011


Author: pcc
Date: Fri Mar 18 17:38:29 2011
New Revision: 127915

URL: http://llvm.org/viewvc/llvm-project?rev=127915&view=rev
Log:
Add support for language-specific address spaces.  On top of that,
add support for the OpenCL __private, __local, __constant and
__global address spaces, as well as the __read_only, _read_write and
__write_only image access specifiers.  Patch originally by ARM;
language-specific address space support by myself.

Added:
    cfe/trunk/include/clang/Basic/AddressSpaces.h
    cfe/trunk/include/clang/Basic/OpenCL.h
    cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
    cfe/trunk/test/Parser/opencl-image-access.cl
Modified:
    cfe/trunk/include/clang/AST/ASTContext.h
    cfe/trunk/include/clang/Basic/LangOptions.h
    cfe/trunk/include/clang/Basic/TargetInfo.h
    cfe/trunk/include/clang/Basic/TokenKinds.def
    cfe/trunk/include/clang/Driver/CC1Options.td
    cfe/trunk/include/clang/Parse/Parser.h
    cfe/trunk/include/clang/Sema/AttributeList.h
    cfe/trunk/lib/AST/ASTContext.cpp
    cfe/trunk/lib/Basic/TargetInfo.cpp
    cfe/trunk/lib/CodeGen/CGCall.cpp
    cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
    cfe/trunk/lib/CodeGen/CGDecl.cpp
    cfe/trunk/lib/CodeGen/CGExpr.cpp
    cfe/trunk/lib/CodeGen/CGExprConstant.cpp
    cfe/trunk/lib/CodeGen/CodeGenModule.cpp
    cfe/trunk/lib/CodeGen/CodeGenTypes.cpp
    cfe/trunk/lib/Frontend/CompilerInvocation.cpp
    cfe/trunk/lib/Parse/ParseDecl.cpp
    cfe/trunk/lib/Sema/AttributeList.cpp
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/lib/Sema/SemaType.cpp

Modified: cfe/trunk/include/clang/AST/ASTContext.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/ASTContext.h?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/ASTContext.h (original)
+++ cfe/trunk/include/clang/AST/ASTContext.h Fri Mar 18 17:38:29 2011
@@ -14,6 +14,7 @@
 #ifndef LLVM_CLANG_AST_ASTCONTEXT_H
 #define LLVM_CLANG_AST_ASTCONTEXT_H
 
+#include "clang/Basic/AddressSpaces.h"
 #include "clang/Basic/IdentifierTable.h"
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/OperatorKinds.h"
@@ -311,6 +312,9 @@
   llvm::OwningPtr<CXXABI> ABI;
   CXXABI *createCXXABI(const TargetInfo &T);
 
+  /// \brief The logical -> physical address space map.
+  const LangAS::Map &AddrSpaceMap;
+
   friend class ASTDeclReader;
 
 public:
@@ -1295,6 +1299,21 @@
   QualType getFloatingTypeOfSizeWithinDomain(QualType typeSize,
                                              QualType typeDomain) const;
 
+  unsigned getTargetAddressSpace(QualType T) const {
+    return getTargetAddressSpace(T.getQualifiers());
+  }
+
+  unsigned getTargetAddressSpace(Qualifiers Q) const {
+    return getTargetAddressSpace(Q.getAddressSpace());
+  }
+
+  unsigned getTargetAddressSpace(unsigned AS) const {
+    if (AS < LangAS::Offset || AS >= LangAS::Offset + LangAS::Count)
+      return AS;
+    else
+      return AddrSpaceMap[AS - LangAS::Offset];
+  }
+
 private:
   // Helper for integer ordering
   unsigned getIntegerRank(const Type *T) const;

Added: cfe/trunk/include/clang/Basic/AddressSpaces.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AddressSpaces.h?rev=127915&view=auto
==============================================================================
--- cfe/trunk/include/clang/Basic/AddressSpaces.h (added)
+++ cfe/trunk/include/clang/Basic/AddressSpaces.h Fri Mar 18 17:38:29 2011
@@ -0,0 +1,44 @@
+//===--- AddressSpaces.h - Language-specific address spaces -----*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+//  This file provides definitions for the various language-specific address
+//  spaces.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_BASIC_ADDRESSSPACES_H
+#define LLVM_CLANG_BASIC_ADDRESSSPACES_H
+
+namespace clang {
+
+namespace LangAS {
+
+/// This enum defines the set of possible language-specific address spaces.
+/// It uses a high starting offset so as not to conflict with any address
+/// space used by a target.
+enum ID {
+  Offset = 0xFFFF00,
+
+  opencl_global = Offset,
+  opencl_local,
+  opencl_constant,
+
+  Last,
+  Count = Last-Offset
+};
+
+/// The type of a lookup table which maps from language-specific address spaces
+/// to target-specific ones.
+typedef unsigned Map[Count];
+
+}
+
+}
+
+#endif

Modified: cfe/trunk/include/clang/Basic/LangOptions.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.h?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.h (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.h Fri Mar 18 17:38:29 2011
@@ -124,6 +124,8 @@
   unsigned DefaultFPContract : 1; // Default setting for FP_CONTRACT
   // FIXME: This is just a temporary option, for testing purposes.
   unsigned NoBitFieldTypeAlign : 1;
+  unsigned FakeAddressSpaceMap : 1; // Use a fake address space map, for
+                                    // testing languages such as OpenCL.
 
   unsigned MRTD : 1;            // -mrtd calling convention
 

Added: cfe/trunk/include/clang/Basic/OpenCL.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenCL.h?rev=127915&view=auto
==============================================================================
--- cfe/trunk/include/clang/Basic/OpenCL.h (added)
+++ cfe/trunk/include/clang/Basic/OpenCL.h Fri Mar 18 17:38:29 2011
@@ -0,0 +1,28 @@
+//===--- OpenCL.h - OpenCL enums --------------------------------*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+//  This file defines some OpenCL-specific enums.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_BASIC_OPENCL_H
+#define LLVM_CLANG_BASIC_OPENCL_H
+
+namespace clang {
+
+/// Names for the OpenCL image access qualifiers (OpenCL 1.1 6.6).
+enum OpenCLImageAccess {
+  CLIA_read_only = 1,
+  CLIA_write_only = 2,
+  CLIA_read_write = 3
+};
+
+}
+
+#endif

Modified: cfe/trunk/include/clang/Basic/TargetInfo.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/TargetInfo.h (original)
+++ cfe/trunk/include/clang/Basic/TargetInfo.h Fri Mar 18 17:38:29 2011
@@ -19,6 +19,7 @@
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/ADT/Triple.h"
 #include "llvm/Support/DataTypes.h"
+#include "clang/Basic/AddressSpaces.h"
 #include <cassert>
 #include <vector>
 #include <string>
@@ -78,6 +79,7 @@
   const llvm::fltSemantics *FloatFormat, *DoubleFormat, *LongDoubleFormat;
   unsigned char RegParmMax, SSERegParmMax;
   TargetCXXABI CXXABI;
+  const LangAS::Map *AddrSpaceMap;
 
   unsigned HasAlignMac68kSupport : 1;
   unsigned RealTypeUsesObjCFPRet : 3;
@@ -530,6 +532,11 @@
   virtual const char *getStaticInitSectionSpecifier() const {
     return 0;
   }
+
+  const LangAS::Map &getAddressSpaceMap() const {
+    return *AddrSpaceMap;
+  }
+
 protected:
   virtual uint64_t getPointerWidthV(unsigned AddrSpace) const {
     return PointerWidth;

Modified: cfe/trunk/include/clang/Basic/TokenKinds.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TokenKinds.def?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/TokenKinds.def (original)
+++ cfe/trunk/include/clang/Basic/TokenKinds.def Fri Mar 18 17:38:29 2011
@@ -252,7 +252,7 @@
 KEYWORD(namespace                   , KEYCXX)
 KEYWORD(new                         , KEYCXX)
 KEYWORD(operator                    , KEYCXX)
-KEYWORD(private                     , KEYCXX)
+KEYWORD(private                     , KEYCXX|KEYOPENCL)
 KEYWORD(protected                   , KEYCXX)
 KEYWORD(public                      , KEYCXX)
 KEYWORD(reinterpret_cast            , KEYCXX)
@@ -350,6 +350,19 @@
 KEYWORD(__kernel                    , KEYOPENCL)
 ALIAS("kernel", __kernel            , KEYOPENCL)
 KEYWORD(vec_step                    , KEYOPENCL)
+KEYWORD(__private                   , KEYOPENCL)
+KEYWORD(__global                    , KEYOPENCL)
+KEYWORD(__local                     , KEYOPENCL)
+KEYWORD(__constant                  , KEYOPENCL)
+ALIAS("global", __global            , KEYOPENCL)
+ALIAS("local", __local              , KEYOPENCL)
+ALIAS("constant", __constant        , KEYOPENCL)
+KEYWORD(__read_only                 , KEYOPENCL)
+KEYWORD(__write_only                , KEYOPENCL)
+KEYWORD(__read_write                , KEYOPENCL)
+ALIAS("read_only", __read_only      , KEYOPENCL)
+ALIAS("write_only", __write_only    , KEYOPENCL)
+ALIAS("read_write", __read_write    , KEYOPENCL)
 
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)

Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Fri Mar 18 17:38:29 2011
@@ -513,6 +513,8 @@
   HelpText<"Ignore bit-field types when aligning structures">;
 def traditional_cpp : Flag<"-traditional-cpp">,
   HelpText<"Enable some traditional CPP emulation">;
+def ffake_address_space_map : Flag<"-ffake-address-space-map">,
+  HelpText<"Use a fake address space map; OpenCL testing purposes only">;
 
 //===----------------------------------------------------------------------===//
 // Header Search Options

Modified: cfe/trunk/include/clang/Parse/Parser.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Parse/Parser.h?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Parse/Parser.h (original)
+++ cfe/trunk/include/clang/Parse/Parser.h Fri Mar 18 17:38:29 2011
@@ -1546,6 +1546,7 @@
   void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
   void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
   void ParseOpenCLAttributes(ParsedAttributes &attrs);
+  void ParseOpenCLQualifiers(DeclSpec &DS);
 
   void ParseTypeofSpecifier(DeclSpec &DS);
   void ParseDecltypeSpecifier(DeclSpec &DS);

Modified: cfe/trunk/include/clang/Sema/AttributeList.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/AttributeList.h?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/AttributeList.h (original)
+++ cfe/trunk/include/clang/Sema/AttributeList.h Fri Mar 18 17:38:29 2011
@@ -18,6 +18,7 @@
 #include "llvm/Support/Allocator.h"
 #include "clang/Sema/Ownership.h"
 #include "clang/Basic/SourceLocation.h"
+#include "clang/AST/Expr.h"
 #include <cassert>
 
 namespace clang {
@@ -76,6 +77,13 @@
                                 declspec, cxx0x);
         return Mem;
       }
+
+    AttributeList* CreateIntegerAttribute(ASTContext &C, IdentifierInfo *Name,
+                                          SourceLocation TokLoc, int Arg) {
+      Expr* IArg = IntegerLiteral::Create(C, llvm::APInt(32, (uint64_t)Arg),
+                                          C.IntTy, TokLoc);
+      return Create( Name, TokLoc, 0, TokLoc, 0, TokLoc, &IArg, 1, 0);
+    }
   };
   
   enum Kind {             // Please keep this list alphabetized.
@@ -135,6 +143,7 @@
     AT_ns_consumed,             // Clang-specific.
     AT_ns_consumes_self,        // Clang-specific.
     AT_objc_gc,
+    AT_opencl_image_access,     // OpenCL-specific.
     AT_opencl_kernel_function,  // OpenCL-specific.
     AT_overloadable,       // Clang-specific.
     AT_ownership_holds,    // Clang-specific.

Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Fri Mar 18 17:38:29 2011
@@ -190,6 +190,22 @@
   return 0;
 }
 
+static const LangAS::Map &getAddressSpaceMap(const TargetInfo &T,
+                                             const LangOptions &LOpts) {
+  if (LOpts.FakeAddressSpaceMap) {
+    // The fake address space map must have a distinct entry for each
+    // language-specific address space.
+    static const unsigned FakeAddrSpaceMap[] = {
+      1, // opencl_global
+      2, // opencl_local
+      3  // opencl_constant
+    };
+    return FakeAddrSpaceMap;
+  } else {
+    return T.getAddressSpaceMap();
+  }
+}
+
 ASTContext::ASTContext(const LangOptions& LOpts, SourceManager &SM,
                        const TargetInfo &t,
                        IdentifierTable &idents, SelectorTable &sels,
@@ -204,7 +220,8 @@
   sigjmp_bufDecl(0), BlockDescriptorType(0), BlockDescriptorExtendedType(0),
   cudaConfigureCallDecl(0),
   NullTypeSourceInfo(QualType()),
-  SourceMgr(SM), LangOpts(LOpts), ABI(createCXXABI(t)), Target(t),
+  SourceMgr(SM), LangOpts(LOpts), ABI(createCXXABI(t)),
+  AddrSpaceMap(getAddressSpaceMap(t, LOpts)), Target(t),
   Idents(idents), Selectors(sels),
   BuiltinInfo(builtins),
   DeclarationNames(*this),
@@ -806,7 +823,8 @@
     Align = Target.getPointerAlign(0);
     break;
   case Type::BlockPointer: {
-    unsigned AS = cast<BlockPointerType>(T)->getPointeeType().getAddressSpace();
+    unsigned AS = getTargetAddressSpace(
+        cast<BlockPointerType>(T)->getPointeeType());
     Width = Target.getPointerWidth(AS);
     Align = Target.getPointerAlign(AS);
     break;
@@ -815,13 +833,14 @@
   case Type::RValueReference: {
     // alignof and sizeof should never enter this code path here, so we go
     // the pointer route.
-    unsigned AS = cast<ReferenceType>(T)->getPointeeType().getAddressSpace();
+    unsigned AS = getTargetAddressSpace(
+        cast<ReferenceType>(T)->getPointeeType());
     Width = Target.getPointerWidth(AS);
     Align = Target.getPointerAlign(AS);
     break;
   }
   case Type::Pointer: {
-    unsigned AS = cast<PointerType>(T)->getPointeeType().getAddressSpace();
+    unsigned AS = getTargetAddressSpace(cast<PointerType>(T)->getPointeeType());
     Width = Target.getPointerWidth(AS);
     Align = Target.getPointerAlign(AS);
     break;
@@ -1468,7 +1487,7 @@
   // the target.
   llvm::APInt ArySize(ArySizeIn);
   ArySize =
-    ArySize.zextOrTrunc(Target.getPointerWidth(EltTy.getAddressSpace()));
+    ArySize.zextOrTrunc(Target.getPointerWidth(getTargetAddressSpace(EltTy)));
 
   llvm::FoldingSetNodeID ID;
   ConstantArrayType::Profile(ID, EltTy, ArySize, ASM, IndexTypeQuals);

Modified: cfe/trunk/lib/Basic/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/TargetInfo.cpp (original)
+++ cfe/trunk/lib/Basic/TargetInfo.cpp Fri Mar 18 17:38:29 2011
@@ -11,6 +11,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "clang/Basic/AddressSpaces.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/LangOptions.h"
 #include "llvm/ADT/APFloat.h"
@@ -19,6 +20,8 @@
 #include <cstdlib>
 using namespace clang;
 
+static const LangAS::Map DefaultAddrSpaceMap = { 0 };
+
 // TargetInfo Constructor.
 TargetInfo::TargetInfo(const std::string &T) : Triple(T) {
   // Set defaults.  Defaults are set for a 32-bit RISC platform, like PPC or
@@ -64,6 +67,9 @@
 
   // Default to using the Itanium ABI.
   CXXABI = CXXABI_Itanium;
+
+  // Default to an empty address space map.
+  AddrSpaceMap = &DefaultAddrSpaceMap;
 }
 
 // Out of line virtual dtor for TargetInfo.

Modified: cfe/trunk/lib/CodeGen/CGCall.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGCall.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGCall.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGCall.cpp Fri Mar 18 17:38:29 2011
@@ -627,7 +627,8 @@
     assert(!RetAI.getIndirectAlign() && "Align unused on indirect return.");
     ResultType = llvm::Type::getVoidTy(getLLVMContext());
     const llvm::Type *STy = ConvertType(RetTy, IsRecursive);
-    ArgTys.push_back(llvm::PointerType::get(STy, RetTy.getAddressSpace()));
+    unsigned AS = Context.getTargetAddressSpace(RetTy);
+    ArgTys.push_back(llvm::PointerType::get(STy, AS));
     break;
   }
 

Modified: cfe/trunk/lib/CodeGen/CGDebugInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDebugInfo.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDebugInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDebugInfo.cpp Fri Mar 18 17:38:29 2011
@@ -462,8 +462,8 @@
   // Bit size, align and offset of the type.
   // Size is always the size of a pointer. We can't use getTypeSize here
   // because that does not return the correct value for references.
-  uint64_t Size = 
-    CGM.getContext().Target.getPointerWidth(PointeeTy.getAddressSpace());
+  unsigned AS = CGM.getContext().getTargetAddressSpace(PointeeTy);
+  uint64_t Size = CGM.getContext().Target.getPointerWidth(AS);
   uint64_t Align = CGM.getContext().getTypeAlign(Ty);
 
   return 

Modified: cfe/trunk/lib/CodeGen/CGDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDecl.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDecl.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDecl.cpp Fri Mar 18 17:38:29 2011
@@ -179,7 +179,8 @@
     new llvm::GlobalVariable(CGM.getModule(), LTy,
                              Ty.isConstant(getContext()), Linkage,
                              CGM.EmitNullConstant(D.getType()), Name, 0,
-                             D.isThreadSpecified(), Ty.getAddressSpace());
+                             D.isThreadSpecified(),
+                             CGM.getContext().getTargetAddressSpace(Ty));
   GV->setAlignment(getContext().getDeclAlign(&D).getQuantity());
   if (Linkage != llvm::GlobalValue::InternalLinkage)
     GV->setVisibility(CurFn->getVisibility());
@@ -222,7 +223,7 @@
                                   OldGV->getLinkage(), Init, "",
                                   /*InsertBefore*/ OldGV,
                                   D.isThreadSpecified(),
-                                  D.getType().getAddressSpace());
+                           CGM.getContext().getTargetAddressSpace(D.getType()));
     GV->setVisibility(OldGV->getVisibility());
     
     // Steal the name of the old global
@@ -289,7 +290,8 @@
   // FIXME: It is really dangerous to store this in the map; if anyone
   // RAUW's the GV uses of this constant will be invalid.
   const llvm::Type *LTy = CGM.getTypes().ConvertTypeForMem(D.getType());
-  const llvm::Type *LPtrTy = LTy->getPointerTo(D.getType().getAddressSpace());
+  const llvm::Type *LPtrTy =
+    LTy->getPointerTo(CGM.getContext().getTargetAddressSpace(D.getType()));
   DMEntry = llvm::ConstantExpr::getBitCast(GV, LPtrTy);
 
   // Emit global variable debug descriptor for static vars.
@@ -724,7 +726,8 @@
 
     // Get the element type.
     const llvm::Type *LElemTy = ConvertTypeForMem(Ty);
-    const llvm::Type *LElemPtrTy = LElemTy->getPointerTo(Ty.getAddressSpace());
+    const llvm::Type *LElemPtrTy =
+      LElemTy->getPointerTo(CGM.getContext().getTargetAddressSpace(Ty));
 
     llvm::Value *VLASize = EmitVLASize(Ty);
 

Modified: cfe/trunk/lib/CodeGen/CGExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGExpr.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGExpr.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGExpr.cpp Fri Mar 18 17:38:29 2011
@@ -726,7 +726,7 @@
     // Cast to the access type.
     const llvm::Type *PTy = llvm::Type::getIntNPtrTy(getLLVMContext(),
                                                      AI.AccessWidth,
-                                                    ExprType.getAddressSpace());
+                              CGM.getContext().getTargetAddressSpace(ExprType));
     Ptr = Builder.CreateBitCast(Ptr, PTy);
 
     // Perform the load.

Modified: cfe/trunk/lib/CodeGen/CGExprConstant.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGExprConstant.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGExprConstant.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGExprConstant.cpp Fri Mar 18 17:38:29 2011
@@ -800,7 +800,7 @@
                                      E->getType().isConstant(CGM.getContext()),
                                      llvm::GlobalValue::InternalLinkage,
                                      C, ".compoundliteral", 0, false,
-                                     E->getType().getAddressSpace());
+                          CGM.getContext().getTargetAddressSpace(E->getType()));
       return C;
     }
     case Expr::DeclRefExprClass: {

Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Mar 18 17:38:29 2011
@@ -1060,7 +1060,7 @@
     Ty = getTypes().ConvertTypeForMem(ASTTy);
 
   const llvm::PointerType *PTy =
-    llvm::PointerType::get(Ty, ASTTy.getAddressSpace());
+    llvm::PointerType::get(Ty, getContext().getTargetAddressSpace(ASTTy));
 
   llvm::StringRef MangledName = getMangledName(D);
   return GetOrCreateLLVMGlobal(MangledName, PTy, D);
@@ -1239,7 +1239,8 @@
   // from the type of the global (this happens with unions).
   if (GV == 0 ||
       GV->getType()->getElementType() != InitType ||
-      GV->getType()->getAddressSpace() != ASTTy.getAddressSpace()) {
+      GV->getType()->getAddressSpace() !=
+        getContext().getTargetAddressSpace(ASTTy)) {
 
     // Move the old entry aside so that we'll create a new one.
     Entry->setName(llvm::StringRef());

Modified: cfe/trunk/lib/CodeGen/CodeGenTypes.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenTypes.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenTypes.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenTypes.cpp Fri Mar 18 17:38:29 2011
@@ -270,14 +270,16 @@
     QualType ETy = RTy.getPointeeType();
     llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
     PointersToResolve.push_back(std::make_pair(ETy, PointeeType));
-    return llvm::PointerType::get(PointeeType, ETy.getAddressSpace());
+    unsigned AS = Context.getTargetAddressSpace(ETy);
+    return llvm::PointerType::get(PointeeType, AS);
   }
   case Type::Pointer: {
     const PointerType &PTy = cast<PointerType>(Ty);
     QualType ETy = PTy.getPointeeType();
     llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
     PointersToResolve.push_back(std::make_pair(ETy, PointeeType));
-    return llvm::PointerType::get(PointeeType, ETy.getAddressSpace());
+    unsigned AS = Context.getTargetAddressSpace(ETy);
+    return llvm::PointerType::get(PointeeType, AS);
   }
 
   case Type::VariableArray: {
@@ -402,7 +404,8 @@
     const QualType FTy = cast<BlockPointerType>(Ty).getPointeeType();
     llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
     PointersToResolve.push_back(std::make_pair(FTy, PointeeType));
-    return llvm::PointerType::get(PointeeType, FTy.getAddressSpace());
+    unsigned AS = Context.getTargetAddressSpace(FTy);
+    return llvm::PointerType::get(PointeeType, AS);
   }
 
   case Type::MemberPointer: {

Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Fri Mar 18 17:38:29 2011
@@ -672,6 +672,8 @@
     Res.push_back("-fconstant-string-class");
     Res.push_back(Opts.ObjCConstantStringClass);
   }
+  if (Opts.FakeAddressSpaceMap)
+    Res.push_back("-ffake-address-space-map");
 }
 
 static void PreprocessorOptsToArgs(const PreprocessorOptions &Opts,
@@ -1483,6 +1485,7 @@
   Opts.FastRelaxedMath = Args.hasArg(OPT_cl_fast_relaxed_math);
   Opts.OptimizeSize = 0;
   Opts.MRTD = Args.hasArg(OPT_mrtd);
+  Opts.FakeAddressSpaceMap = Args.hasArg(OPT_ffake_address_space_map);
 
   // FIXME: Eliminate this dependency.
   unsigned Opt = getOptimizationLevel(Args, IK, Diags);

Modified: cfe/trunk/lib/Parse/ParseDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseDecl.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseDecl.cpp (original)
+++ cfe/trunk/lib/Parse/ParseDecl.cpp Fri Mar 18 17:38:29 2011
@@ -13,6 +13,7 @@
 
 #include "clang/Parse/Parser.h"
 #include "clang/Parse/ParseDiagnostic.h"
+#include "clang/Basic/OpenCL.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ParsedTemplate.h"
 #include "clang/Sema/PrettyDeclStackTrace.h"
@@ -311,6 +312,56 @@
   }
 }
 
+void Parser::ParseOpenCLQualifiers(DeclSpec &DS) {
+  SourceLocation Loc = Tok.getLocation();
+  switch(Tok.getKind()) {
+    // OpenCL qualifiers:
+    case tok::kw___private:
+    case tok::kw_private: 
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute( 
+          Actions.getASTContext(), 
+          PP.getIdentifierInfo("address_space"), Loc, 0));
+      break;
+      
+    case tok::kw___global:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(),
+          PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_global));
+      break;
+      
+    case tok::kw___local:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(),
+          PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_local));
+      break;
+      
+    case tok::kw___constant:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(),
+          PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_constant));
+      break;
+      
+    case tok::kw___read_only:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(), 
+          PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_read_only));
+      break;
+      
+    case tok::kw___write_only:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(), 
+          PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_write_only));
+      break;
+      
+    case tok::kw___read_write:
+      DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+          Actions.getASTContext(),
+          PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_read_write));
+      break;
+    default: break;
+  }
+}
+
 void Parser::DiagnoseProhibitedAttributes(ParsedAttributesWithRange &attrs) {
   Diag(attrs.Range.getBegin(), diag::err_attributes_not_allowed)
     << attrs.Range;
@@ -1446,6 +1497,20 @@
       ParseDecltypeSpecifier(DS);
       continue;
 
+    // OpenCL qualifiers:
+    case tok::kw_private: 
+      if (!getLang().OpenCL)
+        goto DoneWithDeclSpec;
+    case tok::kw___private:
+    case tok::kw___global:
+    case tok::kw___local:
+    case tok::kw___constant:
+    case tok::kw___read_only:
+    case tok::kw___write_only:
+    case tok::kw___read_write:
+      ParseOpenCLQualifiers(DS);
+      break;
+      
     case tok::less:
       // GCC ObjC supports types like "<SomeProtocol>" as a synonym for
       // "id<SomeProtocol>".  This is hopelessly old fashioned and dangerous,
@@ -1697,6 +1762,20 @@
     ParseDecltypeSpecifier(DS);
     return true;
 
+  // OpenCL qualifiers:
+  case tok::kw_private: 
+    if (!getLang().OpenCL)
+      return false;
+  case tok::kw___private:
+  case tok::kw___global:
+  case tok::kw___local:
+  case tok::kw___constant:
+  case tok::kw___read_only:
+  case tok::kw___write_only:
+  case tok::kw___read_write:
+    ParseOpenCLQualifiers(DS);
+    break;
+
   // C++0x auto support.
   case tok::kw_auto:
     if (!getLang().CPlusPlus0x)
@@ -2269,10 +2348,22 @@
 bool Parser::isTypeQualifier() const {
   switch (Tok.getKind()) {
   default: return false;
+
+    // type-qualifier only in OpenCL
+  case tok::kw_private:
+    return getLang().OpenCL;
+
     // type-qualifier
   case tok::kw_const:
   case tok::kw_volatile:
   case tok::kw_restrict:
+  case tok::kw___private:
+  case tok::kw___local:
+  case tok::kw___global:
+  case tok::kw___constant:
+  case tok::kw___read_only:
+  case tok::kw___read_write:
+  case tok::kw___write_only:
     return true;
   }
 }
@@ -2400,7 +2491,19 @@
   case tok::kw___w64:
   case tok::kw___ptr64:
   case tok::kw___pascal:
+
+  case tok::kw___private:
+  case tok::kw___local:
+  case tok::kw___global:
+  case tok::kw___constant:
+  case tok::kw___read_only:
+  case tok::kw___read_write:
+  case tok::kw___write_only:
+
     return true;
+
+  case tok::kw_private:
+    return getLang().OpenCL;
   }
 }
 
@@ -2413,6 +2516,9 @@
   switch (Tok.getKind()) {
   default: return false;
 
+  case tok::kw_private:
+    return getLang().OpenCL;
+
   case tok::identifier:   // foo::bar
     // Unfortunate hack to support "Class.factoryMethod" notation.
     if (getLang().ObjC1 && NextToken().is(tok::period))
@@ -2522,6 +2628,15 @@
   case tok::kw___ptr64:
   case tok::kw___forceinline:
   case tok::kw___pascal:
+
+  case tok::kw___private:
+  case tok::kw___local:
+  case tok::kw___global:
+  case tok::kw___constant:
+  case tok::kw___read_only:
+  case tok::kw___read_write:
+  case tok::kw___write_only:
+
     return true;
   }
 }
@@ -2627,6 +2742,21 @@
       isInvalid = DS.SetTypeQual(DeclSpec::TQ_restrict, Loc, PrevSpec, DiagID,
                                  getLang());
       break;
+
+    // OpenCL qualifiers:
+    case tok::kw_private: 
+      if (!getLang().OpenCL)
+        goto DoneWithTypeQuals;
+    case tok::kw___private:
+    case tok::kw___global:
+    case tok::kw___local:
+    case tok::kw___constant:
+    case tok::kw___read_only:
+    case tok::kw___write_only:
+    case tok::kw___read_write:
+      ParseOpenCLQualifiers(DS);
+      break;
+
     case tok::kw___w64:
     case tok::kw___ptr64:
     case tok::kw___cdecl:

Modified: cfe/trunk/lib/Sema/AttributeList.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/AttributeList.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/AttributeList.cpp (original)
+++ cfe/trunk/lib/Sema/AttributeList.cpp Fri Mar 18 17:38:29 2011
@@ -94,6 +94,7 @@
     .Case("unavailable", AT_unavailable)
     .Case("overloadable", AT_overloadable)
     .Case("address_space", AT_address_space)
+    .Case("opencl_image_access", AT_opencl_image_access)
     .Case("always_inline", AT_always_inline)
     .Case("returns_twice", IgnoredAttribute)
     .Case("vec_type_hint", IgnoredAttribute)

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Fri Mar 18 17:38:29 2011
@@ -2724,6 +2724,7 @@
   case AttributeList::AT_IBOutletCollection:
       HandleIBOutletCollection(D, Attr, S); break;
   case AttributeList::AT_address_space:
+  case AttributeList::AT_opencl_image_access:
   case AttributeList::AT_objc_gc:
   case AttributeList::AT_vector_size:
   case AttributeList::AT_neon_vector_type:

Modified: cfe/trunk/lib/Sema/SemaType.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=127915&r1=127914&r2=127915&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaType.cpp (original)
+++ cfe/trunk/lib/Sema/SemaType.cpp Fri Mar 18 17:38:29 2011
@@ -13,6 +13,7 @@
 
 #include "clang/Sema/SemaInternal.h"
 #include "clang/Sema/Template.h"
+#include "clang/Basic/OpenCL.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/CXXInheritance.h"
 #include "clang/AST/DeclObjC.h"
@@ -2965,6 +2966,41 @@
   return true;
 }
 
+/// Handle OpenCL image access qualifiers: read_only, write_only, read_write
+static void HandleOpenCLImageAccessAttribute(QualType& CurType,
+                                             const AttributeList &Attr,
+                                             Sema &S) {
+  // Check the attribute arguments.
+  if (Attr.getNumArgs() != 1) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 1;
+    Attr.setInvalid();
+    return;
+  }
+  Expr *sizeExpr = static_cast<Expr *>(Attr.getArg(0));
+  llvm::APSInt arg(32);
+  if (sizeExpr->isTypeDependent() || sizeExpr->isValueDependent() ||
+      !sizeExpr->isIntegerConstantExpr(arg, S.Context)) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_not_int)
+      << "opencl_image_access" << sizeExpr->getSourceRange();
+    Attr.setInvalid();
+    return;
+  }
+  unsigned iarg = static_cast<unsigned>(arg.getZExtValue());
+  switch (iarg) {
+  case CLIA_read_only:
+  case CLIA_write_only:
+  case CLIA_read_write:
+    // Implemented in a separate patch
+    break;
+  default:
+    // Implemented in a separate patch
+    S.Diag(Attr.getLoc(), diag::err_attribute_invalid_size)
+      << sizeExpr->getSourceRange();
+    Attr.setInvalid();
+    break;
+  }
+}
+
 /// HandleVectorSizeAttribute - this attribute is only applicable to integral
 /// and float scalars, although arrays, pointers, and function return values are
 /// allowed in conjunction with this construct. Aggregates with this attribute
@@ -3119,6 +3155,10 @@
                                "neon_polyvector_type");
       break;
 
+    case AttributeList::AT_opencl_image_access:
+      HandleOpenCLImageAccessAttribute(type, attr, state.getSema());
+      break;
+
     FUNCTION_TYPE_ATTRS_CASELIST:
       // Never process function type attributes as part of the
       // declaration-specifiers.

Added: cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces.cl?rev=127915&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-spaces.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/address-spaces.cl Fri Mar 18 17:38:29 2011
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+void f__p(__private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void f__g(__global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void f__l(__local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void f__c(__constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+
+
+void fp(private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void fg(global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void fl(local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void fc(constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+

Added: cfe/trunk/test/Parser/opencl-image-access.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Parser/opencl-image-access.cl?rev=127915&view=auto
==============================================================================
--- cfe/trunk/test/Parser/opencl-image-access.cl (added)
+++ cfe/trunk/test/Parser/opencl-image-access.cl Fri Mar 18 17:38:29 2011
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 %s -fsyntax-only
+
+typedef void* image2d_t;
+
+__kernel void f__ro(__read_only image2d_t a) { }
+
+__kernel void f__wo(__write_only image2d_t a) { }
+
+__kernel void f__rw(__read_write image2d_t a) { }
+
+
+__kernel void fro(read_only image2d_t a) { }
+
+__kernel void fwo(write_only image2d_t a) { }
+
+__kernel void frw(read_write image2d_t a) { }





More information about the cfe-commits mailing list