Skip to content
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

[WIP] [SYCL] Move static const variables to constant address space #158

Closed
wants to merge 1 commit into from
Closed

[WIP] [SYCL] Move static const variables to constant address space #158

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
7 changes: 7 additions & 0 deletions clang/lib/CodeGen/CodeGenModule.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3396,6 +3396,13 @@ LangAS CodeGenModule::GetGlobalVarAddressSpace(const VarDecl *D) {
return AddrSpace;
}

if (LangOpts.SYCLIsDevice) {
if (D && D->getType().isConstQualified())
return LangAS::opencl_constant;

return LangAS::opencl_global;
Copy link
Contributor

Choose a reason for hiding this comment

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

We use sycl_* address spaces in all other places and with this change we mix two potentially different mappings in the single program.
I think we should first agree on using opencl_* address space map for SYCL device code and if so, replace all usages of sycl_* address spaces with opencl_* and remove sycl_* address spaces from the map.
@Naghasan, @keryell do you see any potential issues with using opencl_* address spaces?

Copy link
Contributor

@Naghasan Naghasan May 23, 2019 •
edited
Loading

Choose a reason for hiding this comment

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

TBH, I never fully understood the rational behind the introduction of the sycl_* address spaces. I think it was made to allow this c278543#diff-54e8e44f96ab947511166d795bad8022R477 but I find this suspicious as it is easy to trick the compiler (see #40).

Apart from the constant address space, I don't see any particular issue in moving from sycl_* asp to opencl_* asp.

The device compiler needs to convert the outlined code into an OpenCL kernel, so reusing LangAS::opencl_* seems a rational choice to me as it will play better with existing code in clang (especially that it will need to obey to the same set of rules). For the library, the exposure of the address space is pushed to the multi_ptr and accessor classes. You will probably have to do adjustments in the header files, but those classes are also there to handle the host/device boundary in a transparent way for the user.

That being said, there is the question of handling the (OpenCL) constant address space. As you can't make it overlap with the generic address space, this becomes tricky. But that's an exception that can be approached in different ways.

Note on the patch itself

considering this code

static const int bar = 42;

void work(int*);

[...]
cgh.single_task([=]() {
  int foo = 21;
  work(&foo);
  work(&bar);
});

this patch will tranform it into something like this opencl C code:

constant const int bar = 42;

void work(generic int*);

__kernel void ker() {
  int foo = 21;
  work(&foo); // ok
  work(&bar); // error: constant and generic address space does not overlap
});

Copy link
Contributor Author

Choose a reason for hiding this comment

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

// error: constant and generic address space does not overlap

Right, I noticed that as well. Please treat this patch as WIP for now.

Copy link
Contributor

Choose a reason for hiding this comment

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

NOTE: even if we use opencl_* it doesn't mean we enforce OpenCL rules. A lot of OpenCL rules are enabled only for OpenCL language, which we can't use in for SYCL unfortunately.
We might need to go over this checks that enable them for SYCL too by adding additional case or removing language check at all.
+@AnastasiaStulova

Copy link
Contributor

Choose a reason for hiding this comment

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

Btw, I don't think it is correct to deduce const to __constant. Variables that are const qualified can be initialized at RT. But __constant required compile time constant initializers.

In a way it is a bit closer to constexpr.

NOTE: even if we use opencl_* it doesn't mean we enforce OpenCL rules. A lot of OpenCL rules are enabled only for OpenCL language, which we can't use in for SYCL unfortunately.
We might need to go over this checks that enable them for SYCL too by adding additional case or removing language check at all.
+@AnastasiaStulova

I would prefer to remove OpenCL checks completely. They normally come from the old code anyway. I never understood why it was written like that. If we have OpenCL address spaces surely we are compiling OpenCL code (or may be SYCL device code now too!).

I was also thinking if it helps we could add some sort of common language mode (i.e. CLDevice) that we could use for common logic between OpenCL and SYCL. And then derive OpenCL and SYCL from CLDevice. Not sure how helpful it is though but it could be one options.

}

if (LangOpts.CUDA && LangOpts.CUDAIsDevice) {
if (D && D->hasAttr<CUDAConstantAttr>())
return LangAS::cuda_constant;
Expand Down
6 changes: 6 additions & 0 deletions clang/test/CodeGenSYCL/address-space-new.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,11 +3,17 @@


void test() {
static const int foo = 0x42;
// CHECK-DEFAULT: @_ZZ4testvE3foo = internal addrspace(2) constant i32 66, align 4
// CHECK-NEW: @_ZZ4testvE3foo = internal addrspace(2) constant i32 66, align 4

int i = 0;
int *pptr = &i;
// CHECK-DEFAULT: store i32* %i, i32** %pptr
// CHECK-NEW: %[[GEN:[0-9]+]] = addrspacecast i32* %i to i32 addrspace(4)*
// CHECK-NEW: store i32 addrspace(4)* %[[GEN]], i32 addrspace(4)** %pptr

*pptr = foo;
}


Expand Down