[SYCL][GPU] Reducing stack frame usage for the "ranges and index space identifiers"

**The Issue**
Actually, most of the SYCL objects such as `id`, `range`, `item`, `nd_range`, `nd_item` are implemented using [that](https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/detail/array.hpp) array class based on a stack array.

On GPUs, the arrays are stored in the stack frame which resides in the *global memory* and the latency is thus the worst. Optimisations described [here](https://developer.nvidia.com/blog/fast-dynamic-indexing-private-arrays-cuda/) are available if and only if the array is accessed using constant indices. If so, the array will be stored in registers and the access pretty much instantaneous. That is currently what's going on.

When using variable indices to access the array, that object is systematically sent to the stack frame. If it happens to be in a `struct` or `class`, the whole object will be stored on the stack.

**To Reproduce**

Let use the following code, compiled with the CUDA back-end and with `-O3`:

```C++
#include

class kernel_name;

struct my_struct { // sizeof == 80
int32_t i = 0;
int32_t j = 0;
sycl::id<3> some_coordinates{0, 0, 0};
sycl::id<3> some_more_coordinates{0, 0, 0};
sycl::id<3> even_more_coordinates{0, 0, 0};
};

int main() {
sycl::queue q = sycl::queue(sycl::gpu_selector{});
std::cout << "Sizeof my_struct " << sizeof(my_struct) << '\n'; // Sizeof my_struct 80
volatile auto ptr = sycl::malloc_shared(1, q);
q.submit([&](sycl::handler &cgh) {
cgh.single_task([=]() {
my_struct data{};
int i = *ptr;
*ptr = data.some_coordinates[3]; // 8 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
// *ptr = data.some_coordinates[i%3]; // 88 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
});
}).wait();
return 0;
}
```

When compiling with the line ` *ptr = data.some_coordinates[3];`, we get only 8 bytes of stack frame usage (which corresponds to the `ptr`). But if we access `id<3>` with `i%3`, the compiler cannot perform "registerization" (by assigning the array to the registers) as it don't know which register to access. The whole structure is sent to the stack and we get `88 bytes stack frame` from `ptxas`.

The same issue is happening with the sycl vectors, but it is probably less harmful.

**Solution(s)?**

I haven't found a solution that doesn't break the ABI and has no overhead. With ABI breakage I tried this : https://github.com/Michoumichmich/llvm/tree/register_array_optimisation which compiles down to the same kernel if the indices are constant, and if they are not, the data is still kept in the registers which solves the issue.
I haven't tested openCL on CPU.

**Environment (please complete the following information):**

- OS: RHEL
- Target device and vendor: GPU
- DPC++ version: c357af1c5701e418879160b44e6b006cbb8e2a5d