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

[asavonic]

Previously when ENASBLE_INFER_AS is enabled (see: 1a3a536 [SYCL]
Optionally override addrspace map for SYCL device code), static
variables were put into generic address space.

Now, if a static variable has a const qualifier, it will go to
constant address space.

Non-const variables with static storage are not supported by the SYCL
1.2.1 specification (see s6.3 "Language restrictions for kernels"),
but this patch handles this case as well by putting such variables
into global address space (we cannot do anything better in CodeGen
anyway).

Signed-off-by: Andrew Savonichev [email protected]

[bader]

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?

[Naghasan]

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
});

[asavonic]

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

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

[bader]

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

[AnastasiaStulova]

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.

[asavonic]

Ok, __constant address space is probably a bad choice here: it cannot be casted to generic, thus it breaks the code that have pointers const static. OpenCL doesn't do this as well: static const variables in OpenCL gets into global address space.

For example:

__kernel void k() {
  static const int foo = 0x221;
}

gets compiled to:

@k.foo = internal addrspace(1) constant i32 545, align 4

Obvious solution would be to follow the OpenCL approach here, but I need to verify that the rest of the code plays nicely with this.

[keryell]

Perhaps you can keep this as an optional feature, when you want to pack as much as you can in some __constant address-space for architectural limitation or performance reason...

[asavonic]

Implemented as part of #242