移行されたコードのデバッグ: ランタイムの動作#

注

CodePin 機能は実験的なものであり、将来のリリースで完成する予定です。

状況によって、移行された SYCL* プログラムの実行時の動作が元の CUDA* プログラムと異なることがあります。これには、以下のような原因が考えられます。

ハードウェア間の演算精度の違い
CUDA* API と SYCL* API のセマンティクスの違い
自動移行中に発生したエラー

CodePin は、実行時の動作のこのような不一致のデバッグ労力を軽減するインテル® DPC++ 互換性ツールの機能です。CodePin が有効である場合、インテル® DPC++ 互換性ツールは CUDA* プログラムを SYCL* に移行しますが、CUDA* プログラムをインストルメントしたバージョンも生成します。

このインストルメント化されたコードは、選択された API またはカーネル呼び出しの前後の関連する変数のデータをレポートにダンプします。CUDA* プログラムと SYCL* プログラムから生成されたレポートを比較して、実行時の動作の相違となる原因を特定します。

CodePin を有効にする#

CodePin を有効にするには、–enable-codepin オプションを使用します。インストルメントされたプログラムは、dpct_output_codepin_cuda および dpct_output_codepin_sycl フォルダーに配置されます。

例#

次の CUDA* コードの例には、vectorAdd カーネル呼び出しの前の cudaMemcpy() に問題があります。コピーされるサイズが vectorSize * sizeof(int3) ではなく vectorSize * 12 としてハードコードされているため、移行された SYCL* プログラムは正しく機能しません。これは、int3 は sycl::int3 に移行されますが、sycl::int3 のサイズが 12 バイトではなく 16 バイトであるためです。

//example.cu 
#include <iostream> 
__global__ void vectorAdd(int3 *a, int3 *result) { 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    result[tid].x = a[tid].x + 1; 
    result[tid].y = a[tid].y + 1; 
    result[tid].z = a[tid].z + 1; 
} 

int main() { 
    const int vectorSize = 4; 
    int3 h_a[vectorSize], h_result[vectorSize]; 
    int3 *d_a, *d_result; 
    for (int i = 0; i < vectorSize; ++i) 
        h_a[i] = make_int3(1, 2, 3); 

    cudaMalloc((void **)&d_a, vectorSize * sizeof(int3)); 
    cudaMalloc((void **)&d_result, vectorSize * sizeof(int3)); 

    // Copy host vectors to device 
    // !! Using 12 instead of "sizeof(int3)」 
    cudaMemcpy(d_a, h_a, vectorSize * 12, cudaMemcpyHostToDevice); 

    // Launch the CUDA kernel 
    vectorAdd<<<1, 4>>>(d_a, d_result); 

    // Copy result from device to host 
    cudaMemcpy(h_result, d_result, vectorSize * sizeof(int3), 
        cudaMemcpyDeviceToHost); 

    // Print the result 
    for (int i = 0; i < vectorSize; ++i) { 
        std::cout << "Result[" << i << "]: (「
             << h_result[i].x << ", " << h_result[i].y << ", " << h_result[i].z << ")\n"; 
    }
 }

 
/* 
Execution Result: 
    Result[0]: (2, 3, 4) 
    Result[1]: (2, 3, 4) 
    Result[2]: (2, 3, 4) 
    Result[3]: (2, 3, 4) 
*/

この問題をデバッグするには、CodePin を有効にして CUDA* プログラムを移行します。

dpct example.cu --enable-codepin

移行後、次の 2 つのファイルが作成されます。dpct_output_codepin_sycl/example.dp.cpp と dpct_output_codepin_cuda/example.cu。

workspace
 ├── example.cu
 ├── dpct_output_codepin_sycl
 │     ├── example.dp.cpp
 │     ├── generated_schema.hpp
 │     └── MainSourceFiles.yaml
 ├── dpct_output_codepin_cuda
 │     ├── example.cu
 │     └── generated_schema.hpp

