オフロード処理のデバッグ

異なるランタイムまたは計算デバイスで実行

オフロードプログラムが正常に実行されなかったか、生成された結果が正しくない場合、比較的容易な正当性の確認方法は、OpenMP* アプリケーションでは LIBOMPTARGET_PLUGIN と OMP_TARGET_OFFLOAD 環境変数を、また SYCL* アプリケーションでは ONEAPI_DEVICE_SELECTOR 環境変数を使用して、別のランタイム (OpenCL* とレベルゼロ) または計算デバイス (CPU と GPU) でアプリケーションを実行することです。異なるランタイム間で再現されるエラーは、ほとんどの場合、ランタイムの問題として排除できます。そしてデバイス間で再現されるエラーの大部分は、不良ハードウェアの問題を排除できます。

CPU 実行をデバッグ

オフロードコードを CPU で実行するには、"ホスト" 実装と OpenCL* の CPU バージョン:という 2 つのオプションがあります。"ホスト" 実装は、オフロードコードのネイティブ実装であり、オフロードされないコードと同じようにデバッグできます。OpenCL* の CPU バージョンは、OpenCL* ランタイムとコード生成プロセスを通過しますが、最終的にインテル® TBB ランタイムで実行される通常の並列コードになります。繰り返しますが、これにより慣れ親しんだアセンブリーと並列処理メカニズムのデバッグ環境が提供されています。ポインターはスタック全体をアクセスでき、データを直接参照できます。また、オペレーティング・システム・プロセスの通常の上限を超えるメモリー制限はありません。

CPU オフロード実行のエラーを検出して修正すると、GPU オフロード実行で発生するエラーよりもはるかに少ない労力でエラーを解決でき、GPU や他のアクセラレーターが接続されたシステムを利用する必要もなくなります。

OpenMP* アプリケーションで "ホスト" 実装を適用するには、“target” または “device” 構造を削除して、通常のホスト OpenMP* コードに置き換えます。LIBOMPTARGET_PLUGIN=OPENCL が設定され、GPU オフロードが無効化されると、オフロードコードは OpenMP* ランタイムで実行され、TBB が並列処理を行います。

SYCL* アプリケーションで ONEAPI_DEVICE_SELECTOR=host を設定すると "ホスト" デバイスはシングルスレッドで実行を行います。これは、データ競合やデッドロックなどスレッド化の問題が実行エラーの原因であるか判断するのに役立ちます。ONEAPI_DEVICE_SELECTOR=opencl:cpu に設定すると、CPU の OpenCL* ランタイムが使用され、TBB が並列処理を行います。

インテル® ディストリビューションの GDB* を使用して互換性のある GPU での実行をデバッグ

