# HG changeset patch
# User Justin Holewinski <justin.holewinski@gmail.com>
# Date 1318615158 14400
# Node ID e5fa5cb973fd98c7f7b1e94b17131f9055fcb070
# Parent  6b46db7a9c71bf71d533651cf149cc5b539adb91
Extend TBAA to differentiate between memory spaces in OpenCL-mode

diff --git a/lib/CodeGen/CodeGenTBAA.cpp b/lib/CodeGen/CodeGenTBAA.cpp
--- a/lib/CodeGen/CodeGenTBAA.cpp
+++ b/lib/CodeGen/CodeGenTBAA.cpp
@@ -22,13 +22,17 @@
 #include "llvm/Metadata.h"
 #include "llvm/Constants.h"
 #include "llvm/Type.h"
+#include "llvm/ADT/StringExtras.h"
 using namespace clang;
 using namespace CodeGen;
 
 CodeGenTBAA::CodeGenTBAA(ASTContext &Ctx, llvm::LLVMContext& VMContext,
                          const LangOptions &Features, MangleContext &MContext)
   : Context(Ctx), VMContext(VMContext), Features(Features), MContext(MContext),
-    Root(0), Char(0) {
+    Root(0) {
+  for (unsigned i = 0; i < LangAS::Count+1; ++i) {
+    CharCache[i] = 0;
+  }
 }
 
 CodeGenTBAA::~CodeGenTBAA() {
@@ -45,15 +49,50 @@
   return Root;
 }
 
-llvm::MDNode *CodeGenTBAA::getChar() {
+llvm::MDNode *CodeGenTBAA::getChar(Qualifiers Quals) {
   // Define the root of the tree for user-accessible memory. C and C++
   // give special powers to char and certain similar types. However,
   // these special powers only cover user-accessible memory, and doesn't
   // include things like vtables.
-  if (!Char)
-    Char = getTBAAInfoForNamedType("omnipotent char", getRoot());
 
-  return Char;
+  unsigned CacheIndex = getCacheIndex(Quals);
+
+  if (CharCache[CacheIndex] == 0) {
+    std::string Name = "omnipotent char";
+    // If we are in OpenCL mode, add an address space identifier to the name.
+    if (Context.getLangOptions().OpenCL) {
+      switch (Quals.getAddressSpace()) {
+      default:
+        llvm_unreachable("Unknown OpenCL address space");
+      case 0:
+        Name += " - OCL default";
+        break;
+      case LangAS::opencl_global:
+        Name += " - OCL global";
+        break;
+      case LangAS::opencl_local:
+        Name += " - OCL local";
+        break;
+      case LangAS::opencl_constant:
+        Name += " - OCL constant";
+        break;
+      }
+    }
+    CharCache[CacheIndex] = getTBAAInfoForNamedType(Name, getRoot());
+  }
+  return CharCache[CacheIndex];
+}
+
+unsigned CodeGenTBAA::getCacheIndex(Qualifiers Quals) {
+  unsigned AS = 0;
+  if (Context.getLangOptions().OpenCL) {
+    AS = Quals.getAddressSpace();
+    if (AS < LangAS::Offset || AS >= LangAS::Offset + LangAS::Count)
+     return AS;
+    else
+      return AS - LangAS::Offset;
+  }
+  return AS;
 }
 
 /// getTBAAInfoForNamedType - Create a TBAA tree node with the given string
@@ -94,16 +133,24 @@
   return false;
 }
 
-llvm::MDNode *
-CodeGenTBAA::getTBAAInfo(QualType QTy) {
+llvm::MDNode *CodeGenTBAA::getTBAAInfo(QualType QTy) {
   // If the type has the may_alias attribute (even on a typedef), it is
   // effectively in the general char alias class.
   if (TypeHasMayAlias(QTy))
-    return getChar();
+    return getChar(QTy.getQualifiers());
+  else
+    return getTBAAInfo(QTy.split());
+}
 
-  const Type *Ty = Context.getCanonicalType(QTy).getTypePtr();
+llvm::MDNode *CodeGenTBAA::getTBAAInfo(SplitQualType QTy) {
+  // Retrieve the canonical type and the qualifiers for this QualType.
+  const Type *Ty = Context.getCanonicalType(QTy.first);
+  Qualifiers Quals = QTy.second;
+  unsigned CacheIndex = getCacheIndex(Quals);
+  llvm::DenseMap<const Type *, llvm::MDNode *>& Cache =
+    MetadataCache[CacheIndex];
 
-  if (llvm::MDNode *N = MetadataCache[Ty])
+  if (llvm::MDNode *N = Cache[Ty])
     return N;
 
   // Handle builtin types.
@@ -118,26 +165,28 @@
     case BuiltinType::Char_S:
     case BuiltinType::UChar:
     case BuiltinType::SChar:
-      return getChar();
+      return getChar(Quals);
 
     // Unsigned types can alias their corresponding signed types.
+    // We copy over the qualifiers so the right metadata is emitted in OpenCL
+    // mode.
     case BuiltinType::UShort:
-      return getTBAAInfo(Context.ShortTy);
+      return getTBAAInfo(SplitQualType(Context.ShortTy.getTypePtr(), Quals));
     case BuiltinType::UInt:
-      return getTBAAInfo(Context.IntTy);
+      return getTBAAInfo(SplitQualType(Context.IntTy.getTypePtr(), Quals));
     case BuiltinType::ULong:
-      return getTBAAInfo(Context.LongTy);
+      return getTBAAInfo(SplitQualType(Context.LongTy.getTypePtr(), Quals));
     case BuiltinType::ULongLong:
-      return getTBAAInfo(Context.LongLongTy);
+      return getTBAAInfo(SplitQualType(Context.LongLongTy.getTypePtr(), Quals));
     case BuiltinType::UInt128:
-      return getTBAAInfo(Context.Int128Ty);
+      return getTBAAInfo(SplitQualType(Context.Int128Ty.getTypePtr(), Quals));
 
     // Treat all other builtin types as distinct types. This includes
     // treating wchar_t, char16_t, and char32_t as distinct from their
     // "underlying types".
     default:
-      return MetadataCache[Ty] =
-               getTBAAInfoForNamedType(BTy->getName(Features), getChar());
+      return Cache[Ty] =
+               getTBAAInfoForNamedType(BTy->getName(Features), getChar(Quals));
     }
   }
 
@@ -145,8 +194,7 @@
   // TODO: Implement C++'s type "similarity" and consider dis-"similar"
   // pointers distinct.
   if (Ty->isPointerType())
-    return MetadataCache[Ty] = getTBAAInfoForNamedType("any pointer",
-                                                       getChar());
+    return Cache[Ty] = getTBAAInfoForNamedType("any pointer", getChar(Quals));
 
   // Enum types are distinct types. In C++ they have "underlying types",
   // however they aren't related for TBAA.
@@ -157,7 +205,7 @@
     // members into a single identifying MDNode.
     if (!Features.CPlusPlus &&
         ETy->getDecl()->getTypedefNameForAnonDecl())
-      return MetadataCache[Ty] = getChar();
+      return Cache[Ty] = getChar(Quals);
 
     // In C++ mode, types have linkage, so we can rely on the ODR and
     // on their mangled names, if they're external.
@@ -165,7 +213,7 @@
     // decl with local linkage or no linkage?
     if (Features.CPlusPlus &&
         ETy->getDecl()->getLinkage() != ExternalLinkage)
-      return MetadataCache[Ty] = getChar();
+      return Cache[Ty] = getChar(Quals);
 
     // TODO: This is using the RTTI name. Is there a better way to get
     // a unique string for a type?
@@ -173,9 +221,9 @@
     llvm::raw_svector_ostream Out(OutName);
     MContext.mangleCXXRTTIName(QualType(ETy, 0), Out);
     Out.flush();
-    return MetadataCache[Ty] = getTBAAInfoForNamedType(OutName, getChar());
+    return Cache[Ty] = getTBAAInfoForNamedType(OutName, getChar(Quals));
   }
 
   // For now, handle any other kind of type conservatively.