dpct_output_codepin_sycl/example.dp.cpp は移行され、インストルメントされた SYCL* プログラムです。

//dpct_output_codepin_sycl/example.dp.cpp 
#include <dpct/dpct.hpp> 
#include <sycl/sycl.hpp>

#include "generated_schema.hpp" 
#include <dpct/codepin/codepin.hpp> 
#include <iostream> 

void vectorAdd(sycl::int3 *a, sycl::int3 *result, 
            const sycl::nd_item<3> &item_ct1) { 
    int tid = item_ct1.get_group(2) * item_ct1.get_local_range(2) + 
                item_ct1.get_local_id(2); 
    result[tid].x() = a[tid].x() + 1; 
    result[tid].y() = a[tid].y() + 1; 
    result[tid].z() = a[tid].z() + 1; 
} 

int main() { 
    sycl::device dev_ct1; 
    sycl::queue q_ct1(dev_ct1, 
                        sycl::property_list{sycl::property::queue::in_order()}); 
    const int vectorSize = 4; 
    sycl::int3 h_a[vectorSize], h_result[vectorSize]; 
    sycl::int3 *d_a, *d_result; 
    for (int i = 0; i < vectorSize; ++i) 
        h_a[i] = sycl::int3(1, 2, 3); 

    d_a = sycl::malloc_device<sycl::int3>(vectorSize, q_ct1); 
    dpct::experimental::get_ptr_size_map()[*((void **)&d_a)] = 
        vectorSize * sizeof(sycl::int3); 

    d_result = sycl::malloc_device<sycl::int3>(vectorSize, q_ct1); dpct::experimental::get_ptr_size_map()[*((void **)&d_result)] = 
        vectorSize * sizeof(sycl::int3); 

    // Copy host vectors to device 
    q_ct1.memcpy(d_a, h_a, vectorSize * 12); 

    // Launch the CUDA kernel 
    dpct::experimental::gen_prolog_API_CP( 
        "example.cu:38:3(SYCL)", &q_ct1, VAR_SCHEMA_0, (long *)&d_a, VAR_SCHEMA_1, (long *)&d_result); 
    q_ct1.parallel_for( sycl::nd_range<3>(sycl::range<3>(1, 1, 4), 
        sycl::range<3>(1, 1, 4)), 
        [=](sycl::nd_item<3> item_ct1) { vectorAdd(d_a, d_result, item_ct1); }); 

    // Copy result from device to host 
    dpct::experimental::gen_epilog_API_CP( 
        "example.cu:38:3(SYCL)", &q_ct1, 
        VAR_SCHEMA_0, (long *)&d_a, VAR_SCHEMA_1, (long *)&d_result); 

    q_ct1.memcpy(h_result, d_result, vectorSize * sizeof(sycl::int3)).wait(); 

    // Print the result 
    for (int i = 0; i < vectorSize; ++i) { 
        std::cout << "Result[" << i << "]: (" << h_result[i].x() << ", " << 
                h_result[i].y() << ", " << h_result[i].z() << ")\n"; 
    } 
} 

/* Execution Result: 
    Result[0]: (2, 3, 4) 
    Result[1]: (2, 3, 4) 
    Result[2]: (2, 3, 4) 
    Result[3]: (1, 1, 1) <--- incorrect result 
*/

dpct_output_codepin_cuda/example.cu は移行され、インストルメントされた CUDA* プログラムです。

//dpct_output_codepin_cuda/example.cu 
#include "generated_schema.hpp" 
#include <dpct/codepin/codepin.hpp> 
#include <iostream>

__global__ void vectorAdd(int3 *a, int3 *result) {
   int tid = blockIdx.x * blockDim.x + threadIdx.x;
    result[tid].x = a[tid].x + 1;
    result[tid].y = a[tid].y + 1;
    result[tid].z = a[tid].z + 1; }
 
