From 5ad3ebaefd65263ddd6adfe1c309819f9d19e5dc Mon Sep 17 00:00:00 2001
From: Peter Collingbourne <peter@pcc.me.uk>
Date: Tue, 8 Nov 2011 02:52:58 +0000
Subject: [PATCH] TypePrinter: print OpenCL address space names.  Patch by
 Richard Membarth, test case by myself.

llvm-svn: 144063
---
 clang/lib/AST/TypePrinter.cpp           | 18 +++++++++++++++---
 clang/test/SemaOpenCL/address-spaces.cl | 13 +++++++++++++
 clang/test/SemaOpenCL/local.cl          |  6 ------
 3 files changed, 28 insertions(+), 9 deletions(-)
 create mode 100644 clang/test/SemaOpenCL/address-spaces.cl
 delete mode 100644 clang/test/SemaOpenCL/local.cl

diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp
index ec6cb48bf866..e11e5065728c 100644
--- a/clang/lib/AST/TypePrinter.cpp
+++ b/clang/lib/AST/TypePrinter.cpp
@@ -1162,9 +1162,21 @@ void Qualifiers::getAsStringInternal(std::string &S,
   AppendTypeQualList(S, getCVRQualifiers());
   if (unsigned addrspace = getAddressSpace()) {
     if (!S.empty()) S += ' ';
-    S += "__attribute__((address_space(";
-    S += llvm::utostr_32(addrspace);
-    S += ")))";
+    switch (addrspace) {
+      case LangAS::opencl_global:
+        S += "__global";
+        break;
+      case LangAS::opencl_local:
+        S += "__local";
+        break;
+      case LangAS::opencl_constant:
+        S += "__constant";
+        break;
+      default:
+        S += "__attribute__((address_space(";
+        S += llvm::utostr_32(addrspace);
+        S += ")))";
+    }
   }
   if (Qualifiers::GC gc = getObjCGCAttr()) {
     if (!S.empty()) S += ' ';
diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
new file mode 100644
index 000000000000..6ab10b3507f7
--- /dev/null
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -0,0 +1,13 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+
+__constant int ci = 1;
+
+__kernel void foo(__global int *gip) {
+  __local int li;
+  __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
+
+  int *ip;
+  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}}
+}
diff --git a/clang/test/SemaOpenCL/local.cl b/clang/test/SemaOpenCL/local.cl
deleted file mode 100644
index 8637cfff30d1..000000000000
--- a/clang/test/SemaOpenCL/local.cl
+++ /dev/null
@@ -1,6 +0,0 @@
-// 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}}
-}