Skip to content

[SYCL] Re-use OpenCL address space attributes for SYCL #1039

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed

[SYCL] Re-use OpenCL address space attributes for SYCL #1039

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
15 changes: 6 additions & 9 deletions clang/include/clang/AST/Type.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -486,15 +486,12 @@ class Qualifiers {
/// Returns true if the address space in these qualifiers is equal to or
/// a superset of the address space in the argument qualifiers.
bool isAddressSpaceSupersetOf(Qualifiers other) const {

return
isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace()) ||
(!hasAddressSpace() &&
(other.getAddressSpace() == LangAS::sycl_private ||
other.getAddressSpace() == LangAS::sycl_local ||
other.getAddressSpace() == LangAS::sycl_global ||
other.getAddressSpace() == LangAS::sycl_constant ||
other.getAddressSpace() == LangAS::sycl_generic));
return isAddressSpaceSupersetOf(getAddressSpace(),
other.getAddressSpace()) ||
(!hasAddressSpace() &&
(other.getAddressSpace() == LangAS::opencl_private ||
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems strange. Why do you need this downstream change here?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What are you trying to achieve?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'd like to avoid adding AS qualifiers to all types and treat all objects of "unqualified" types to be allocated in "generic" address space. This allows us to emit valid SPIR-V files for the SYCL code, which inter-operates regular C++ code satisfying SYCL kernel restrictions.
We qualify with address spaces only types required by the SPIR-V spec (e.g. kernel parameters, program scope variables, etc). All other objects (not qualified explicitly by user via SYCL pointer classes), are residing in "generic" address space. This approach allows us altering C++ type system.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In OpenCL (and other languages too) we do qualifiers deduction during parsing. All pointers without explicit address space are in generic in OpenCL v2.0. I guess you could just do the same.

This function determines qualifier relations regardless of exact type. Also this function should just have the same behaviour as the one above since they are to be used interchangeably. I think we just added this function to allow calling it as a member but it just called the helper. So I don't understand why these two have different logic in your downstream changes. They should return the same result for the same pair of qualifiers.

Copy link
Contributor

@AnastasiaStulova AnastasiaStulova Jan 23, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We have recently changed how deduction works:
https://reviews.llvm.org/D65744

My guess is you just need to alter deduceOpenCLPointeeAddrSpace slightly and make sure you invoke it in the same places.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

All pointers in AST have the same address space (default).

I presume if pointer address space was provided explicitly you would keep it?

This way we follow standard C++ type system, so every valid C++ code can be compiled and executed.

What about C++ features that can't be easily supported by OpenCL devices like virtual functions?

In Clang CG, concrete address space pointers appear: alloca produces private pointers, kernel arguments can be global or local. Since AST assumes that all pointers have the same type, most of these concrete pointers are immediately casted to generic.

I don't see how they appear in CodeGen if you didn't have them in AST at all? I presume you have some address spaces somewhere in AST? Also how about address spaces that can't convert to generic like constant. Do you support this?

Later in the compilation flow we use llvm::InferAddressSpaces pass to optimize addrspace casts where the algorithm can identify the real address space behind a generic pointer.

That makes sense. I guess the limitation is if you compile separate translation modules or don't inline function calls then the inference will be limited.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

All pointers in AST have the same address space (default).

I presume if pointer address space was provided explicitly you would keep it?

Yes, but the only way to do that is through sycl::multi_ptr (and sycl::accessor, forgot about it) classes. User never writes

This way we follow standard C++ type system, so every valid C++ code can be compiled and executed.

What about C++ features that can't be easily supported by OpenCL devices like virtual functions?

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

In Clang CG, concrete address space pointers appear: alloca produces private pointers, kernel arguments can be global or local. Since AST assumes that all pointers have the same type, most of these concrete pointers are immediately casted to generic.

I don't see how they appear in CodeGen if you didn't have them in AST at all? I presume you have some address spaces somewhere in AST?

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

Also how about address spaces that can't convert to generic like constant. Do you support this?

Constant is not fully supported as a raw pointer. It cannot be converted to generic, so if we allow it, then we risk emitting LLVM IR that cannot be lowered to SPIR-V.
Constant can still be used with sycl::multi_ptr and sycl::accessor classes though.

Later in the compilation flow we use llvm::InferAddressSpaces pass to optimize addrspace casts where the algorithm can identify the real address space behind a generic pointer.

That makes sense. I guess the limitation is if you compile separate translation modules or don't inline function calls then the inference will be limited.

Right. More limitations come if a pointer is stored in memory and then loaded elsewhere.
Unless mem2reg eliminates this store/load, it is difficult to prove anything.

Copy link
Contributor

@AnastasiaStulova AnastasiaStulova Jan 24, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

Right this is great! But I was thinking of generic supported by multiple vendors. Does it mean that SYCL will define a new type of devices that it supports i.e. not just any OpenCL device?

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

I am thinking of some language related aspect that you might miss to diagnose with this strategy. Let's say if user code creates pointers:

loc_ptr1 - pointer to local addr space
loc_ptr2 - pointer to local addr space
glob_ptr - pointer to global addr space

then are they cast to generic (no address space) straight away and then if they appear in the following statement:

loc_ptr1 = somevariable ? loc_ptr2 : glob_ptr;

if you don't have address spaces any more at this point you can't really provide an error saying that the code is likely illegal?

Are you aware of the cases like that and how severe they might be?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We do support function pointers on CPU and GPU as an extension:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/SPIRV/SPV_INTEL_function_pointers.asciidoc

Right this is great! But I was thinking of generic supported by multiple vendors. Does it mean that SYCL will define a new type of devices that it supports i.e. not just any OpenCL device?

Since OpenCL doesn't have this in the core specification, it is hard (impossible?) to provide a solution that works everywhere (with reasonable performance). But the extension is fairly generic, so I guess any vendor interested in this feature can implement it.

Sorry, I should've say this better.
You're correct, some address spaces are present in AST, but they are "hidden" from a user inside C++ classes, so there is no legal way to get their type.
Alloca as I mentioned, is always addrspace(0) from data layout, this address space does not come from AST.

I am thinking of some language related aspect that you might miss to diagnose with this strategy. Let's say if user code creates pointers:

loc_ptr1 - pointer to local addr space
loc_ptr2 - pointer to local addr space
glob_ptr - pointer to global addr space

Address space qualifiers are not exposed in the SYCL language, so there is no way for a user to create a qualified raw pointer. (unless they use internal attributes directly; nobody should do this)

then are they cast to generic (no address space) straight away and then if they appear in the following statement:

loc_ptr1 = somevariable ? loc_ptr2 : glob_ptr;

if you don't have address spaces any more at this point you can't really provide an error saying that the code is likely illegal?

This is perfectly legal, if you write it without address space qualifiers.
Consider this (not real SYCL code, but close enough):

void main() {
  global_ptr<int> glob_ptr;
  local_ptr<int> loc_ptr;
  int cond = random();
  parallel_for([]() {
    // kernel code
    int *glob = glob_ptr.get();  // glob_ptr.get returns an int* in default AS (== generic)
    int *loc= loc_ptr.get();     // loc_ptr.get also returns an int* in default AS
    int *ptr = cond? glob : loc; // ptr is also default
    *ptr = 42; // this is going to be a store by a generic pointer, unless the compiler
               // figures out what `cond' is.
  }
}

In LLVM IR this looks like:

define kernel(i32 addrspace(1)* glob_ptr, i32 addrspace(3)* loc_ptr, i32 cond) {
    %glob = addrspacecast %glob_ptr to i32 addrspace(4)*
    %loc = addrspacecast %loc_ptr to i32 addrspace(4)*
    %ptr = select %cond,  i32 addrspace(4)* %glob,  i32 addrspace(4)* %loc
    store  i32 addrspace(4)* %ptr, 42
}

Again, this a valid code for a device that supports generic address space. If a device doesn't support it, then compilation terminates (or gives a diagnostic) from the device backend/middle-end compiler (it is not as pretty as the a diagnostic from Clang). Although, for less capable devices, it is probably better to use C++ classes instead of raw pointers where possible.

Are you aware of the cases like that and how severe they might be?

Control flow, stored/loaded pointers from memory, arrays of pointers in different address spaces. I suspect these cases are pretty common for generic C++ code, and there is a runtime overhead for supporting them. User can optimize them manually (by using sycl::multi_ptr class), or rely on the compiler.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In my original code I didn't use generic pointer. By adding generic you removed the ability to diagnose the behaviour that is illegal i.e. assigning local address space pointer into global.

Even in OpenCL 2.0 compatible devices address spaces are used for performance reasons so developers might want to avoid generic as much as possible. Then it is really helpful to get diagnostics from the compiler as debugging such issues can be very painful to nearly impossible on accelerators.

Ok I generally understand the reasoning behind your design. I just find it a bit unfortunate that you remove address spaces from AST early so you can't benefit from the power of semantical analysis fully. I don't know whether developers will find it valuable that the compiler compiles everything but then they get little help in what goes wrong.

other.getAddressSpace() == LangAS::opencl_local ||
other.getAddressSpace() == LangAS::opencl_global));
}

/// Determines if these qualifiers compatibly include another set.
Expand Down
8 changes: 0 additions & 8 deletions clang/include/clang/Basic/AddressSpaces.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,14 +42,6 @@ enum class LangAS : unsigned {
cuda_constant,
cuda_shared,

sycl_global,
sycl_local,
sycl_constant,
sycl_private,
// Likely never used, but useful in the future to reserve the spot in the
// enum.
sycl_generic,

// Pointer size and extension address spaces.
ptr32_sptr,
ptr32_uptr,
Expand Down
2 changes: 0 additions & 2 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10292,8 +10292,6 @@ def err_builtin_launder_invalid_arg : Error<
"'__builtin_launder' is not allowed">;

// SYCL-specific diagnostics
def err_sycl_attribute_address_space_invalid : Error<
"address space is outside the valid range of values">;
def err_sycl_kernel_name_class_not_top_level : Error<
"kernel name class and its template argument classes' declarations can only "
"nest in a namespace: %0">;
Expand Down
18 changes: 0 additions & 18 deletions clang/include/clang/Sema/ParsedAttr.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -534,24 +534,6 @@ class ParsedAttr final
}
}