int main() { 
    const int vectorSize = 4;
     int3 h_a[vectorSize], h_result[vectorSize]; 
    int3 *d_a, *d_result; 
    for (int i = 0; i < vectorSize; ++i) 
        h_a[i] = make_int3(1, 2, 3); 

    cudaMalloc((void **)&d_a, vectorSize * sizeof(int3)); 
    dpct::experimental::get_ptr_size_map()[*((void **)&d_a)] = 
        vectorSize * sizeof(int3); 
    cudaMalloc((void **)&d_result, vectorSize * sizeof(int3)); 
    dpct::experimental::get_ptr_size_map()[*((void **)&d_result)] = 
        vectorSize * sizeof(int3); 

// Copy host vectors to device 
cudaMemcpy(d_a, h_a, vectorSize * 12, cudaMemcpyHostToDevice); 

    // Launch the CUDA kernel 
    dpct::experimental::gen_prolog_API_CP( 
        "example.cu:38:3", 0, VAR_SCHEMA_0, 
        (long *)&d_a, VAR_SCHEMA_1, (long *)&d_result); 
    vectorAdd<<<1, 4>>>(d_a, d_result); 

    // Copy result from device to host 
    dpct::experimental::gen_epilog_API_CP( 
        "example.cu:38:3", 0, VAR_SCHEMA_0, 
        (long *)&d_a, VAR_SCHEMA_1, (long *)&d_result); 
    cudaMemcpy(h_result, d_result, vectorSize * sizeof(int3), 
            cudaMemcpyDeviceToHost); 

    // Print the result 
    for (int i = 0; i < vectorSize; ++i) { 
        std::cout << "Result[" << i << "]: (「
                 << h_result[i].x << ", " << h_result[i].y << ", " << h_result[i].z << ")\n"; 
    }
 }

 /* 
Execution Result: 
    Result[0]: (2, 3, 4) 
    Result[1]: (2, 3, 4) 
    Result[2]: (2, 3, 4) 
    Result[3]: (2, 3, 4) 
*/

dpct_output_codepin_sycl/example.dp.cpp と dpct_output_codepin_cuda/example.cu をビルドし、ビルドされたバイナリーを実行すると、次の実行ログファイルが生成されます。

