sycl/dpc++ accessing global variable using global pointer

Question

I am trying to access a global variable from inside the sycl kernel. A code that uses this pattern and its output is as follows.

#include<CL/sycl.hpp>
using namespace sycl;
int g_var = 22;
int * const g_ptr = &g_var;
int main() {
    queue q{host_selector{}};
    g_var = 27;
    std::cout<<"global var changed : " <<g_var<<" "<<*g_ptr<<"\n";
    q.submit( [=] (handler &h) {
        stream os(1024, 128, h);
        h.single_task([=] () {
            os<<"global var : "<<*g_ptr<<"\n";
        });
    }).wait();
    return 0;
}

Its output is as follows.

$$ dpcpp test.cpp; ./a.out
global var changed : 27 27
global var : 22

Even though g_var is changed to 27, it is printed with the initial value of 22 inside the kernel. Is this the intended behavior?

In general lambdas does not create a copy of global variables. Does dpc++ compiler create a copy of the global variable inside the device or does it propagate the constant value during compilation so as to not access the global memory at runtime?

Answer 1

SYCL divides execution and data through the "host" and "device." The data on the "host" lives on the CPU normally, and needs to be transferred to the "device" (GPU usuall) in order to be accessed for use in a SYCL kernel.

SYCL uses buffers and accessors, or "Unified Shared Memory" to transfer and provide access to data for the "device" side code (i.e. the code running on the GPU). So your pointer is never sent over to the device, and therefore never modified.

So instead of using int * const g_ptr = &g_var; you should use a buffer or USM. There are some lessons on these for buffers and USM in SYCL Academy.

E.g. using buffers and accessors:

int a = 18, b = 24, r = 0;

  auto defaultQueue = sycl::queue{};

  {
    auto bufA = sycl::buffer{&a, sycl::range{1}};
    auto bufB = sycl::buffer{&b, sycl::range{1}};
    auto bufR = sycl::buffer{&r, sycl::range{1}};

    defaultQueue
        .submit([&](sycl::handler &cgh) {
          auto accA = sycl::accessor{bufA, cgh, sycl::read_only};
          auto accB = sycl::accessor{bufB, cgh, sycl::read_only};
          auto accR = sycl::accessor{bufR, cgh, sycl::write_only};

          cgh.single_task<scalar_add>([=] { accR[0] = accA[0] + accB[0]; });
        })
        .wait();
  }

And using USM:

 auto usmQueue = sycl::queue{usm_selector{}, asyncHandler};

 usmQueue.memcpy(devicePtrA, a, sizeof(float) * dataSize).wait();
 usmQueue.memcpy(devicePtrB, b, sizeof(float) * dataSize).wait();

 usmQueue
     .parallel_for<vector_add>(sycl::range{dataSize},
                                  [=](sycl::id<1> idx) {
                                    auto globalId = idx[0];
                                    devicePtrR[globalId] =
                                        devicePtrA[globalId] +
                                        devicePtrB[globalId];
                                  })
     .wait();

 usmQueue.memcpy(r, devicePtrR, sizeof(float) * dataSize).wait();

 sycl::free(devicePtrA, usmQueue);
 sycl::free(devicePtrB, usmQueue);
 sycl::free(devicePtrR, usmQueue);

 usmQueue.throw_asynchronous();