/// If this is an OpenCL addr space attribute returns its SYCL representation
/// in LangAS, otherwise returns default addr space.
LangAS asSYCLLangAS() const {
switch (getKind()) {
case ParsedAttr::AT_OpenCLConstantAddressSpace:
return LangAS::sycl_constant;
case ParsedAttr::AT_OpenCLGlobalAddressSpace:
return LangAS::sycl_global;
case ParsedAttr::AT_OpenCLLocalAddressSpace:
return LangAS::sycl_local;
case ParsedAttr::AT_OpenCLPrivateAddressSpace:
return LangAS::sycl_private;
case ParsedAttr::AT_OpenCLGenericAddressSpace:
default:
return LangAS::Default;
}
}

AttributeCommonInfo::Kind getKind() const { return getParsedKind(); }
};

Expand Down
5 changes: 0 additions & 5 deletions clang/lib/AST/ASTContext.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -836,11 +836,6 @@ static const LangASMap *getAddressSpaceMap(const TargetInfo &T,
5, // cuda_device
6, // cuda_constant
7, // cuda_shared
1, // sycl_global
3, // sycl_local
2, // sycl_constant
0, // sycl_private
4, // sycl_generic
8, // ptr32_sptr
9, // ptr32_uptr
10 // ptr64
Expand Down
4 changes: 0 additions & 4 deletions clang/lib/AST/TypePrinter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1792,16 +1792,12 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
case LangAS::Default:
return "";
case LangAS::opencl_global:
case LangAS::sycl_global:
return "__global";
case LangAS::opencl_local:
case LangAS::sycl_local:
return "__local";
case LangAS::opencl_private:
case LangAS::sycl_private:
return "__private";
case LangAS::opencl_constant:
case LangAS::sycl_constant:
return "__constant";
case LangAS::opencl_generic:
return "__generic";
Expand Down
10 changes: 0 additions & 10 deletions clang/lib/Basic/Targets/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,11 +48,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
Global, // cuda_device
Constant, // cuda_constant
Local, // cuda_shared
Global, // sycl_global
Local, // sycl_local
Constant, // sycl_constant
Private, // sycl_private
Generic, // sycl_generic
Generic, // ptr32_sptr
Generic, // ptr32_uptr
Generic // ptr64
Expand All @@ -68,11 +63,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
Global, // cuda_device
Constant, // cuda_constant
Local, // cuda_shared
Global, // sycl_global
Local, // sycl_local
Constant, // sycl_constant
Private, // sycl_private
Generic, // sycl_generic
Generic, // ptr32_sptr
Generic, // ptr32_uptr
Generic // ptr64
Expand Down
6 changes: 0 additions & 6 deletions clang/lib/Basic/Targets/NVPTX.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,12 +33,6 @@ static const unsigned NVPTXAddrSpaceMap[] = {
1, // cuda_device
4, // cuda_constant
3, // cuda_shared
1, // sycl_global
3, // sycl_local
4, // sycl_constant
0, // sycl_private
// FIXME: generic has to be added to the target
0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0 // ptr64
Expand Down
18 changes: 3 additions & 15 deletions clang/lib/Basic/Targets/SPIR.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,11 +33,6 @@ static const unsigned SPIRAddrSpaceMap[] = {
0, // cuda_device
0, // cuda_constant
0, // cuda_shared
1, // sycl_global
3, // sycl_local
2, // sycl_constant
0, // sycl_private
4, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0 // ptr64
Expand All @@ -53,11 +48,6 @@ static const unsigned SYCLAddrSpaceMap[] = {
0, // cuda_device
0, // cuda_constant
0, // cuda_shared
1, // sycl_global
3, // sycl_local
2, // sycl_constant
0, // sycl_private
4, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0 // ptr64
Expand All @@ -70,11 +60,9 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo {
TLSSupported = false;
VLASupported = false;
LongWidth = LongAlign = 64;
if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
AddrSpaceMap = &SYCLAddrSpaceMap;
} else {
AddrSpaceMap = &SPIRAddrSpaceMap;
}
AddrSpaceMap = (Triple.getEnvironment() == llvm::Triple::SYCLDevice)
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems like an unrelated change. A good one, just not related.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not sure how to interpret this...
Do you want to me revert it? Add to the PR description? No actions implied?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, sorry for not being clear. I guess the answer is: It depends on where this lands. I'd rather you do a separate commit for this formatting change here, but if you're submitting to llvm community, I'd rather the ternary version just be part of the patch.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This patch is to SYCL project.
First, I'd like to commit it here and see if it breaks anything. So far internal testing do not exposed any issues.
In addition to that I'd like to get high-level feedback on the approach before moving to the LLVM project.

WRT using different address space map for SYCL. I don't know if we break anything by mapping "default" address space to "4", which SPIR-V converter interprets as generic. It might that other programming models targeting SPIR rely on existing mapping.
In theory OpenCL is not supposed to use "default" address space, so it should be okay to use SYCLAddrSpaceMap in all cases.
As it was a separate patch, I suggest doing this experiment separately. Does it sounds okay to you?

Tagging other SYCL implementers: @Naghasan, @keryell, @illuhad,

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Having lang address space in target information seems strange. Ideally the IR should be emitted as language agnostic as possible. What are you trying to achieve with this separate address space map?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Having lang address space in target information seems strange. Ideally the IR should be emitted as language agnostic as possible. What are you trying to achieve with this separate address space map?

I'm trying to map "pointer w/o address space qualifiers" to "generic" address space. Currently it's mapped to "private", which doesn't seems right. As mentioned in the previous comment if all other languages do not rely on that behavior, I can remove "language dependent" customization.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Well OpenCL 2.0 does exactly the same. All pointer are in generic address space. Why don't you just add generic on AST while parsing the types?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@erichkeane answered this in this comment: https://github.com/intel/llvm/pull/1039/files#r370236039.

? &SYCLAddrSpaceMap
: &SPIRAddrSpaceMap;
UseAddrSpaceMapMangling = true;
HasLegalHalfType = true;
HasFloat16 = true;
Expand Down
6 changes: 0 additions & 6 deletions clang/lib/Basic/Targets/TCE.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,12 +40,6 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
0, // cuda_device
0, // cuda_constant
0, // cuda_shared
3, // sycl_global
4, // sycl_local
5, // sycl_constant
0, // sycl_private
// FIXME: generic has to be added to the target
0, // sycl_generic
0, // ptr32_sptr
0, // ptr32_uptr
0, // ptr64
Expand Down
5 changes: 0 additions & 5 deletions clang/lib/Basic/Targets/X86.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,11 +32,6 @@ static const unsigned X86AddrSpaceMap[] = {
0, // cuda_device
0, // cuda_constant
0, // cuda_shared
0, // sycl_global
0, // sycl_local
0, // sycl_constant
0, // sycl_private
0, // sycl_generic
270, // ptr32_sptr
271, // ptr32_uptr
272 // ptr64
Expand Down
28 changes: 3 additions & 25 deletions clang/lib/Sema/SemaType.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5963,35 +5963,14 @@ static bool BuildAddressSpaceIndex(Sema &S, LangAS &ASIdx,
llvm::APSInt max(addrSpace.getBitWidth());
max =
Qualifiers::MaxAddressSpace - (unsigned)LangAS::FirstTargetAddressSpace;

if (addrSpace > max) {
S.Diag(AttrLoc, diag::err_attribute_address_space_too_high)
<< (unsigned)max.getZExtValue() << AddrSpace->getSourceRange();
return false;
}

if (S.LangOpts.SYCLIsDevice && (addrSpace >= 4)) {
S.Diag(AttrLoc, diag::err_sycl_attribute_address_space_invalid)
<< AddrSpace->getSourceRange();
return false;
}

ASIdx = getLangASFromTargetAS(
static_cast<unsigned>(addrSpace.getZExtValue()));

if (S.LangOpts.SYCLIsDevice) {
ASIdx =
[](unsigned AS) {
switch (AS) {
case 0: return LangAS::sycl_private;
case 1: return LangAS::sycl_global;
case 2: return LangAS::sycl_constant;
case 3: return LangAS::sycl_local;
case 4: default: llvm_unreachable("Invalid SYCL AS");
}
}(static_cast<unsigned>(ASIdx) -
static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
}
ASIdx =
getLangASFromTargetAS(static_cast<unsigned>(addrSpace.getZExtValue()));
return true;
}

Expand Down Expand Up @@ -6117,8 +6096,7 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
Attr.setInvalid();
} else {
// The keyword-based type attributes imply which address space to use.
ASIdx = S.getLangOpts().SYCLIsDevice ?
Attr.asSYCLLangAS() : Attr.asOpenCLLangAS();
ASIdx = Attr.asOpenCLLangAS();
if (ASIdx == LangAS::Default)
llvm_unreachable("Invalid address space");

Expand Down
39 changes: 2 additions & 37 deletions clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ void foo(int * Data) {}
// CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
void foo2(int * Data) {}
// CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
void foo(__attribute__((address_space(3))) int * Data) {}
void foo(__attribute__((opencl_local)) int * Data) {}
// CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* %

template<typename T>
Expand All @@ -18,12 +18,11 @@ void tmpl(T t){}

void usages() {
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)*
__attribute__((address_space(1))) int *GLOB;
__attribute__((opencl_global)) int *GLOB;
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)*
__attribute__((opencl_local)) int *LOC;
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)*
int *NoAS;

// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32*
__attribute__((opencl_private)) int *PRIV;

Expand Down Expand Up @@ -94,57 +93,23 @@ void usages() {
// CHECK-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* %

void usages2() {
__attribute__((address_space(0))) int *PRIV_NUM;
// CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32*
__attribute__((address_space(0))) int *PRIV_NUM2;
// CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32*
__attribute__((opencl_private)) int *PRIV;
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32*
__attribute__((address_space(1))) int *GLOB_NUM;
// CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
__attribute__((opencl_global)) int *GLOB;
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
__attribute__((address_space(2))) int *CONST_NUM;
// CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
__attribute__((opencl_constant)) int *CONST;
// CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
__attribute__((address_space(3))) int *LOCAL_NUM;
// CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*
__attribute__((opencl_local)) int *LOCAL;
// CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*

bar(*PRIV_NUM);
// CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]]
// CHECK-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]])
bar(*PRIV_NUM2);
// CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]]
// CHECK-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]])
bar(*PRIV);
// CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]]
// CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]])
bar(*GLOB_NUM);
// CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]]
// CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]])
bar(*GLOB);
// CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]]
// CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]])
bar(*CONST_NUM);
// CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]]
// CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]])
bar(*CONST);
// CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]]
// CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]])
bar2(*LOCAL_NUM);
// CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]]
// CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)*
// CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]])
bar2(*LOCAL);
// CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]]
// CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)*
Expand Down
4 changes: 2 additions & 2 deletions clang/test/SemaOpenCLCXX/address-space-lambda.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,8 +31,8 @@ __kernel void test_qual() {
//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
auto priv2 = []() __generic {};
priv2();
auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
priv3(); //expected-error{{no matching function for call to object of type}}
auto priv3 = []() __global {}; //ex pected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //ex pected-note{{conversion candidate of type 'void (*)()'}}
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is happening here? The change to expected disables these tests, right? We dont want that, do we?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is happening here?

I'm not sure. I hope Anastasia can help to understand. #1039 (review).

The change to expected disables these tests, right?

Yes.

We dont want that, do we?

I'd like to understand why this patch affected the test and fix either the test or OpenCL mode.

priv3(); //ex pected-error{{no matching function for call to object of type}}
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@AnastasiaStulova, I'd like to check with you if this is a bug in C++ for OpenCL compiler or not.
NOTE, this patch enables conversion to from/to unqualified types and as OpenCL qualifies all the types it should not be affected.
All OpenCL tests pass except this one test case.
Is it a bug in function pointer type inference for lambdas?

Copy link
Contributor

@AnastasiaStulova AnastasiaStulova Jan 23, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Can you explain why you think this is a bug?

In OpenCL all local variables are deduced to be in __private addr space so it is the case for function object priv3, however its call operator has __global qualifier and hence can only be used with function objects in __global addr space because no other addr space can implicitly be converted to __global. So giving an error on the call to this call operator seems correct.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the issue might be due to your downstream change in isAddressSpaceSupersetOf that I don't entirely understand.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

isAddressSpaceSupersetOf allow conversion between unqualified types and types qualified with opencl_global, opencl_local and opencl_private address space qualifiers.
As I mentioned above it should have zero impact on OpenCL compiler, which is supposed to qualify ALL types with address space qualifier.

In OpenCL all local variables are deduced to be in __private addr space so it is the case for function object priv3, however its call operator has __global qualifier and hence can only be used with function objects in __global addr space because no other addr space can implicitly be converted to __global. So giving an error on the call to this call operator seems correct.

It doesn't seem to work as you described if address space is deduced correctly, the mismatch will still be reported. It looks like compiler relies somehow on isAddressSpaceSupersetOf in deduction process and priv3 type is not qualified with the address space. I'll check this assumption.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It doesn't seem to work as you described if address space is deduced correctly, the mismatch will still be reported.

Sorry what doesn't work as I describe?

It looks like compiler relies somehow on isAddressSpaceSupersetOf in deduction process and priv3 type is not qualified with the address space. I'll check this assumption.

isAddressSpaceSupersetOf is not used in deduction it is used in call operator overloading resolution.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do you know where isAddressSpaceSupersetOf called from for this line?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It's not how lambda parsing works in C++. The initializer is parsed separately and there is no danger of declaring it this way unless it's being called.

I would suggest to check the comments on the review too that might help to clarify the topic:
https://reviews.llvm.org/D70242
https://reviews.llvm.org/D69938

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

https://reviews.llvm.org/D69938 - the discussion started with an assumption that deduced address space for lambda object can be changed to legalize function call operator usages - https://reviews.llvm.org/D69938#1737196, but ended with an open question if it's really possible - https://reviews.llvm.org/D69938#1755709.
Later the conclusion was to use "opencl_generic" for temporaries rather than "default" - https://reviews.llvm.org/D69938#1759241. Why does current implementation use default address space instead of opencl_generic?

It's still not clear whether it's the right approach to detect an error at the function call operator use or it should be done for lambda variable declaration. The case where address space of a lambda object is changed to make use of inconsistent declaration is not covered by tests.

Copy link
Contributor

@Fznamznon Fznamznon Apr 3, 2020 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Alright, I dug in. This fail basically happens because It seems that lambdas have predefined conversions to function pointer form:

`-LambdaExpr 0x563552014f70 <col:16, col:31> '(lambda at address-space-lambda.cl:34:16)'
  |-CXXRecordDecl 0x563552014928 <col:16> col:16 implicit class definition
  | |-DefinitionData lambda pass_in_registers empty standard_layout trivially_copyable literal can_const_default_init
  | | |-DefaultConstructor defaulted_is_constexpr
  | | |-CopyConstructor simple trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveConstructor exists simple trivial needs_implicit
  | | |-CopyAssignment trivial has_const_param needs_implicit implicit_has_const_param
  | | |-MoveAssignment
  | | `-Destructor simple irrelevant trivial
  | |-CXXMethodDecl 0x563552014a60 <col:19, col:31> col:16 used constexpr operator() 'void () const __global' inline
  | | `-CompoundStmt 0x563552014b10 <col:30, col:31>
  | |-[[THIS ONE]] CXXConversionDecl 0x563552014e08 <col:16> col:16 implicit used constexpr operator void (*)() 'void (*() const noexcept)()' inline
  | | `-CompoundStmt 0x563552044e78 <col:16>
  | |   `-ReturnStmt 0x563552044e68 <col:16>
  | |     `-ImplicitCastExpr 0x563552044e50 <col:16> 'void (*)()' <FunctionToPointerDecay>
  | |       `-DeclRefExpr 0x563552044e30 <col:16> 'void ()' lvalue CXXMethod 0x563552014eb8 '__invoke' 'void ()'
  | |-CXXMethodDecl 0x563552014eb8 <col:16> col:16 implicit used __invoke 'void ()' static inline
  | | `-CompoundStmt 0x563552044e20 <col:16>
  | `-CXXDestructorDecl 0x563552014fa0 <col:16> col:16 implicit referenced ~ 'void () __generic noexcept' inline default trivial
  `-CompoundStmt 0x563552014b10 <col:30, col:31>

C++ requres it:

The closure type for a lambda-expression with no lambda-capture has a public non-virtual non-explicit const conversion function to pointer to function having the same parameter and return types as the closure type’s function call operator.

But as you can see from the AST dump this conversion operator converts this lambda to a poiner to a function without address space qualifiers at all, i,e, everything what this function accepts can be in Default address space. Whereas calling operator of such lambda has __global qualifier.
After we allow conversion between between Default address space and others this conversion is used by compiler to make such call of this lambda valid:

`-CallExpr 0x563552044ec8 <line:35:3, col:9> 'void':'void'
  `-ImplicitCastExpr 0x563552044eb0 <col:3> 'void (*)()' <UserDefinedConversion>
    `-CXXMemberCallExpr 0x563552044e90 <col:3> 'void (*)()'
      `-MemberExpr 0x563552044df0 <col:3> '<bound member function type>' .operator void (*)() [[THE CONVERSION!!!!]] 0x563552014e08
        `-ImplicitCastExpr 0x563552044dd8 <col:3> 'const (lambda at address-space-lambda.cl:34:16)' lvalue <AddressSpaceConversion>
          `-DeclRefExpr 0x563552044d70 <col:3> '__private (lambda at address-space-lambda.cl:34:16)':'__private (lambda at address-space-lambda.cl:34:16)' lvalue Var 0x563552014820 'priv3' '__private (lambda

I'm not an expert in OpenCL and C++ for OpenCL but, if the philosophy of OpenCL address spaces is qualify ALL types with address space qualifier and if C++ for OpenCL should work like regular C++ for language features like lambdas, I think conversion from lambda object to function pointer must convert to a pointer to a function with the same address space qualifiers as operator () has. So, IMO this is a bug in C++ for OpenCL. @AnastasiaStulova , could you please provide your opinion on this?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Submitted a bug https://bugs.llvm.org/show_bug.cgi?id=45472 .


__constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}}
const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
Expand Down
6 changes: 1 addition & 5 deletions clang/test/SemaSYCL/address-space-parameter-conversions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@ void tmpl(T *t){}
void usages() {
__attribute__((opencl_global)) int *GLOB;
__attribute__((opencl_private)) int *PRIV;
__attribute__((address_space(3))) int *LOC;
__attribute__((opencl_local)) int *LOC;
int *NoAS;

bar(*GLOB);
Expand Down Expand Up @@ -53,10 +53,6 @@ void usages() {

// expected-error@+1{{address space is negative}}
__attribute__((address_space(-1))) int *TooLow;
// expected-error@+1{{address space is outside the valid range of values}}
__attribute__((address_space(6))) int *TooHigh;
// expected-error@+1{{address space is outside the valid range of values}}
__attribute__((address_space(4))) int *TriedGeneric;
// expected-error@+1{{unknown type name '__generic'}}
__generic int *IsGeneric;

Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaTemplate/address_space-dependent.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ void neg() {

template <long int I>
void tooBig() {
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388590)}}
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388595)}}
}

template <long int I>
Expand Down
Loading