インストルメントされた CUDA* プログラムのレポート	インストルメントされた移行後の SYCL* プログラムのレポート
1[ 2 { 3 "ID": "example.cu:26:3:prolog", 4 "Free Device Memory": "16374562816", 5 "Total Device Memory": "16882663424", 6 "Elapse Time(ms)": "0", 7 "CheckPoint": { 8 "d_a": { 9 "Type": "Pointer", 10 "Data": [ 11 { 12 "Type": "int3", 13 "Data": [ 14 { 15 "x": { 16 "Type": "int", 17 "Data": [ 18 1 19 ] 20 } 21 }, 22 { 23 "y": { 24 "Type": "int", 25 "Data": [ 26 2 27 ] 28 } 29 }, 30 ...	1[ 2 { 3 "ID": "example.cu:26:3:prolog", 4 "Free Device Memory": "0", 5 "Total Device Memory": "31023112192", 6 "Elapse Time(ms)": "0", 7 "CheckPoint": { 8 "d_a": { 9 "Type": "Pointer", 10 "Data": [ 11 { 12 "Type": "sycl::int3", 13 "Data": [ 14 { 15 "x": { 16 "Type": "int", 17 "Data": [ 18 1 19 ] 20 } 21 }, 22 { 23 "y": { 24 "Type": "int", 25 "Data": [ 26 2 27 ] 28 } 29 }, 30 ...

インストルメントされた CUDA* プログラムのレポート

インストルメントされた移行後の SYCL* プログラムのレポート

[ 
    { 
       "ID": "example.cu:26:3:prolog", 
       "Free Device Memory": "16374562816", 
       "Total Device Memory": "16882663424", 
       "Elapse Time(ms)": "0", 
       "CheckPoint": { 
          "d_a": { 
             "Type": "Pointer", 
           "Data": [ 
               { 
                   "Type": "int3", 
                   "Data": [ 
                       { 
                          "x": { 
                              "Type": "int", 
                              "Data": [ 
                                  1 
                              ] 
                          } 
                       }, 
                       { 
                              "y": { 
                                   "Type": "int", 
                                   "Data": [ 
                                   2 
                                   ] 
                              } 
                       }, 
 ...

[ 
    { 
    "ID": "example.cu:26:3:prolog", 
    "Free Device Memory": "0", 
    "Total Device Memory": "31023112192", 
    "Elapse Time(ms)": "0", 
    "CheckPoint": { 
       "d_a": { 
          "Type": "Pointer", 
        "Data": [ 
           { 
               "Type": "sycl::int3", 
               "Data": [ 
                  { 
                     "x": { 
                        "Type": "int", 
                        "Data": [ 
                           1 
                        ] 
                     } 
                  }, 
                  { 
                      "y": { 
                         "Type": "int", 
                         "Data": [ 
                         2 
                      ] 
                  } 
              }, 
 ...

このレポートは、CUDA* プログラムと SYCL* プログラムの実行時の動作がどこで相違するか特定するのに役立ちます。

CodePin の結果を解析する#

codepin-report.py (dpct/c2s –codepin-report でもトリガー可能) は、CUDA* と SYCL* コードの両方から実行ログファイルを入力し、自動解析を行う互換性ツールの機能です。codepin-report.py は、データ値の不一致を識別し、実行の統計データを報告できます。

ユーザーは、次の形式で浮動小数点の比較許容値を指定できます。

{ 
     "bf16_abs_tol": 9.77E-04,, 
     "fp16_abs_tol":      9.77E-04,, 
     "float_abs_tol": 1.19E-04,, 
     "double_abs_tol": 2.22E-04,, 
     "rel_tol": 1e-3 
}

最初の 4 つの項目 “bf16_abs_tol”、“bf16_abs_tol”、“bf16_abs_tol” および “bf16_abs_tol” は、対応するタイプの絶対許容範囲です。最後の “rel_tol” は、比率値で表される相対許容範囲です。

codepin-report.py は、次のコマンドラインで CUDA* および SYCL* コードの両方から生成された実行ログファイルを使用します。 codepin-report.py [-h] --instrumented-cuda-log <file path> --instrumented-sycl-log <file path> [--floating-point-comparison-epsilon <file path>]

以下は解析レポートの例です。

CodePin Summary 
Totally APIs count, 2 
Consistently APIs count, 2 
Most Time-consuming Kernel(CUDA), example.cu:26:3:epilog, time:8.2316 
Most Time-consuming Kernel(SYCL), example.cu:26:3:epilog, time:10.2575 
Peak Device Memory Used(CUDA), 508100608 
Peak Device Memory Used(SYCL), 31023112192 
CUDA Meta Data ID, SYCL Meta Data ID, Type, Detail 
example.cu:26:3:prolog,example.cu:26:3:prolog,Data value,[WARNING: METADATA MISMATCH] The pair of prolog data example.cu:26:3:prolog are mismatched, 
and the corresponding pair of epilog data matches. This mismatch may be caused by the initialized memory or argument used in the API example.cu.

インテル® DPC++
互換性ツール・
デベロッパー・ガイド
およびリファレンス

移行されたコードのデバッグ: ランタイムの動作

説明

移行されたコードのデバッグ: ランタイムの動作#

CodePin を有効にする#

例#

CodePin の結果を解析する#

インテル® DPC++互換性ツール・デベロッパー・ガイドおよびリファレンス

移行されたコードのデバッグ: ランタイムの動作

説明

移行されたコードのデバッグ: ランタイムの動作#

CodePin を有効にする#

例#

CodePin の結果を解析する#

インテル® DPC++
互換性ツール・
デベロッパー・ガイド
およびリファレンス