jchlanda created this revision.
Herald added subscribers: kerbowa, Anastasia, jvesely.
Herald added a project: All.
jchlanda requested review of this revision.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
Use target's address space map to handle cases when both language and
target address spaces are provided. In such case, attempt language to
target translation and only then perform the calculation.
The main motivation is to be able to use language address spaces as
inputs for builtins, which are defined in terms of target address space
(as discussed here: https://reviews.llvm.org/D112718) and hence the
definition of builtins with generic address space pointers that would
allow any other address space pointers inputs (bar constant).
This patch attempts to find a happy medium between not recognising target
address spaces at all (current state) and allowing all uses of it, based on
the assumption that users must know better. What it does not to is to
provide a bidirectional translation mechanism, which I'm not sure could ever
be done, with the current address space implementation (use of 0, the value
of default, etc).
Based on OpenCL rules, this patch follows the conversion guidelines for
`generic` and `constant` address space pointers as described here:
https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_API.html#_memory_model
Repository:
rG LLVM Github Monorepo
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
@@ -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/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<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, 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/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<LangAS>(
+ (*ASMap)[static_cast<unsigned>(LangAS::opencl_generic)] +
+ static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+ LangAS Constant = static_cast<LangAS>(
+ (*ASMap)[static_cast<unsigned>(LangAS::opencl_constant)] +
+ static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+ if (IsATargetAS)
+ B = static_cast<LangAS>(
+ (*ASMap)[static_cast<unsigned>(B)] +
+ static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+ else
+ A = static_cast<LangAS>(
+ (*ASMap)[static_cast<unsigned>(A)] +
+ static_cast<unsigned>(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() ||
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