-  return MetadataCache[Ty] = getChar();
+  return Cache[Ty] = getChar(Quals);
 }
diff --git a/lib/CodeGen/CodeGenTBAA.h b/lib/CodeGen/CodeGenTBAA.h
--- a/lib/CodeGen/CodeGenTBAA.h
+++ b/lib/CodeGen/CodeGenTBAA.h
@@ -15,8 +15,10 @@
 #ifndef CLANG_CODEGEN_CODEGENTBAA_H
 #define CLANG_CODEGEN_CODEGENTBAA_H
 
+#include "clang/AST/ASTContext.h"
 #include "clang/Basic/LLVM.h"
 #include "llvm/ADT/DenseMap.h"
+#include "llvm/Metadata.h"
 
 namespace llvm {
   class LLVMContext;
@@ -28,6 +30,7 @@
   class LangOptions;
   class MangleContext;
   class QualType;
+  class Qualifiers;
   class Type;
 
 namespace CodeGen {
@@ -41,11 +44,11 @@
   const LangOptions &Features;
   MangleContext &MContext;
 
+  llvm::MDNode *Root;
+  llvm::MDNode *CharCache[LangAS::Count+1];
+
   /// MetadataCache - This maps clang::Types to llvm::MDNodes describing them.
-  llvm::DenseMap<const Type *, llvm::MDNode *> MetadataCache;
-
-  llvm::MDNode *Root;
-  llvm::MDNode *Char;
+  llvm::DenseMap<const Type *, llvm::MDNode *> MetadataCache[LangAS::Count+1];
 
   /// getRoot - This is the mdnode for the root of the metadata type graph
   /// for this translation unit.
@@ -53,7 +56,11 @@
 
   /// getChar - This is the mdnode for "char", which is special, and any types
   /// considered to be equivalent to it.
-  llvm::MDNode *getChar();
+  llvm::MDNode *getChar(Qualifiers Quals);
+
+  /// getQualifierIndex - Returns the index into internal caches for the given
+  /// Qualifiers object.
+  unsigned getCacheIndex(Qualifiers Quals);
 
   llvm::MDNode *getTBAAInfoForNamedType(StringRef NameStr,
                                         llvm::MDNode *Parent,
@@ -68,6 +75,7 @@
   /// getTBAAInfo - Get the TBAA MDNode to be used for a dereference
   /// of the given type.
   llvm::MDNode *getTBAAInfo(QualType QTy);
+  llvm::MDNode *getTBAAInfo(SplitQualType QTy);
 };
 
 }  // end namespace CodeGen
diff --git a/test/CodeGenOpenCL/tbaa.cl b/test/CodeGenOpenCL/tbaa.cl
new file mode 100644
--- /dev/null
+++ b/test/CodeGenOpenCL/tbaa.cl
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+__kernel void foo(__global float* a, __local float* b) {
+// CHECK: load float addrspace(2)* %{{.*}}, align 4, !tbaa !1
+// CHECK: store float %{{.*}}, float addrspace(1)* %{{.*}}, align 4, !tbaa !4
+  a[0] = b[0];
+}
+
+// CHECK: !1 = metadata !{metadata !"float", metadata !2}
+// CHECK: !2 = metadata !{metadata !"omnipotent char - OCL local", metadata !3}
+// CHECK: !3 = metadata !{metadata !"Simple C/C++ TBAA", null}
+// CHECK: !4 = metadata !{metadata !"float", metadata !5}
+// CHECK: !5 = metadata !{metadata !"omnipotent char - OCL default", metadata !3}