Index: clang/include/clang/AST/Type.h =================================================================== --- clang/include/clang/AST/Type.h +++ clang/include/clang/AST/Type.h @@ -477,7 +477,45 @@ /// every address space is a superset of itself. /// CL2.0 adds: /// __generic is a superset of any address space except for __constant. - static bool isAddressSpaceSupersetOf(LangAS A, LangAS B) { + /// If ASMap is provided and address spaces are given in both language and + /// target form the function will attempt to convert language to target + /// address space. + static bool isAddressSpaceSupersetOf(LangAS A, LangAS B, + const LangASMap *ASMap = nullptr, + bool IsSYCLOrOpenCL = false) { + if (ASMap) { + bool IsATargetAS = false; + bool IsBTargetAS = false; + if (A > LangAS::FirstTargetAddressSpace) + IsATargetAS = true; + if (B > LangAS::FirstTargetAddressSpace) + IsBTargetAS = true; + if (IsATargetAS ^ IsBTargetAS) { + LangAS Generic = static_cast( + (*ASMap)[static_cast(LangAS::opencl_generic)] + + static_cast(LangAS::FirstTargetAddressSpace)); + LangAS Constant = static_cast( + (*ASMap)[static_cast(LangAS::opencl_constant)] + + static_cast(LangAS::FirstTargetAddressSpace)); + if (IsATargetAS) + B = static_cast( + (*ASMap)[static_cast(B)] + + static_cast(LangAS::FirstTargetAddressSpace)); + else + A = static_cast( + (*ASMap)[static_cast(A)] + + static_cast(LangAS::FirstTargetAddressSpace)); + // When dealing with target AS return true if: + // * A is equal to B, or + // * in OpenCL or SYCL and A is generic and B is not constant (making + // sure that constant and generic are in target address spaces). + if (IsSYCLOrOpenCL) + return A == B || + (A == Generic && B != Constant && Generic != Constant); + return A == B; + } + } + // Address spaces must match exactly. return A == B || // Otherwise in OpenCLC v2.0 s6.5.5: every address space except @@ -514,8 +552,11 @@ /// Determines if these qualifiers compatibly include another set. /// Generally this answers the question of whether an object with the other /// qualifiers can be safely used as an object with these qualifiers. - bool compatiblyIncludes(Qualifiers other) const { - return isAddressSpaceSupersetOf(other) && + bool compatiblyIncludes(Qualifiers other, const LangASMap *ASMap = nullptr, + bool IsSYCLOrOpenCL = false) { + return isAddressSpaceSupersetOf(this->getAddressSpace(), + other.getAddressSpace(), ASMap, + IsSYCLOrOpenCL) && // ObjC GC qualifiers can match, be added, or be removed, but can't // be changed. (getObjCGCAttr() == other.getObjCGCAttr() || !hasObjCGCAttr() || Index: clang/lib/Sema/SemaCast.cpp =================================================================== --- clang/lib/Sema/SemaCast.cpp +++ clang/lib/Sema/SemaCast.cpp @@ -2600,16 +2600,21 @@ bool Nested = false; unsigned DiagID = diag::err_typecheck_incompatible_address_space; DestPtr = Self.getASTContext().getCanonicalType(DestType.getTypePtr()), - SrcPtr = Self.getASTContext().getCanonicalType(SrcType.getTypePtr()); + SrcPtr = Self.getASTContext().getCanonicalType(SrcType.getTypePtr()); + const LangASMap &ASMap = + Self.getASTContext().getTargetInfo().getAddressSpaceMap(); while (isa(DestPtr) && isa(SrcPtr)) { const PointerType *DestPPtr = cast(DestPtr); const PointerType *SrcPPtr = cast(SrcPtr); QualType DestPPointee = DestPPtr->getPointeeType(); QualType SrcPPointee = SrcPPtr->getPointeeType(); - if (Nested - ? DestPPointee.getAddressSpace() != SrcPPointee.getAddressSpace() - : !DestPPointee.isAddressSpaceOverlapping(SrcPPointee)) { + LangAS DestAS = DestPPointee.getAddressSpace(); + LangAS SrcAS = SrcPPointee.getAddressSpace(); + const bool OverlappingAS = + Qualifiers::isAddressSpaceSupersetOf(DestAS, SrcAS, &ASMap, true) || + Qualifiers::isAddressSpaceSupersetOf(SrcAS, DestAS, &ASMap, true); + if (Nested ? DestAS != SrcAS : !OverlappingAS) { Self.Diag(OpRange.getBegin(), DiagID) << SrcType << DestType << Sema::AA_Casting << SrcExpr.get()->getSourceRange(); Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -9198,17 +9198,27 @@ rhq.removeObjCLifetime(); } - if (!lhq.compatiblyIncludes(rhq)) { + const bool IsSYCLOrOpenCL = + S.getLangOpts().OpenCL || S.getLangOpts().SYCLIsDevice; + const LangASMap &ASMap = S.Context.getTargetInfo().getAddressSpaceMap(); + if (!lhq.compatiblyIncludes(rhq, &ASMap)) { + const bool AddressSpaceSuperset = Qualifiers::isAddressSpaceSupersetOf( + lhq.getAddressSpace(), rhq.getAddressSpace(), &ASMap, IsSYCLOrOpenCL); + // Treat address-space mismatches as fatal. - if (!lhq.isAddressSpaceSupersetOf(rhq)) + if (!AddressSpaceSuperset) return Sema::IncompatiblePointerDiscardsQualifiers; + // In OpenCL/SYCL don't issue discard qualifier warning if address spaces + // overlap. + else if (AddressSpaceSuperset && IsSYCLOrOpenCL) + ; // keep Compatible + // It's okay to add or remove GC or lifetime qualifiers when converting to // and from void*. - else if (lhq.withoutObjCGCAttr().withoutObjCLifetime() - .compatiblyIncludes( - rhq.withoutObjCGCAttr().withoutObjCLifetime()) - && (lhptee->isVoidType() || rhptee->isVoidType())) + else if (lhq.withoutObjCGCAttr().withoutObjCLifetime().compatiblyIncludes( + rhq.withoutObjCGCAttr().withoutObjCLifetime()) && + (lhptee->isVoidType() || rhptee->isVoidType())) ; // keep old // Treat lifetime mismatches as fatal. Index: clang/test/Sema/address_space_type_casts_amdgpu.cl =================================================================== --- /dev/null +++ clang/test/Sema/address_space_type_casts_amdgpu.cl @@ -0,0 +1,38 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s + +void __builtins_AS_3(__attribute__((address_space(3))) int *); + +// Check calling a function using address space 3 (local for AMD) pointer works +// with __local. +__kernel void ker(__local int *IL) { + __builtins_AS_3(IL); +} + +// Check casting __local to address space 3 (local for AMD) pointer works. +__kernel void ker_2(__global int *Array, int N) { + __local int IL; + __attribute__((address_space(3))) int *I3; + I3 = (__attribute__((address_space(3))) int *)&IL; + Array[N] = *I3; +} + +// Check casting __local to address space 5 (private for AMD) pointer errors. +__kernel void ker_3(__global int *Array, int N) { + __local int IP; + __attribute__((address_space(5))) int *I5; + I5 = (__attribute__((address_space(5))) int *)&IP; // expected-error {{casting '__local int *' to type '__attribute__((address_space(5))) int *' changes address space of pointer}} + Array[N] = *I5; +} + +// Check casting of address_space(3) to __generic pointer works. +__kernel void ker_4(__global int *Array, int N, __attribute__((address_space(3))) int *AS3_ptr) { + __generic int *IG; + IG = AS3_ptr; +} + +// Check casting of address_space(4) (__constant) to __generic pointer fails. +__kernel void ker_5(__global int *Array, int N, __attribute__((address_space(4))) int *AS4_ptr) { + __generic int *IG; + IG = AS4_ptr; // expected-error {{assigning '__attribute__((address_space(4))) int *__private' to '__generic int *__private' changes address space of pointer}} +} Index: clang/test/Sema/address_space_type_casts_default.cl =================================================================== --- /dev/null +++ clang/test/Sema/address_space_type_casts_default.cl @@ -0,0 +1,34 @@ +// REQUIRES: x86-registered-target +// RUN: %clang_cc1 -cl-std=CL2.0 -verify -pedantic -fsyntax-only %s + +// The same as address_space_type_cast_amdgpu.cl, but as x86 does not provide +// ASMap all cases should error out. + +void __builtins_AS_3(__attribute__((address_space(3))) int *); // expected-note {{passing argument to parameter here}} + +// No relatioship between address_space(3) and __local on x86. +__kernel void ker(__local int *IL) { + __builtins_AS_3(IL); // expected-error {{passing '__local int *__private' to parameter of type '__attribute__((address_space(3))) int *' changes address space of pointer}} +} + +// No relatioship between address_space(3) and __local on x86. +__kernel void ker_2(__global int *Array, int N) { + __local int IL; + __attribute__((address_space(3))) int *I3; + I3 = (__attribute__((address_space(3))) int *)&IL; // expected-error {{casting '__local int *' to type '__attribute__((address_space(3))) int *' changes address space of pointer}} + Array[N] = *I3; +} + +// No relatioship between address_space(5) and __private on x86. +__kernel void ker_3(__global int *Array, int N) { + __private int IP; + __attribute__((address_space(5))) int *I5; + I5 = (__attribute__((address_space(5))) int *)&IP; // expected-error {{casting '__private int *' to type '__attribute__((address_space(5))) int *' changes address space of pointer}} + Array[N] = *I5; +} + +// Without ASMap compiler can't tell if address_space(3) is not equal to __constant, fail. +__kernel void ker_4(__global int *Array, int N, __attribute__((address_space(3))) int *AS3_ptr) { + __generic int *IG; + IG = AS3_ptr; // expected-error {{assigning '__attribute__((address_space(3))) int *__private' to '__generic int *__private' changes address space of pointer}} +} Index: clang/test/SemaOpenCL/atomic-ops.cl =================================================================== --- clang/test/SemaOpenCL/atomic-ops.cl +++ clang/test/SemaOpenCL/atomic-ops.cl @@ -67,12 +67,12 @@ bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}} - (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}} + (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}} - (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}} + (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // Pointers to different address spaces are allowed. bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); Index: clang/test/SemaOpenCL/numbered-address-space.cl =================================================================== --- clang/test/SemaOpenCL/numbered-address-space.cl +++ clang/test/SemaOpenCL/numbered-address-space.cl @@ -2,11 +2,16 @@ // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s void test_numeric_as_to_generic_implicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) { - generic int* generic_ptr = as3_ptr; // FIXME: This should error + generic int *generic_ptr = as3_ptr; +} + +// AS 4 is constant on AMDGPU, casting it to generic is illegal. +void test_numeric_as_const_to_generic_implicit_cast(__attribute__((address_space(4))) int *as4_ptr, float src) { + generic int *generic_ptr = as4_ptr; // expected-error{{initializing '__generic int *__private' with an expression of type '__attribute__((address_space(4))) int *__private' changes address space of pointer}} } void test_numeric_as_to_generic_explicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) { - generic int* generic_ptr = (generic int*) as3_ptr; // Should maybe be valid? + generic int *generic_ptr = (generic int *)as3_ptr; } void test_generic_to_numeric_as_implicit_cast(void) { @@ -20,12 +25,12 @@ } void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) { - generic int* generic_ptr = as3_ptr; // FIXME: This should error - volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float*) generic_ptr, src, 0, 0, false); // expected-error {{passing '__attribute__((address_space(3))) float *' to parameter of type '__local float *' changes address space of pointer}} + generic int *generic_ptr = as3_ptr; + // This is legal, as address_space(3) corresponds to local on amdgpu. + volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float *)generic_ptr, src, 0, 0, false); } void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) { - generic int* generic_ptr = as3_ptr; + generic int *generic_ptr = as3_ptr; volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *__private' to parameter of type '__local float *' changes address space of pointer}} } - Index: clang/test/SemaOpenCL/predefined-expr.cl =================================================================== --- clang/test/SemaOpenCL/predefined-expr.cl +++ clang/test/SemaOpenCL/predefined-expr.cl @@ -2,7 +2,7 @@ // RUN: %clang_cc1 %s -verify -cl-std=CL2.0 void f() { - char *f1 = __func__; //expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}} - constant char *f2 = __func__; //expected-warning{{initializing '__constant char *__private' with an expression of type 'const __constant char[2]' discards qualifiers}} + char *f1 = __func__; // expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}} + constant char *f2 = __func__; constant const char *f3 = __func__; } Index: clang/test/SemaOpenCL/vector-conv.cl =================================================================== --- clang/test/SemaOpenCL/vector-conv.cl +++ clang/test/SemaOpenCL/vector-conv.cl @@ -16,7 +16,8 @@ e = (constant int4)i; e = (private int4)i; - private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}} - global int4 *global_ptr = const_global_ptr; // expected-warning {{initializing '__global int4 *__private' with an expression of type 'const __global int4 *__private' discards qualifiers}} +private + int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}} + global int4 *global_ptr = const_global_ptr; global_ptr = (global int4 *)const_global_ptr; }