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 += ' ';
--- /dev/null
+// 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}}
+}
+++ /dev/null
-// 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}}
-}