インテル® ディストリビューションの GDB* については、インテル® ディストリビューションの GDB* の導入ガイドに詳しく記載されています (Linux* ホスト (英語) | Windows* ホスト (英語)。有用なコマンドについては、インテル® ディストリビューションの GDB* の「リファレンス・シート」 (英語) で簡単に説明されています。ただし、GDB* を使用して GPU アプリケーションをデバッグする方法は、ホストでの手順とは若干異なるため (一部のコマンドの使い方が異なり、見慣れない出力が表示されることがあります)、それらの違いの一部をここで紹介します。

Linux* ホストでインテル® ディストリビューションの GDB* を使用したデバッグのチュートリアル (英語) では、SYCL* プログラムのデバッグセッションを開始し、カーネル内にブレークポイントを設定し、プログラムを実行して GPU にオフロードしローカル値を出力し、スレッドの SIMD レーン 5 を切り替えて変数を再度出力するサンプルのデバッグセッションを紹介しています。

通常の GDB* と同様に、command <CMD> には GDB の help <CMD>> コマンドを使用して、<CMD> の情報テキストを読み取ります。以下に例を示します:

(gdb) help info threads 
Display currently known threads. Usage: info threads [OPTION]... [ID]... If ID is given, it is a space-separated list of IDs of threads to display. Otherwise, all threads are displayed. 
Options: 
  -gid 
    Show global thread IDs.

GDB* でインフェリア―、スレッド、および SIMD レーンの参照

アプリケーションのスレッドは、デバッガーによって一覧表示できます。表示される情報には、スレッド ID とスレッドが停止している位置が含まれます。GPU スレッドの場合、デバッガーはアクティブな SIMD レーンも表示します。

上記の例では、GDB* の "info threads" コマンドでスレッドを表示していますが、見慣れない形式で情報が示されることがあります:

 Id Target Id Frame 
 1.1 Thread <id omitted> <frame omitted> 
 1.2 Thread <id omitted> <frame omitted> 
* 2.1:1 Thread 1073741824 <frame> at array-transform.cpp:61 
 2.1:[3 5 7] Thread 1073741824 <frame> at array-transform.cpp:61 
 2.2:[1 3 5 7] Thread 1073741888 <frame> at array-transform.cpp:61 
 2.3:[1 3 5 7] Thread 1073742080 <frame> at array-transform.cpp:61

GDB* は次の形式でスレッドを表示します: <inferior_number>.<thread_number>:<SIMD Lane/s>

例えば、スレッド ID: "2.3:[1 3 5 7]" は、インフェリアー 2 で実行されるスレッド 3 の SIMD レーン1、3、5 および 7 を意味します。

GDB* 用語の "inferior (インフェリアー)" は、でデバッグされるプロセスを指します。GPU にオフロードを行うプログラムのデバッグセッションには、通常 2 つのインフェリアーがあります。プログラムのホストを示す 1 つの "ネイティブ" インフェリアー (上記のインフェリアー 1) と、GPU デバイスを示すもう 1つの "リモート" インフェリアー (上記のインフェリアー 2) です。インテル® ディストリビューションの GDB* は自動的に GPU インフェリアーを生成するため、特に操作は必要ありません。

式の値を出力すると、式は現在のスレッドの現在の SIMD レーンのコンテキストで評価されます。"thread 3:4 "、"thread :6 "、または "thread 7 " などの "thread" コマンドを使用して、スレッドと SIMD レーンを切り替えることができます。最初のコマンドは、スレッド3 と SIMD レーン 4 に切り替えます。2 番目のコマンドは、現在のスレッドで SIMD レーン 6 に切り替えます。3 番目のコマンドは、スレッド 7 に切り替えます。選択されるデフォルトレーンは、以前に選択したレーン (アクティブであれば)、またはスレッド内で最初にアクティブになったレーンのどちらかになります。

"thread apply コマンド" は、同様に広域または集中的である可能性があります (これにより、変数を調査するコマンドからの出力を制限しやすくなります)。SIMD レーンのデバッグの詳細と例については、Linux* ホストでインテル® ディストリビューションの GDB* を使用したデバッグのチュートリアル (英語) を参照してください。

GDB* のスレッドと下位クラスの詳細については、https://sourceware.org/gdb/current/onlinedocs/gdb/Threads.html (英語) および https://sourceware.org/gdb/current/onlinedocs/gdb/Inferiors-Connections-and-Programs.html#Inferiors-Connections-and-Programs (英語) を参照してください。

スケジューラーの制御

デフォルトでは、スレッドがブレークポイントに到達すると、デバッガーはブレーク・ポイント・ヒット・イベントをユーザーに通知する前にすべてのスレッドを停止します。これは GDB* のすべて停止モードです。非停止モードでは、他のスレッドが実行される間、スレッドの停止イベントが表示されます。

すべて停止モードでは、スレッドが再開されると (例: continue コマンドで通常のように再開する、または next コマンドでステップ実行する場合)、他のすべてのスレッドも再開されます。スレッド化されたアプリケーションで複数のブレークポイントが設定されていると、ブレークポイントに到達した次のスレッドが続くスレッドではない可能性があるため、混乱を招く可能性があります。

set scheduler-locking コマンドを使用することで、現在のスレッドが再開された時に他のスレッドが再開されないように制御することができます。これは、現在のスレッドのみが命令を実行している時に、他のスレッドの介入を避けるのに有効です。help set scheduler-locking と入力すると、利用可能なオプションが表示されます。詳細は、https://sourceware.org/gdb/current/onlinedocs/gdb/Thread-Stops.html (英語) をご覧ください。SIMD レーンは個別に再開できないことに注意してください。これは、ベースとなるスレッドと共に再開されます。

非停止モードのデフォルトでは、現在のスレッドのみが再開されます。すべてのスレッドを再開するには、continue コマンドで "-a" フラグを指定します。

1 つ以上のスレッド/レーンの情報をダンプ (Thread Apply)

プログラム状態を調査するコマンドは、通常、現在のスレッドの現在の SIMD レーンのコンテキストに適用されます。複数のコンテキストの値を調査することが必要なこともあります。そのような場合、thread apply コマンドを使用します。例えば、以下はスレッド 2,5 の SIMD レーン 3-5 に対して print element コマンドを実行します:

(gdb) thread apply 2.5:3-5 print element

同様に、以下は、現在のスレッドの SIMD レーン 3、5、および 6 のコンテキストに対し print element コマンドを実行します:

(gdb) thread apply :3 :5 :6 print element

ブレークポイント停止後の GPU コードのステップ実行

GPU にオフロードされたカーネル内で停止するには、カーネル内のソース行にブレークポイントを設定するだけです。GPU スレッドがそのソース行に到達すると、デバッガーは実行を停止してブレークポイントの到達を示します。ソース行単位でスレッドをステップ実行するには、step または next コマンドを使用します。step コマンドは関数にステップインし、next コマンドは関数呼び出しをステップオーバーします。ステップ実行する前に、他のスレッドの介入を避けるため set scheduler-locking step を設定することを推奨します。

インテル® ディストリビューションの GDB* で使用する SYCL* 実行形式をビルド

ホスト・アプリケーションのデバッグと同様に、GPU でデバッグ可能なバイナリーを作成するには、いくつかの追加フラグを指定する必要があります。詳細については、Linux* ホストでのインテル® ディストリビューションの GDB* 導入ガイド (英語) をご覧ください。

ジャストインタイム (JIT) コンパイルを行う際にスムーズなデバッグを可能にするには、-g フラグを指定してコンパイラーのデバッグ情報生成を有効にし、アプリケーションのホストと JIT コンパイルカーネルの両方で -O0 フラグを指定して最適化を無効にします。カーネルのフラグはリンク時に適用されます。以下に例を示します:

• プログラムをコンパイルするには、次のコマンドを使用します: icpx -fsycl -g -O0 -c myprogram.cpp

• プログラムをリンクするには、次のコマンドを使用します: icpx -fsycl -g -O0 myprogram.o

CMake を使用してプログラムをビルドする場合、CMAKE_BUILD_TYPE に Debug タイプを使用し、CMAKE_CXX_FLAGS_DEBUG 変数に -O0 を追加します。例: set (CMAKE_CXX_FLAGS_DEBUG "${CMAKE_CXX_FLAGS_DEBUG} -O0")

デバッグ向けにビルドされたアプリケーションは、通常の “リリース” ビルドで最適化されたアプリケーションよりも起動に時間がかかる場合があります。そのため、デバッガーで起動すると、プログラムの実行速度が遅くなったように感じられることがあります。これにより問題が生じる場合、大規模なアプリケーションの開発者は、プログラムの実行時ではなくビルド時に JIT オフロードコードを事前 (AOT) コンパイルすることを推奨します (これには、-g -O0 を使用するとビルドに時間がかかる場合があります)。詳細については、コンパイル手順の概要をご覧ください。

GPU 向けの事前コンパイルを行う場合、ターゲットデバイスに対応したデバイスタイプを指定する必要があります。次のコマンドを使用して、現在のマシンで利用可能な GPU デバイスオプションを確認できます: ocloc compile --help

さらに、カーネルのデバッグモードを有効にします。次の AOT コンパイルの例は、KBL デバイスをターゲットにしています:

dpcpp -g -O0 -fsycl-targets=spir64_gen-unknown-unknown-sycldevice \ 
-Xs "-device kbl -internal_options -cl-kernel-debug-enable -options -cl-opt-disable" myprogram.cpp

インテル® ディストリビューションの GDB* で使用する OpenMP* 実行形式のビルド

プログラムのコンパイルとリンクに -g -O0 フラグを使用します。以下に例を示します:

icpx -fiopenmp -O0 -fopenmp-targets=spir64 -c -g myprogram.cpp
icpx -fiopenmp -O0 -fopenmp-targets=spir64 -g myprogram.o

次の環境変数を設定して最適化を無効にし、カーネルのデバッグ情報を有効にします:

export LIBOMPTARGET_OPENCL_COMPILATION_OPTIONS="-g -cl-opt-disable" 
export LIBOMPTARGET_LEVEL0_COMPILATION_OPTIONS="-g -cl-opt-disable"

GPU 実行をデバッグ

オフロードプログラムでよくある問題は、プログラムが実行されず、追加情報をほとんど持たない OpenCL* エラーが生成されることです。OpenCL* アプリケーションのインターセプト・レイヤーと onetrace、ze_tracer、および cl_tracer を使用して、このエラーに関する詳細情報を取得できます。これは、開発者が問題の原因を特定するのに役立ちます。

OpenCL* アプリケーションのインターセプト・レイヤー

このライブラリーを使用する場合、Buildlogging、ErrorLogging、および USMChecking=1 オプションを使用してエラーの原因を特定できます。

1. 次のテキストを含む clintercept.conf ファイルをホーム・ディレクトリーに作成します:

   SimpleDumpProgramSource=1 
   CallLogging=1 
   LogToFile=1 
   //KernelNameHashTracking=1 
   BuildLogging=1 
   ErrorLogging=1 
   USMChecking=1 
   //ContextCallbackLogging=1 
   // Profiling knobs 
   KernelInfoLogging=1 
   DevicePerformanceTiming=1 
   DevicePerformanceTimeLWSTracking=1 
   DevicePerformanceTimeGWSTracking=1

2. 次のように cliloader を使用してアプリケーションを実行します:

   <OCL_Intercept_Install_Dir>/bin/cliloader/cliloader -d ./<app_name> <app_args>

3. ~CLIntercept_Dump/<app_name> ディレクトリーで次の結果を確認します:

   • clintercept_report.txt: プロファイルの結果

   • clintercept_log.txt: OpenCL* の問題をデバッグする際に使用される OpenCL* 呼び出しログ

次のテキストは、ランタイムエラーが発生したプログラムで生成されたログファイルの例の一部です: CL_INVALID_ARG_VALUE (-50)

...<<<< clSetKernelArgMemPointerINTEL -> CL_SUCCESS 
>>>> clGetKernelInfo( _ZTSZZ10outer_coreiP5mesh_i16dpct_type_1c0e3516dpct_type_60257cS2_S2_S2_S2_S2_S2_S2_S2_fS2_S2_S2_S2_iENKUlRN2cl4sycl7handlerEE197->45clES6_EUlNS4_7nd_itemILi3EEEE225->13 ): param_name = CL_KERNEL_CONTEXT (1193) 
<<<< clGetKernelInfo -> CL_SUCCESS 
>>>> clSetKernelArgMemPointerINTEL( _ZTSZZ10outer_coreiP5mesh_i16dpct_type_1c0e3516dpct_type_60257cS2_S2_S2_S2_S2_S2_S2_S2_fS2_S2_S2_S2_iENKUlRN2cl4sycl7handlerEE197->45clES6_EUlNS4_7nd_itemILi3EEEE225->13 ): kernel = 0xa2d51a0, index = 3, value = 0x41995e0 
mem pointer 0x41995e0 is an UNKNOWN pointer and no device support shared system pointers! ERROR! clSetKernelArgMemPointerINTEL returned CL_INVALID_ARG_VALUE (-50) 
<<<< clSetKernelArgMemPointerINTEL -> CL_INVALID_ARG_VALUE

この例は、次の値がエラーのデバッグに役立ちます:

• ZTSZZ10outer_coreiP5mesh

• index = 3, value = 0x41995e0

このデータによりどのカーネルに問題があるか、またどの引数に問題があるかが分かり、その理由を特定できます。

onetrace、ze_tracer、および cl_tracer

OpenCL* アプリケーションのインターセプト・レイヤーと同様に、onetrace、ze_tracer および cl_tracer ツールはレベルゼロのランタイムエラーの原因を検出するのに役立ちます。

onetrace または ze_tracer ツールを使用してレベルゼロにおける問題の根本的な原因を特定します (cl_tracer は、OpenCL* の問題の同様の原因を特定するのに使用されます)。

1. 呼び出しログモードでアプリケーションを実行します。ツールの出力をファイルにリダイレクトすることを推奨します。

   ./onetrace -c ./<app_name> <app_args> [2> log.txt]

“ze_tracer” のコマンドも同様です。“onetrace” を “ze_tracer” に置き換えるだけです。

1. 呼び出しトレースを確認してエラーを特定します (log.txt)。以下に例を示します:

   >>>> [102032049] zeKernelCreate: hModule = 0x55a68c762690 desc = 0x7fff865b5570 {29 0 0 GEMM} phKernel = 0x7fff865b5438 (hKernel = 0) 
   <<<< [102060428] zeKernelCreate [28379 ns] hKernel = 0x55a68c790280 -> ZE_RESULT_SUCCESS (0) 
   ...>>>> [102249951] zeKernelSetGroupSize: hKernel = 0x55a68c790280 groupSizeX = 256 groupSizeY = 1 groupSizeZ = 1 
   <<<< [102264632] zeKernelSetGroupSize [14681 ns] -> ZE_RESULT_SUCCESS (0) 
   >>>> [102278558] zeKernelSetArgumentValue: hKernel = 0x55a68c790280 argIndex = 0 argSize = 8 pArgValue = 0x7fff865b5440 
   <<<< [102294960] zeKernelSetArgumentValue [16402 ns] -> ZE_RESULT_SUCCESS (0) 
   >>>> [102308273] zeKernelSetArgumentValue: hKernel = 0x55a68c790280 argIndex = 1 argSize = 8 pArgValue = 0x7fff865b5458 
   <<<< [102321981] zeKernelSetArgumentValue [13708 ns] -> ZE_RESULT_ERROR_INVALID_ARGUMENT (2013265924) 
   >>>> [104428764] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 2 argSize = 8 pArgValue = 0x7ffe289c7e60 
   <<<< [104442529] zeKernelSetArgumentValue [13765 ns] -> ZE_RESULT_SUCCESS (0) 
   >>>> [104455176] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 3 argSize = 4 pArgValue = 0x7ffe289c7e2c 
   <<<< [104468472] zeKernelSetArgumentValue [13296 ns] -> ZE_RESULT_SUCCESS (0) 
   ...

   この例のログには次のデータが示されています:

   • 問題の原因となるレベルゼロ API 呼び出し (zeKernelSetArgumentValue)

   • 問題の理由 (ZE_RESULT_ERROR_INVALID_ARGUMENT)

   • この引数インデックス (argIndex = 1)

   • 無効な値の場所 (pArgValue = 0x7fff865b5458)

   • カーネルハンドル (hKernel = 0x55a68c790280)、この問題が検出されたカーネル名を示します (GEMM)

"ファイルへのリダイレクト" オプションを省略して、すべての出力 (アプリケーションの出力 + ツールの出力) を 1 つのストリームにダンプすることで、より多くの情報を取得できます。単一のストリームにダンプを行うことで、アプリケーションの出力に関連するエラーの原因を特定するのに役立つことがあります (例えば、アプリケーションの初期化と計算の最初のフェーズでエラーが発生しているなどが分かります)。

Level Zero Matrix Multiplication (matrix size: 1024 x 1024, repeats 4 times) 
Target device: Intel® Graphics [0x3ea5]
 ...>>>> [104131109] zeKernelCreate: hModule = 0x55af5f39ca10 desc = 0x7ffe289c7f80 {29 0 0 GEMM} phKernel = 0x7ffe289c7e48 (hKernel = 0) 
<<<< [104158819] zeKernelCreate [27710 ns] hKernel = 0x55af5f3ca600 -> ZE_RESULT_SUCCESS (0)
 ...>>>> [104345820] zeKernelSetGroupSize: hKernel = 0x55af5f3ca600 groupSizeX = 256 groupSizeY = 1 groupSizeZ = 1 
<<<< [104360082] zeKernelSetGroupSize [14262 ns] -> ZE_RESULT_SUCCESS (0) 
>>>> [104373679] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 0 argSize = 8 pArgValue = 0x7ffe289c7e50 
<<<< [104389443] zeKernelSetArgumentValue [15764 ns] -> ZE_RESULT_SUCCESS (0) 
>>>> [104402448] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 1 argSize = 8 pArgValue = 0x7ffe289c7e68 
<<<< [104415871] zeKernelSetArgumentValue [13423 ns] -> ZE_RESULT_ERROR_INVALID_ARGUMENT (2013265924) 
>>>> [104428764] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 2 argSize = 8 pArgValue = 0x7ffe289c7e60 
<<<< [104442529] zeKernelSetArgumentValue [13765 ns] -> ZE_RESULT_SUCCESS (0) 
>>>> [104455176] zeKernelSetArgumentValue: hKernel = 0x55af5f3ca600 argIndex = 3 argSize = 4 pArgValue = 0x7ffe289c7e2c 
<<<< [104468472] zeKernelSetArgumentValue [13296 ns] -> ZE_RESULT_SUCCESS (0)
 ...Matrix multiplication time: 0.0427564 sec 
Results are INCORRECT with accuracy: 1
 ...Matrix multiplication time: 0.0430995 sec 
Results are INCORRECT with accuracy: 1
 ... Total execution time: 0.381558 sec

正当性

オフロードコードは、接続された計算デバイスで大量の情報を効率良く処理するカーネル、または一部の入力パラメーターから大量の情報を生成する際に利用されます。それらのかカーネルがクラッシュすることなく実行されている場合、その多くはプログラム実行のかなり後で正しい結果が得られていないことを学習したことを意味します。

そのような場合、どのカーネルが誤った結果を生成しているか特定するのは困難です。誤った結果を生成するカーネルを特定する方法として、プログラムを 2 回実行することが考えられます。最初はホストベースの実装を実行し、2 回目はオフロード実装を実行してすべてのカーネル (多くは個々のファイル) の入力と出力を取得します。次に、結果を比較して、どのカーネルが予期しない結果を生成しているか確認します (特定のイプシロンで、オフロード・ハードウェアの操作順序やネイティブの精度が異なるため、結果が最後の 1 もしくは 2 桁ホストコードとは異なる可能性があります)。

誤った結果を生成するカーネルが特定されたら、インテル® ディストリビューションの GDB* を使用して原因を調査します。基本情報と詳細なドキュメントへのリンクについては、Linux* ホストでインテル® ディストリビューションの GDB* を使用したデバッグのチュートリアル (英語) を参照してください。

SYCL* と OpenMP* はどちらもオフロードされたカーネル内で標準の print メカニズム（SYCL* と C++ OpenMP* オフロードでは printf、Fortran OpenMP* オフロードでは print *, ... など）を利用できます。これを使用して実行中の動作を確認できます。スレッドと SIMD レーンの出力をプリントし、同期メカニズムを追加して、プリントされた情報がプリント時に一貫性を持つようになることを検討してください。ストリームクラスを使用した SYCL* で同様のことを行う例は、 oneAPI GPU 最適化ガイド (PDF) に記載されています。OpenMP* オフロードには、SYCL* で説明した同様のアプローチを使用できます。

OpenMP* ディレクティブを使用してアプリケーションに並列処理を実装する方法の詳細については、「インテルのツールによる OpenMP* アプリケーションのオフロードと最適化 <https://www.intel.com/content/www/us/en/developer/tools/oneapi/training/offload-optimize-openmp-applications.html>」を参照してください。

ヒント

SYCL カーネルでは printf は冗長的になる可能性があります。簡単にするため次のマクロを追加します:

#ifdef __SYCL_DEVICE_ONLY__ 
  #define CL_CONSTANT __attribute__((opencl_constant)) 
#else 
  #define CL_CONSTANT 
#endif #define PRINTF(format, ...){ \ 
            static const CL_CONSTANT char _format[] = format; \ 
            sycl::ONEAPI::experimental::printf(_format, ## __VA_ARGS__); }

以下に使用例を示します: PRINTF("My integer variable:%d\n", (int) x);

障害

SYCL* または OpenMP* オフロード言語の誤った用法が原因でJIT コンパイルが失敗すると、プログラムはエラーで終了します。

SYCL* では 事前コンパイルされていることを判定できない場合、OpenCL* バックエンドを選択して OpenCL* アプリケーションのインターセプト・レイヤーを使用すると、構文エラーを持つカーネルを特定できることがあります。

ロジックエラーは、実行中にクラッシュが発生したり、エラーメッセージが表示されることがあります。これには以下が含まれます:

• 誤ったコンテキストに属するバッファーをカーネルに渡す場合

• "this" ポインターをクラス要素ではなくカーネルに渡す場合

• デバイスバッファーではなくホストバッファーを渡す場合

• カーネルで使用されなくても、初期化されていないポインターを渡す場合

インテル® ディストリビューションの GDB* （またはネイティブ GDB*）を使用して注意深く調査することで、生成されたすべてのコンテキストのアドレスを記録してオフロードカーネルに渡されるアドレスが正しいコンテキストに属するか確認できます。同様に、変数のアドレスがそれを含むクラスでなく、変数自身のアドレスと一致するか確認できます。

OpenCL* 割り当て用のインターセプト・レイヤーまたは、onetrace/cl_tracer を使用して、適切なバックエンドを選択する方がバッファーとアドレスをトレースするよりも簡単なことがあります。OpenCL* バックエンドを使用する場合、CallLogging、BuildLogging、ErrorLogging、および USMChecking を設定してプログラムを実行すると、コード内のどのエラーが OpenCL* エラーの原因であるかを明らかにする出力が生成されます。

onetrace や ze_tracer の呼び出しログやデバイス・タイムラインを参照すると、レベルゼロのバックエンドからのエラーの原因を理解するのに役だつ追加のエラー情報が得られます。これは、上記の論理エラーを検出するのに役立ちます。

レベルゼロ・バックエンドを使用してデバイスにオフロードする際にコードでエラーが発生する場合、OpenCL* バックエンドを試してみてください。プログラムが正常に機能する場合、レベルゼロのバックエンドにエラーをレポートしてください。デバイス向けの OpenCL* バックエンドでもエラーが再現する場合、OpenCL CPU バックエンドを試します。OpenMP* オフロードでは、OMP_TARGET_OFFLOAD を CPU に設定することで指定できます。SYCL* では、ONEAPI_DEVICE_SELECTOR=opencl:cpu を設定します。CPU　上でのデバッグはかなり容易になり、データのコピーとプログラムのデバイスへの変換によって生じる複雑性も排除できます。

問題が発生する可能性があるロジックの例として、次の SYCL* コードで parallel_for を実装する際に使用されるラムダ関数でキャプチャーされる場合を考えます。

class MyClass { 
private: 
  int *data; 
  int factor; 
  : 
void run() { 
  : 
  auto data2 = data; 
  auto factor2 = factor; 
  { 
    dpct::get_default_queue_wait().submit([&](cl::sycl::handler &cgh) 
    { 
      auto dpct_global_range = grid * block; 
      auto dpct_local_range = block; 
      cgh.parallel_for<dpct_kernel_name<class kernel_855a44>>( 
          cl::sycl::nd_range<1>( 
            cl::sycl::range<1> dpct_global_range.get(0)), 
            cl::sycl::range<1>( dpct_local_range.get(0))), 
            [=](cl::sycl::nd_item<3> item_ct1) 
      { 
            kernel(data, b, factor, LEN, item_ct1); // This blows up 
      }); 
    }); 
  } 
} // run 
} // MyClass

上記のコードでは、[=] がラムダ内で使用される変数を値でコピーするため、プログラムはクラッシュします。この例では、“factor” が実際には “this->factor” であり、“data” が実際には “this->data,” であることが明らかではないかもしれません。したがって、“this” は上記の “data” と “factor” を使用するためにキャプチャーされた変数です。OpenCL* またはレベルゼロでは、"kernel(data, b, factor, LEN, item_ct1)" 呼び出しで不正な引数エラーが原因でクラッシュします。

解決するには、ローカル変数 auto data2 と auto factor2 を使用します。“auto factor2 = factor” は “int factor2 = this->factor” になるので、ラムダ内で [=] を使って factor2 を使用すると “int” がキャプチャーされます。内部セクションを “kernel(data2, b, factor2, LEN, item_ct1);” に書き換えます。

注

この問題は、CUDA* カーネルを移行する際によく発生します。同じ CUDA* カーネルの起動シグネチャーを保持し、コマンドグループとラムダをカーネル内に配置することで問題を解決することもできます。

OpenCL* 割り当てのインターセプト・レイヤーや onetrace または ze_tracer を使用すると、カーネルが 2 つの同一アドレスで呼び出されることが分かり、拡張エラー情報を見ると、重要なデータ構造をオフロードデバイスにコピーしようとしていることを確認できます。

統合共有メモリー (USM) を使用しており、"MyClass" が USM に割り当てられる場合、上記のコードが動作することに注意してください。ただし、USM に "data" のみが割り当てられている場合、前述の理由からプログラムはクラッシュします。

この例では、カーネル呼び出しですべてを変更する必要がないように、ローカルスコープ内で同じ名前の変数を再宣言できることも留意してください。

インテル® Inspector は、このような障害の診断に役立ちます。次の環境変数を設定して、CPU デバイスでオフロードコードのメモリー解析を実行すると、インテル® Inspector は上記の問題の多くを通知します:

• OpenMP*

  • export OMP_TARGET_OFFLOAD=CPU

  • export OMP_TARGET_OFFLOAD=MANDATORY

  • export LIBOMPTARGET_PLUGIN=OPENCL

• SYCL*

  • export ONEAPI_DEVICE_SELECTOR=opencl:cpu

  • または、CPU セレクターを使用してキューを初期化し、OpenCL* CPU デバイスの使用を強制します: cl::sycl::queue Queue(cl::sycl::cpu_selector{});

• 両方

  • export CL_CONFIG_USE_VTUNE=True

  • export CL_CONFIG_USE_VECTORIZER=false

注

コンパイル中に最適化を有効にするとクラッシュする場合があります。最適化を無効にしてクラッシュが解決される場合、デバッグ向けに -g -[最適化 レベル] を指定します。詳細については、インテル® oneAPI DPC++/C++ コンパイラー・デベロッパー・ガイドおよびリファレンス (英語) を参照してください。