DPCT1008#

Message#

<function name> function is not defined in SYCL. This is a hardware-specific feature. Consult with your hardware vendor to find a replacement.

Detailed Help#

The clock function call was not replaced, because it is not defined in SYCL*. Replace this hardware-specific feature with the one offered by your hardware vendor.

Suggestions to Fix#

Consult with your hardware vendor to find a replacement.

For example, this original CUDA* code:

 1__global__ void k(clock_t *timer) {
 2  int tid = threadIdx.x;
 3  if (tid == 0)
 4    timer[0] = clock();
 5
 6  // workload
 7
 8  __syncthreads();
 9  if (tid == 0)
10    timer[1] = clock();
11}
12
13int main() {
14  clock_t *dtimer = NULL;
15  cudaMalloc((void **)&dtimer, sizeof(clock_t) * 2);
16  clock_t timer[2];
17  k<<<1, 128>>>(dtimer);
18  cudaMemcpy(timer, dtimer, sizeof(clock_t) * 2, cudaMemcpyDeviceToHost);
19  cudaFree(dtimer);
20  long double time = timer[1] - timer[0];
21  return 0;
22}

results in the following migrated SYCL code:

 1#include <sycl/sycl.hpp>
 2#include <dpct/dpct.hpp>
 3#include <time.h>
 4void k(clock_t *timer, sycl::nd_item<3> item_ct1) {
 5  int tid = item_ct1.get_local_id(2);
 6  if (tid == 0)
 7    /*
 8    DPCT1008:1: clock function is not defined in SYCL. This is a
 9    hardware-specific feature. Consult with your hardware vendor to find a
10    replacement.
11    */
12    timer[0] = clock();// clock() is used to measure the kernel runtime
13
14  // workload
15
16  item_ct1.barrier();
17  if (tid == 0)
18    /*
19             DPCT1008:2: clock function is not defined in SYCL. This is a
20             hardware-specific feature. Consult with your hardware vendor to find a
21     replacement.
22    */
23    timer[1] = clock();
24}
25
26int main() {
27  dpct::device_ext &dev_ct1 = dpct::get_current_device();
28  sycl::queue &q_ct1 = dev_ct1.default_queue();
29  clock_t *dtimer = NULL;
30  dtimer = sycl::malloc_device<clock_t>(2, q_ct1);
31  clock_t timer[2];
32  q_ct1.parallel_for(
33      sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
34      [=](sycl::nd_item<3> item_ct1) {
35        k(dtimer, item_ct1);
36      });
37  q_ct1.memcpy(timer, dtimer, sizeof(clock_t) * 2).wait();
38  sycl::free(dtimer, q_ct1);
39  long double time = timer[1] - timer[0];
40  return 0;
41}

which is rewritten to:

 1#define DPCT_PROFILING_ENABLED
 2#include <sycl/sycl.hpp>
 3#include <dpct/dpct.hpp>
 4#include <time.h>
 5void k(sycl::nd_item<3> item_ct1) {
 6  // workload
 7}
 8
 9int main() {
10  dpct::device_ext &dev_ct1 = dpct::get_current_device();
11  sycl::queue &q_ct1 = dev_ct1.default_queue();
12  dpct::event_ptr start;
13  dpct::event_ptr end;
14  start = new sycl::event();
15  end = new sycl::event();
16  *start = q_ct1.ext_oneapi_submit_barrier();
17  q_ct1.parallel_for(
18      sycl::nd_range<3>(sycl::range<3>(1, 1, 128), sycl::range<3>(1, 1, 128)),
19      [=](sycl::nd_item<3> item_ct1) {
20        k(item_ct1);
21      });
22  *end = q_ct1.ext_oneapi_submit_barrier();
23  end->wait_and_throw();
24  long double time =
25      (end->get_profiling_info<sycl::info::event_profiling::command_end>() -
26       start
27           ->get_profiling_info<sycl::info::event_profiling::command_start>()) /
28      1000000.0f;  return 0;
29}