Link Copied
Hi Pendlaya,
Apologies for the delayed response. I am wondering if my problem is more fundamental, as I have not been able to overlap copies and kernels even with a single, non-FIFO queue.
The code is as follows:
#include<iostream>
#include<chrono>
#include <sycl/CL/sycl.hpp>
#include <sycl/usm.hpp>
#include <level_zero/ze_api.h>
#include <sycl/ext/oneapi/backend/level_zero.hpp>
cl::sycl::queue *accelerator;
#define accelerator_for2dNB( iter1, num1, iter2, num2, ... ) \
accelerator->submit([&](cl::sycl::handler &cgh) { \
unsigned long nt=8; \
unsigned long unum1 = num1; \
unsigned long unum2 = num2; \
\
unsigned long unum1_use = ((unum1 + nt - 1)/nt) * nt ; /*round up s.t. divisible by nt*/ \
cl::sycl::range<3> local {nt,1,1}; \
cl::sycl::range<3> global{unum1_use,unum2,1}; \
cgh.parallel_for( \
cl::sycl::nd_range<3>(global,local), \
[=] (cl::sycl::nd_item<3> item) /*mutable*/ \
[[intel::reqd_sub_group_size(16)]] \
{ \
auto iter1 = item.get_global_id(0); \
auto iter2 = item.get_global_id(1); \
auto lane = item.get_global_id(2); \
if(iter1<unum1){ __VA_ARGS__ }; \
}); \
});
#define accelerator_barrier(){ accelerator->wait(); }
inline void acceleratorCopySynchronise(void) { accelerator->wait(); }
inline void acceleratorBarrierAll(){ accelerator_barrier(); }
inline void *acceleratorAllocDevice(size_t bytes){ return malloc_device(bytes,*accelerator);};
inline void acceleratorFreeDevice(void *ptr){free(ptr,*accelerator);};
inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes) { accelerator->memcpy(to,from,bytes);}
inline void acceleratorCopyToDevice(void *from,void *to,size_t bytes) { accelerator->memcpy(to,from,bytes); accelerator->wait();}
inline void acceleratorCopyFromDevice(void *from,void *to,size_t bytes){ accelerator->memcpy(to,from,bytes); accelerator->wait();}
inline void acceleratorCopyToDeviceAsync(void *from,void *to,size_t bytes) { accelerator->memcpy(to,from,bytes); }
int main(int argc, char **argv){
std::cout << "START" << std::endl;
int nDevices = 1;
cl::sycl::gpu_selector selector;
cl::sycl::device selectedDevice { selector };
accelerator = new sycl::queue (selectedDevice);
size_t n = 10000000;
size_t b = n*sizeof(double);
double* dd = (double*)acceleratorAllocDevice(b);
double* dh = (double*)malloc(b);
double* rd1 = (double*)acceleratorAllocDevice(10*b);
double* rh1 = (double*)malloc(10*b);
acceleratorCopyToDevice(dh,dd,b);
acceleratorBarrierAll();
for(int i=0;i<10;i++){
acceleratorCopyToDeviceAsync(rh1,rd1,10*b);
accelerator_for2dNB( i, n, dummy,1, {
for(int j=0;j<1000;j++)
dd[i] = dd[i]*dd[i] + j;
});
accelerator_for2dNB( i, n, dummy,1, {
for(int j=0;j<1000;j++)
dd[i] = dd[i]*dd[i] + j;
});
acceleratorBarrierAll();
}
//acceleratorFreeDevice(dd);
//free(dh);
std::cout << "END" << std::endl;
return 0;
}As you see, it launches a non-blocking memcpy to device as well as 2 kernels. In the output I observe the following pattern:
where we see that the memcpy and kernels simply refuse to overlap (the two kernels do, however). I am using the following environment variables:
export ONEAPI_DEVICE_SELECTOR=gpu,level_zero
export EnableImplicitScaling=0
export ForceThreadGroupDispatchSize=0
export SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=0
export SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1
export SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE=0and running on ALCF Sunspot with the default oneapi installation.
I would appreciate any help you can provide.
