intel / intel/llvm

request: extend the sycl_ext_oneapi_device_global extension

Open
#9,965 9 comments 0 reactions 0 assignees View on GitHub
enhancement spec 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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.