Is it possible to copy a class containing pointers to its internal attribute using SYCL and offload it to the graphics card? Basically, I try to reference members to avoid unnecessary memory usage. I managed to have my example run on local memory but when offloading to a graphics card, problems occurred
PI CUDA ERROR:
Value: 700
Name: CUDA_ERROR_ILLEGAL_ADDRESS
Description: an illegal memory access was encountered
Function: operator()
Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:2612
SYCL exception: Native API failed. Native API returns: -999 (Unknown PI error) -999 (Unknown PI error)
PI CUDA ERROR:
Value: 700
Name: CUDA_ERROR_ILLEGAL_ADDRESS
Description: an illegal memory access was encountered
Function: wait
Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:653
PI CUDA ERROR:
Value: 700
Name: CUDA_ERROR_ILLEGAL_ADDRESS
Description: an illegal memory access was encountered
Function: wait
Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:653
PI CUDA ERROR:
Value: 700
Name: CUDA_ERROR_ILLEGAL_ADDRESS
Description: an illegal memory access was encountered
Function: get_next_transfer_stream
Source Location: /root/intel-llvm-mirror/sycl/plugins/cuda/pi_cuda.cpp:513
I understand that this error is related to copying a memory pointer to another device, but given that I've provided an explicit class copy constructor, I would expect the copy to be a deep copy. Does SYCL fundamentally lack support for copying a pointer object to a graphics card? I've attached my code below
#include <array>
#include <iostream>
#include <sycl/sycl.hpp>
#include <vector>
using namespace sycl; // (optional) avoids the need for "sycl::" before SYCL names
typedef struct my1 {
int value = 1;
// Copy constructor that creates a new my1
my1(const my1& other) : value(other.value) {}
// Default constructor
my1() = default;
// Constructor with a value
my1(int v) : value(v) {}
} my1;
typedef struct my2 {
struct my1* a;
struct my1* b;
} my2;
template <std::size_t N>
class MyClass {
public:
MyClass(std::array<my1, N> a, std::array<my1, N> b) {
for (std::size_t i = 0; i < N; i++) {
_a[i].value = a[i].value;
_b[i].value = b[i].value;
_c[i].a = &_a[i];
_c[i].b = &_b[i];
}
}
float addValue(int i) const {
if (i >= N || i >= N)
// throw std::runtime_error("Index out of bounds");
return 0;
return (_c[i].a)->value + (_c[i].b)->value;
}
void modifyAValue(int i, float a) {
if (i >= N)
return;
_a[i].value = a;
}
void modifyBValue(int i, float b) {
if (i >= N)
return;
_b[i].value = b;
}
MyClass(const MyClass& other) {
_c = std::array<my2, N>();
for (std::size_t i = 0; i < N; i++) {
_a[i].value = other._a[i].value;
_b[i].value = other._b[i].value;
_c[i].a = &_a[i];
_c[i].b = &_b[i];
}
}
private:
std::array<my1, N> _a;
std::array<my1, N> _b;
std::array<my2, N> _c;
};
template <std::size_t N>
struct sycl::is_device_copyable<MyClass<N>> : std::true_type {};
int main() {
// Create a SYCL queue
sycl::queue myQueue;
// Create an instance of MyClass
int N = 10;
std::array<my1, 10> a{1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
std::array<my1, 10> b{10, 9, 8, 7, 6, 5, 4, 3, 2, 1};
MyClass<a.size()> myObject(a, b);
std::vector<int> input = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9};
std::vector<float> result(input.size());
// Create a SYCL buffer to hold the result
sycl::buffer<float, 1> resultBuffer(result.data(), sycl::range<1>(input.size()));
sycl::buffer<int, 1> inputBuffer(input.data(), sycl::range<1>(input.size()));
sycl::buffer<MyClass<10>> buffer(&myObject, 1);
try {
// Submit a SYCL kernel
myQueue.submit([&](sycl::handler& cgh) {
auto resultAcc = resultBuffer.template get_access<sycl::access::mode::write>(cgh);
auto inputAcc = inputBuffer.template get_access<sycl::access::mode::read>(cgh);
auto myObjectAcc = buffer.template get_access<sycl::access::mode::read>(cgh);
cgh.parallel_for<class MyKernel>(
sycl::range<1>(input.size()), [=](sycl::id<1> idx) {
// Call a member function on the instance of MyClass
resultAcc[idx] = myObjectAcc[0].addValue(inputAcc[idx]);
});
});
// Wait for the kernel to complete and get the result
myQueue.wait();
myQueue.update_host(resultBuffer.get_access());
} catch (const sycl::exception& e) {
std::cerr << "SYCL exception: " << e.what() << std::endl;
return 1;
} catch (const std::runtime_error& e) {
std::cerr << "Runtime error: " << e.what() << std::endl;
return 1;
}
// Print the result (accumulated sum)
float sum = 0.0f;
for (float val : result) {
std::cout << val << std::endl;
sum += val;
}
std::cout << "Result: " << sum << std::endl;
return 0;
}
In this example, I have a class that contains three arrays. Two of these arrays store information (my1), and the third one stores references to elements from the two information arrays (my2, which holds two pointers to my1 objects).
std::array<my1, N> _a;
std::array<my1, N> _b;
std::array<my2, N> _c;
As I understand it, errors occur because the kernel is attempting to access host memory. This is unexpected, as the constructor should have reassigned the pointer value after the copy operations.
MyClass(const MyClass& other) {
_c = std::array<my2, N>();
for (std::size_t i = 0; i < N; i++)
{
_a[i].value = other._a[i].value;
_b[i].value = other._b[i].value;
_c[i].a = &_a[i];
_c[i].b = &_b[i];
}
I wouldn't mind using USM to allocate device memory, but I would prefer to use pointers or references to access the member variables inside the class. This approach helps avoid unnecessary memory allocation by referencing device memory directly (in short, I'd like to use pointers or references as member variable to access device memory)
After some research, I find it is impossible to emerge into SYCL without using usm or buffer below I have put the example code for this question.