DPCT1012#

メッセージ#

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

説明#

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

修正方法の提案#

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

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

1__global__ void kernel() { 
2...3} 
4 
5void 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* コードに移行されます。

1void kernel() { 
2 ... 3} 
4 
5void 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 
2void kernel() { 
3 ... 4} 
5 
6void 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}