diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -470,6 +470,10 @@ Mask |= qs.Mask; } + /// Languages can have different address space semantics, especially with + /// regards to which AS are consider to be overlapping. ASOffload specifies + /// the target language in which the address space was used. + enum class ASOffload { OpenCL, SYCL, None }; /// Returns true if address space A is equal to or a superset of B. /// OpenCL v2.0 defines conversion rules (OpenCLC v2.0 s6.5.5) and notion of /// overlapping address spaces. @@ -477,7 +481,62 @@ /// 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, + ASOffload ASO = ASOffload::None) { + if (ASMap) { + const bool IsATargetAS = isTargetAddressSpace(A); + const bool IsBTargetAS = isTargetAddressSpace(B); + // Do not attempt conversion if both values are expressed in the same + // way (only work on mixed, languate and target AS). + if (IsATargetAS ^ IsBTargetAS) { + if (!IsATargetAS) + A = getLangASFromTargetAS((*ASMap)[static_cast(A)]); + else + B = getLangASFromTargetAS((*ASMap)[static_cast(B)]); + // In OpenCL and SYCL apply the same rules of address space supersets + // as when dealing with language only values, for other cases only + // return true if both values match exactly. + if (ASOffload::OpenCL == ASO) { + LangAS Generic = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::opencl_generic)]); + LangAS Constant = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::opencl_constant)]); + LangAS Global = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::opencl_global)]); + LangAS GlobalDevice = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::opencl_global_device)]); + LangAS GlobalHost = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::opencl_global_host)]); + return A == B || + (A == Generic && B != Constant && Generic != Constant) || + (A == Global && (B == GlobalDevice || B == GlobalHost)); + } + if (ASOffload::SYCL == ASO) { + LangAS Default = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::Default)]); + LangAS Global = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::sycl_global)]); + LangAS GlobalDevice = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::sycl_global_device)]); + LangAS GlobalHost = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::sycl_global_host)]); + LangAS Private = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::sycl_private)]); + LangAS Local = getLangASFromTargetAS( + (*ASMap)[static_cast(LangAS::sycl_local)]); + return A == B || + (A == Global && (B == GlobalDevice || B == GlobalHost)) || + (A == Default && (B == Private || B == Local || B == Global || + B == GlobalDevice || B == GlobalHost)); + } + 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 +573,10 @@ /// 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, + ASOffload ASO = ASOffload::None) { + return isAddressSpaceSupersetOf(this->getAddressSpace(), + other.getAddressSpace(), ASMap, ASO) && // ObjC GC qualifiers can match, be added, or be removed, but can't // be changed. (getObjCGCAttr() == other.getObjCGCAttr() || !hasObjCGCAttr() || diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp --- a/clang/lib/Sema/SemaCast.cpp +++ b/clang/lib/Sema/SemaCast.cpp @@ -2600,16 +2600,23 @@ 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, clang::Qualifiers::ASOffload::OpenCL) || + Qualifiers::isAddressSpaceSupersetOf( + SrcAS, DestAS, &ASMap, clang::Qualifiers::ASOffload::OpenCL); + if (Nested ? DestAS != SrcAS : !OverlappingAS) { Self.Diag(OpRange.getBegin(), DiagID) << SrcType << DestType << Sema::AA_Casting << SrcExpr.get()->getSourceRange(); diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -9218,17 +9218,33 @@ rhq.removeObjCLifetime(); } - if (!lhq.compatiblyIncludes(rhq)) { + auto ASO = clang::Qualifiers::ASOffload::None; + if (S.getLangOpts().OpenCL) + ASO = clang::Qualifiers::ASOffload::OpenCL; + else if (S.getLangOpts().SYCLIsDevice) + ASO = clang::Qualifiers::ASOffload::SYCL; + + const LangASMap &ASMap = S.Context.getTargetInfo().getAddressSpaceMap(); + if (!lhq.compatiblyIncludes(rhq, &ASMap, ASO)) { + const bool AddressSpaceSuperset = Qualifiers::isAddressSpaceSupersetOf( + lhq.getAddressSpace(), rhq.getAddressSpace(), &ASMap, ASO); + // 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 && + (ASO == clang::Qualifiers::ASOffload::OpenCL || + ASO == clang::Qualifiers::ASOffload::SYCL)) + ; // 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. diff --git a/clang/test/Sema/address_space_type_casts_amdgpu.cl b/clang/test/Sema/address_space_type_casts_amdgpu.cl new file mode 100644 --- /dev/null +++ b/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}} +} diff --git a/clang/test/Sema/address_space_type_casts_default.cl b/clang/test/Sema/address_space_type_casts_default.cl new file mode 100644 --- /dev/null +++ b/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}} +} diff --git a/clang/test/SemaOpenCL/atomic-ops.cl b/clang/test/SemaOpenCL/atomic-ops.cl --- a/clang/test/SemaOpenCL/atomic-ops.cl +++ b/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); diff --git a/clang/test/SemaOpenCL/numbered-address-space.cl b/clang/test/SemaOpenCL/numbered-address-space.cl --- a/clang/test/SemaOpenCL/numbered-address-space.cl +++ b/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}} } - diff --git a/clang/test/SemaOpenCL/predefined-expr.cl b/clang/test/SemaOpenCL/predefined-expr.cl --- a/clang/test/SemaOpenCL/predefined-expr.cl +++ b/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__; } diff --git a/clang/test/SemaOpenCL/vector-conv.cl b/clang/test/SemaOpenCL/vector-conv.cl --- a/clang/test/SemaOpenCL/vector-conv.cl +++ b/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; }