If the values of

lda

and

ldb

are the same, you can use the following code:

dpct::dpct_memcpy((void*)(B), (void*)(A), lda*cols*elemSize, dpct::host_to_device); // For cublasGetMatrix, use dpct::device_to_host

Otherwise, you can copy the columns of the matrix or elements of the vector one by one:

The equivalent of

cublasSetMatrix(rows, cols, elemSize, A, lda, B, ldb)

is:

auto A_backup = A; 
auto B_backup = B;
A = A - lda;
B = B - ldb;
for(int c = 0; c < cols; ++c) {
  A = A + lda;
  B = B + ldb;
  dpct::dpct_memcpy((void*)(B), (void*)(A), rows*elemSize, dpct::host_to_device); // For cublasGetMatrix, use dpct::device_to_host
}
A = A_backup;
B = B_backup;

If the values of

incx

and

incy

are the same, you can use the following code:

dpct::dpct_memcpy((void*)(B), (void*)(A), n*incx*elemSize, dpct::host_to_device); // For cublasGetVector, use dpct::device_to_host

Otherwise, you can copy the elements of the vector one by one:

To replace

cublasGetVector(n, elemSize, x, incx, y, incy)

, use the following snippet:

auto x_backup = x;
auto y_backup = y;
x = x - incx;
y = y - incy;
for(int c = 0; c < n; ++c) {
  x = x + incx;
  y = y + incy;
  dpct::dpct_memcpy((void*)(y), (void*)(x), elemSize, dpct::device_to_host); // for cublasSetVector, use dpct::host_to_device
}
x = x_backup;
y = y_backup;

The <parameter name> could not be evaluated and may be smaller than <other parameter name>.
The <parameter name> is equal to <other parameter name>, but is greater than 1.

handleError(cudaEventRecord(e));

e_ct1 = clock(), handleError(0);

DPC++
currently does not support setting resource limits on devices
DPC++
currently does not support associating USM with a specific queue
DPC++
currently does not support query operations on queues
DPC++
currently does not support capture operations on queues
DPC++
currently does not support configuring shared memory on devices

DPC++
currently does not support setting resource limits on devices
DPC++
currently does not support associating USM with a specific queue
DPC++
currently does not support query operations on queues
DPC++
currently does not support capture operations on queues
DPC++
currently does not support configuring shared memory on devices

XORWOW generator
Sobol generator of 64-bit sequences
Scrambled Sobol generator

... //declarations for device0 to device30
dpct::constant_memory<int, 1> device31(ARRAY_SIZE);
dpct::constant_memory<int, 1> device32(ARRAY_SIZE);

// kernel function declaration
void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32);

void test_function() {
  ... //memcpy operations for device0 to device30
  dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int));
  dpct::dpct_memcpy(device32.get_ptr(), host, ARRAY_SIZE * sizeof(int));
  {
    dpct::get_default_queue().submit([&](sycl::handler &cgh) {
    ... //accessors for device0 to device30, captured by SYCL kernel lambda
    auto device31_acc_ct1 = device31.get_access(cgh);
    auto device32_acc_ct1 = device32.get_access(cgh);
    cgh.parallel_for(
      sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)),
        [=](sycl::nd_item<3> item_ct1) {
          kernel(item_ct1, ... /*arguments device0 to device30*/
                 device31_acc_ct1.get_pointer(), device32_acc_ct1.get_pointer());
        });
    });
  }
}

... //declarations for device0 to device30
dpct::constant_memory<int, 1> device32(ARRAY_SIZE + ARRAY_SIZE); // one buffer instead of two

// kernel function declaration
void kernel(sycl::nd_item<3> item_ct1, int *device0, ..., int *device32);

void test_function() {
  ... //memcpy operations for device0 to device30
  dpct::dpct_memcpy(device31.get_ptr(), host, ARRAY_SIZE * sizeof(int));
  // memory copy with offset:
  dpct::dpct_memcpy(device31.get_ptr() + ARRAY_SIZE, host, ARRAY_SIZE * sizeof(int));
  {
    dpct::get_default_queue().submit([&](sycl::handler &cgh) {
    ... //accessors for device0 to device30, captured by SYCL kernel lambda
    auto device31_acc_ct1 = device31.get_access(cgh); // only one accessor instead of 2
    cgh.parallel_for(
      sycl::nd_range<3>(sycl::range<3>(1), sycl::range<3>(1)),
        [=](sycl::nd_item<3> item_ct1) {
          // last parameter is modified to use offset and the same accessor as previous parameter
          kernel(item_ct1, ... /*arguments device0 to device30*/
                 device31_acc_ct1.get_pointer(), device31_acc_ct1.get_pointer() + ARRAY_SIZE);
        });
    });
  }
}

Not all values of parameters could be evaluated in migration.
The combination of matrix data type and scalar type is unsupported.

Migrated code, which is using unsupported image format:
dpct::image_wrapper<cl::sycl::uint2, 2> tex; //2-channel image is not supported

void test_image(dpct::image_accessor_ext<cl::sycl::uint2, 2> acc){ 
              cl::sycl::uint2 tex_data;
              tex_data = acc.read(0, 0); 
}
int main(){
…
dpct::get_default_queue().submit([&](cl::sycl::handler &cgh) {
               …
              auto acc = tex.get_access(cgh);
              auto smpl = tex.get_sampler();
              …
              cgh.single_task<class dpct_single_kernel>([=]{
                              test_image(dpct::image_accessor_ext<cl::sycl::uint2, 2>(smpl, acc));}
                               );
              });
             …
}

Manually fixed code:
dpct::image_wrapper<cl::sycl::uint4, 2> tex;

void test_image(dpct::image_accessor_ext<cl::sycl::uint4, 2> acc){
              cl::sycl::uint4 tex_data;
              tex_data = acc.read(0, 0); 
}
int main(){
…
dpct::get_default_queue().submit([&](cl::sycl::handler &cgh) {
               …
              auto acc = tex.get_access(cgh);
              auto smpl = tex.get_sampler();
              …
              cgh.single_task<class dpct_single_kernel>([=]{
                              test_image(dpct::image_accessor_ext<cl::sycl::uint4, 2>(smpl, acc));}
                               );
              });
             …
}

dpct::constant_memory<int, 4> array(dimX, dimY, dimZ, dimW);
void kernel(sycl::id<1> idx, dpct::accessor<int, dpct::constant, 4> const_array) {
…
…  = const_array[x][y][z][w];
…
}

Manually fixed code:
dpct::constant_memory<int, 3> array(dimX, dimY, dimZ * dimW);
void kernel(sycl::id<1> idx, dpct::accessor<int, dpct::constant, 3> const_array) {
…
…  = const_array[x][y][w * dimZ + z];
…
}

Original code:
__device__ void foo(const texture<char, 2, cudaReadModeNormalizedFloat> tex);
__device__ void foo (const texture<char, 2, cudaReadModeElementType> tex);

Resulting code:
void foo (dpct::image_accessor_ext<char, 2> tex);
void foo (dpct::image_accessor_ext<char, 2> tex);

#define POW(B, E) pow(B, E)
POW(2.5, 3.1);  //should be migrated to sycl::pow<double>(2.5, 3.1)
POW(2.5, 3); //should be migrated to sycl::pown((float)2.5, 3)