jchlanda updated this revision to Diff 429956.
jchlanda edited the summary of this revision.
jchlanda added a reviewer: Anastasia.
jchlanda added a comment.
Herald added a subscriber: kosarev.
Use helper functions when handling address space values.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D124382/new/
https://reviews.llvm.org/D124382
Files:
clang/include/clang/AST/Type.h
clang/lib/Sema/SemaCast.cpp
clang/lib/Sema/SemaExpr.cpp
clang/test/Sema/address_space_type_casts_amdgpu.cl
clang/test/Sema/address_space_type_casts_default.cl
clang/test/SemaOpenCL/atomic-ops.cl
clang/test/SemaOpenCL/numbered-address-space.cl
clang/test/SemaOpenCL/predefined-expr.cl
clang/test/SemaOpenCL/vector-conv.cl
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;
}
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/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/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/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/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/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ 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.
Index: clang/lib/Sema/SemaCast.cpp
===================================================================
--- clang/lib/Sema/SemaCast.cpp
+++ 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<PointerType>(DestPtr) && isa<PointerType>(SrcPtr)) {
const PointerType *DestPPtr = cast<PointerType>(DestPtr);
const PointerType *SrcPPtr = cast<PointerType>(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();
Index: clang/include/clang/AST/Type.h
===================================================================
--- clang/include/clang/AST/Type.h
+++ 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<unsigned>(A)]);
+ else
+ B = getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(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<unsigned>(LangAS::opencl_generic)]);
+ LangAS Constant = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::opencl_constant)]);
+ LangAS Global = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::opencl_global)]);
+ LangAS GlobalDevice = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::opencl_global_device)]);
+ LangAS GlobalHost = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(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<unsigned>(LangAS::Default)]);
+ LangAS Global = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::sycl_global)]);
+ LangAS GlobalDevice = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::sycl_global_device)]);
+ LangAS GlobalHost = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::sycl_global_host)]);
+ LangAS Private = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(LangAS::sycl_private)]);
+ LangAS Local = getLangASFromTargetAS(
+ (*ASMap)[static_cast<unsigned>(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() ||
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
