racecheck で競合を探す
共有メモリのデータ競合を検出する
「racecheck で競合を探す」はCoddyKit上の無料CUDA Academyレッスンです。 これはレッスン3/4です。 下記で完全なレッスンを無料で読むことができます。その後、ブラウザ内の組み込みコードエディタと24時間対応のAIチューターでハンズオン演習できます。 これはCUDA Academy学習パスの一部であり、ウェブとCoddyKitアプリ全体で進捗が同期されます。 CUDA Academyコースには全4レッスンが含まれています。
データ競合とは
データ競合は、2つのスレッドが同じ共有メモリにアクセスし、そのうち少なくとも一方が書き込みを行うにもかかわらず、アクセス順序が保証されていないときに発生します。😬
競合が見つけにくい理由
競合は、テストでは正しい答えを返すのに、本番では誤った答えを返すことがあります。結果が予測できないスレッドのタイミングに左右されるためです。
racecheckを使う
compute-sanitizerのracecheckツールは共有メモリへのアクセスを分析し、通常の実行では隠れてしまう危険な競合を報告します。
compute-sanitizer --tool racecheck ./reduceバリアの不足
共有配列に書き込んだ後、バリアなしで隣のスロットを読み取ると競合が発生します。隣のスロットへの書き込みがまだ終わっていない可能性があるためです。
tile[tid] = in[i];
out[tid] = tile[tid + 1];解決策：__syncthreads
書き込みと読み取りの間に__syncthreadsを挿入し、ブロック内のすべてのスレッドが書き込みを終えてから読み取るようにします。
tile[tid] = in[i];
__syncthreads();
out[tid] = tile[tid + 1];Write-After-Writeの危険
racecheckはwrite-after-writeの危険も検出します。これは、2つのスレッドが順序付けなしに同じ共有スロットへ格納する状態です。
Read-After-Writeの危険
最も一般的な危険はread-after-writeです。あるスレッドが別のスレッドによる書き込みの途中の値を読み取るため、古いデータを読むことになります。
危険レポートを読む
racecheckは危険の種類に加え、関係する2つのスレッドと共有アドレスを示します。そのため、どのバリアが不足しているかを正確に把握できます。
共有メモリを監視する
racecheckはブロック内の__shared__メモリに注目します。グローバルメモリの競合には、通常、慎重なロジックまたはアトミック操作が必要です。
行情報を追加する
memcheckと同様に、-lineinfoを指定してコンパイルすると、racecheckは各危険が発生した正確なソース行を示せます。
nvcc -lineinfo reduce.cu -o reduce修正を検証する
バリアを追加した後、racecheckを再実行します。危険が0件と報告されれば、共有メモリへのアクセス順序が正しくなったということです。
クイックチェック
racecheckが共有メモリの危険を報告した場合、通常はどのように修正しますか？
まとめ
racecheckを使って共有メモリの危険を見つけ、レポートを読み、__syncthreadsバリアで修正しました。タイミングによるバグを明らかにできました。🎯
AI チューターと学ぶ C++ — 無料
ブラウザでリアルコードを書いて実行し、24/7 の AI チューターから瞬時にサポートを受け、ウェブまたはアプリで続きから学習できます。
- コース
- 30
- レッスン
- 120
よくある質問
「racecheck で競合を探す」レッスンは無料ですか？
はい。「racecheck で競合を探す」の完全なテキストはこのウェブで無料で読めます。インタラクティブに演習し（組み込みコードエディタと24時間対応のAIチューター）、CUDA Academyコースの残りをアンロックするには、CoddyKit PROにアップグレードしてください。 CUDA Academyコースには全4レッスンが含まれています。
「racecheck で競合を探す」で何を学びますか？
共有メモリのデータ競合を検出する ブラウザで直接実行するハンズオンコードでCUDA Academyを演習し、24時間対応のAIチューターがレッスンを進める中での質問に答えます。
CUDA Academyを始めるのに経験は必要ですか？
事前経験は必要ありません。CoddyKitのCUDA Academyは初級者から上級者向けに構成されているため、ここから始めるか最初から始めて、自分のペースで進むことができます。 これはレッスン3/4です。
「racecheck で競合を探す」レッスンにはどのくらい時間がかかりますか？
ほとんどのCoddyKitレッスンは約5～10分かかります。各レッスンはコンパクトでインタラクティブなので、着実に進歩し、ウェブとアプリ全体で正確に前回の場所から再開できます。
このCUDA Academyレッスンでコードを書いて実行できますか？
はい。すべてのCUDA Academyレッスンに組み込みコードエディタが含まれているため、ブラウザでリアルコードを書いて実行し、即座のAIフィードバックを取得できます。ローカル設定は不要です。
このコースのすべてのレッスン
- cuda-gdb でカーネルをステップ実行
- memcheck でリークを見つける
- racecheck で競合を探す
- synccheck で同期エラーを検出