request: extend the sycl_ext_oneapi_device_global extension
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
I'm working on the [Alpaka portability library](https://github.com/alpaka-group/alpaka/) to add support for Intel devices via oneAPI, and I'm looking at the [sycl_ext_oneapi_device_global](https://github.com/intel/llvm/blob/15edb1bdd3cca0d719c2fca546ea8ddc4d26dc6f/sycl/doc/extensions/proposed/sycl_ext_oneapi_device_global.asciidoc) extension to implement the functionality of CUDA `__device__` global variables.
So far it looks like the extension is missing two features:
- the possibility of initialising a variable
- the possibility of getting a per-device (per-context) pointer to the device data
Both are possible using CUDA `__device__` variables, as showcased here:
```c++
#include
#include
__device__ int global = -1;
int main() {
int count = -1;
cudaGetDeviceCount(&count);
std::cout << "Found " << count << " devices\n";
for (int device = 0; device < count; ++device) {
std::cout << '\n';
std::cout << "using device " << device << '\n';
cudaSetDevice(device);
int* address = nullptr;
cudaGetSymbolAddress((void**) &address, global);
std::cout << "the address of `global` on device " << device << " is " << address << '\n';
int value = -1;
// equivalent to
// cudaMemcpyFromSymbol(&value, global, sizeof(device), 0, cudaMemcpyDeviceToHost);
cudaMemcpy(&value, (void**) address, sizeof(device), cudaMemcpyDeviceToHost);
std::cout << "the value of `global` on device " << device << " is " << value << '\n';
std::cout << "setting the value of `global` on device " << device << " to " << device << '\n';
// equivalent to
// cudaMemcpyToSymbol(global, &device, sizeof(device), 0, cudaMemcpyHostToDevice);
cudaMemcpy((void**) address, &device, sizeof(device), cudaMemcpyHostToDevice);
}
for (int device = 0; device < count; ++device) {
std::cout << '\n';
std::cout << "using device " << device << '\n';
cudaSetDevice(device);
int value = -1;
cudaMemcpyFromSymbol(&value, global, sizeof(device), 0, cudaMemcpyDeviceToHost);
std::cout << "the value of `global` on device " << device << " is " << value << '\n';
}
}
```
On a machine with two NVIDIA GPUs this prints something like
```
Found 2 devices
using device 0
the address of `global` on device 0 is 0x7f4706800000
the value of `global` on device 0 is -1
setting the value of `global` on device 0 to 0
using device 1
the address of `global` on device 1 is 0x7f46f4800000
the value of `global` on device 1 is -1
setting the value of `global` on device 1 to 1
using device 0
the value of `global` on device 0 is 0
using device 1
the value of `global` on device 1 is 1
```
Would you consider extending the sycl_ext_oneapi_device_global extension to support these ?
Contributor guide
Assessment
This issue has not been assessed yet.