DPCT1012#

メッセージ#

カーネル実行時間の測定パターンを検出し、SYCL* で時間測定を行うための初期コードを生成しました。要件に応じて、時間の測定方法を変更できます。

詳細な説明#

生成されたコードは、CPU時間を使用してカーネルの実行時間を測定します。要件に応じて、時間の測定方法を変更できます。

修正方法の提案#

ロジックを確認して、必要に応じて調整してください。

例えば、以下のオリジナル CUDA* コードについて考えてみます。

1  __global__ void kernel() { 
2     ...
3  } 
4 
5  void foo() { 
6   cudaEvent_t start; 
7   cudaEvent_t end; 
8   cudaEventCreate(&start); 
9   cudaEventCreate(&end); 
10  cudaEventRecord(start); 
11  kernel<<<1, 1>>>(); 
12  cudaEventRecord(end, 0); 
13  cudaEventSynchronize(end); 
14  float time; 
15  cudaEventElapsedTime(&time, start, end); 
16 }

このコードは、以下の SYCL* コードに移行されます。

1  void kernel() { 
2      ... 
3  } 
4 
5  void foo() { 
6   dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
7   sycl::queue &q_ct1 = dev_ct1.default_queue(); 
8   dpct::event_ptr start; 
9   std::chrono::time_point<std::chrono::steady_clock> start_ct1; 
10  dpct::event_ptr end; 
11  std::chrono::time_point<std::chrono::steady_clock> end_ct1; 
12  start = new sycl::event(); 
13  end = new sycl::event(); 
14  /* 
15  DPCT1012:0: Detected kernel execution time measurement pattern and generated 
16  an initial code for time measurements in SYCL.You can change the way time is 
17  measured depending on your goals.
18  */ 
19  start_ct1 = std::chrono::steady_clock::now(); 
20  *end = q_ct1.parallel_for( 
21  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
22     [=](sycl::nd_item<3> item_ct1) { 
23     kernel(); 
24  }); 
25  /* 
26  DPCT1012:1: Detected kernel execution time measurement pattern and generated 
27  an initial code for time measurements in SYCL.You can change the way time is 
28  measured depending on your goals.
29  */ 
30  end->wait(); 
31  end_ct1 = std::chrono::steady_clock::now(); 
32  float time; 
33  time = std::chrono::duration<float, std::milli>(end_ct1 - start_ct1).count(); 
34 }

このコードは次のように書き換えられます。

1  // User can add `--enable-profiling` option to migrate the code 
2  void kernel() { 
3   ... 
4  } 
5 
6  void foo() { 
7   dpct::device_ext &dev_ct1 = dpct::get_current_device(); 
8   sycl::queue &q_ct1 = dev_ct1.default_queue(); 
9   dpct::event_ptr start; 
10  dpct::event_ptr end; 
11  start = new sycl::event(); 
12  end = new sycl::event(); 
13  *start = q_ct1.ext_oneapi_submit_barrier(); 
14  q_ct1.parallel_for( 
15  sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)), 
16     [=](sycl::nd_item<3> item_ct1) { 
17     kernel(); 
18  }); 
19  *end = q_ct1.ext_oneapi_submit_barrier(); 
20  end->wait_and_throw(); 
21  float time; 
22  time = 
23     (end->get_profiling_info<sycl::info::event_profiling::command_end>() - 
24     start 
25     ->get_profiling_info<sycl::info::event_profiling::command_start>()) / 
26     1000000.0f; 
27 }