Go to Qiita Advent Calendar 2024 Top
LoginSignup
2
1

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

@vyyyv

[cuDF]C++ libcudfのgithubを読み解く　～cudaStreamとは～

Last updated at Posted at 2024-08-06

CUDAのデフォルトストリーム（default stream）とは、CUDAプログラムで特定のストリームを指定しない場合に使用されるストリームのことです。デフォルトストリームには以下の2種類があります。

レガシーデフォルトストリーム (Legacy Default Stream)

レガシーデフォルトストリームは、暗黙的にすべての他のストリームと同期するストリームです。これは、cudaStream_tパラメータとして0を渡すか、ストリームパラメータを指定しない場合に使用されます。レガシーデフォルトストリームの同期動作は以下の通りです:

  • レガシーストリームでアクションが行われると、まずすべてのブロッキングストリームを待機し、アクションがレガシーストリームにキューイングされます。
  • その後、すべてのブロッキングストリームはレガシーストリームを待機します。
  • これにより、デフォルトストリーム上の操作は他のすべてのストリームと同期されます[1][2][3]。

スレッドごとのデフォルトストリーム (Per-Thread Default Stream)

スレッドごとのデフォルトストリームは、スレッドごとにローカルなストリームであり、他のストリームと同期しないストリームです。これにより、スレッドごとのデフォルトストリームは非同期的に動作します。スレッドごとのデフォルトストリームの特徴は以下の通りです:

  • スレッドごとのデフォルトストリームは、他のストリームと同期しません。
  • ただし、レガシーデフォルトストリームと同じプログラム内で使用される場合は、レガシーデフォルトストリームと同期します。
  • スレッドごとのデフォルトストリームは、CU_STREAM_PER_THREADまたはcudaStreamPerThreadハンドルを使用して明示的に使用できます[1][4][8]。

使用例と同期動作

デフォルトストリームを使用する場合、以下のようなコードが典型的です:

cudaMemcpy(d_u1, h_u1, bytes, cudaMemcpyHostToDevice);
kernel_code<<<dimGrid, dimBlock>>>(d_u1, d_u2);
cudaMemcpy(h_u1, d_u1, bytes, cudaMemcpyDeviceToHost);

このコードでは、すべての操作がデフォルトストリームで実行されるため、各操作は前の操作が完了するまで待機します。これにより、操作が順次実行されることが保証されます[7]。

非同期ストリームとの違い

非同期ストリーム（non-blocking stream）は、cudaStreamNonBlockingフラグを使用して作成され、デフォルトストリームと同期しないため、並行して実行されることが可能です。これにより、パフォーマンスの向上が期待できます[3][6]。

デフォルトストリームの理解と適切な使用は、CUDAプログラムのパフォーマンスと正確性に大きく影響します。

Citations:
[1] https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html
[2] https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf
[3] https://on-demand.gputechconf.com/gtc/2014/presentations/S4158-cuda-streams-best-practices-common-pitfalls.pdf
[4] https://leimao.github.io/blog/CUDA-Stream/
[5] https://stackoverflow.com/questions/31458016/in-cuda-is-it-guaranteed-that-the-default-stream-equals-nullptr
[6] https://hpcworld.jp/techcolumn/cuda-stream%E3%82%92%E7%94%A8%E3%81%84%E3%81%9F%E4%B8%A6%E8%A1%8C%E5%87%A6%E7%90%86/
[7] https://people.maths.ox.ac.uk/~gilesm/cuda/lecs/lec6.pdf
[8] https://leimao.github.io/blog/CUDA-Default-Stream/

CUDAのストリーム（Stream）は、CUDAプログラムにおいて非同期的な操作を管理するための概念です。ストリームを使用することで、異なるデータ転送やカーネル実行を並行して行うことができ、GPUのリソースを効率的に利用することが可能になります。

CUDAストリームとは

CUDAストリームは、以下のような非同期操作のシーケンスを管理するためのキューです：

  • カーネルの実行
  • メモリ転送（ホストとデバイス間、デバイス間）
  • メモリの初期化

各ストリームは独立して動作し、異なるストリームにキューイングされた操作は並行して実行される可能性があります。

ストリームの作成と使用

ストリームはcudaStreamCreate関数を使用して作成され、cudaStreamDestroy関数で破棄されます。以下はその例です：

cudaStream_t stream;
cudaStreamCreate(&stream);

// 非同期メモリ転送
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

// 非同期カーネル実行
my_kernel<<<grid, block, 0, stream>>>(d_data);

// ストリームの破棄
cudaStreamDestroy(stream);

ストリームの利点

ストリームを使用することで、以下の利点があります：

  1. 非同期実行: ストリームを使用することで、データ転送とカーネル実行を非同期に行うことができます。これにより、ホストとデバイス間のデータ転送中にGPUがアイドル状態になるのを防ぎ、リソースの効率的な利用が可能になります。

  2. 並行実行: 異なるストリームにキューイングされた操作は並行して実行される可能性があるため、複数のカーネルやデータ転送を同時に実行することができます。これにより、全体的なパフォーマンスが向上します。

  3. 同期の柔軟性: ストリームを使用することで、特定の操作間でのみ同期を取ることができ、全体の同期を避けることができます。これにより、プログラムの柔軟性と効率が向上します。

ストリームがない場合

ストリームを使用しない場合、すべての操作はデフォルトストリームで実行されます。デフォルトストリームでは、各操作は前の操作が完了するまで待機します。これにより、以下のような制約があります：

  • 非効率なリソース利用: データ転送中にGPUがアイドル状態になるため、リソースの利用効率が低下します。
  • 並行実行の欠如: 複数の操作を並行して実行することができないため、全体的なパフォーマンスが低下します。
  • 柔軟性の欠如: すべての操作が順次実行されるため、特定の操作間でのみ同期を取ることができません。

まとめ

CUDAストリームは、非同期的な操作の管理と並行実行を可能にするための重要なツールです。ストリームを使用することで、GPUリソースの効率的な利用、全体的なパフォーマンスの向上、プログラムの柔軟性の向上が期待できます。ストリームを適切に使用することは、CUDAプログラムの最適化において重要な要素です。

スレッドごとのデフォルトストリーム（Per-Thread Default Stream）は、各スレッドに対して独立したデフォルトストリームを提供する機能です。このストリームは他のストリームと同期しないため、非同期処理を効率的に行うことができます。以下に、その使い方を解説します。

スレッドごとのデフォルトストリームの有効化

スレッドごとのデフォルトストリームを有効にするには、コンパイル時に特定のオプションやマクロを設定する必要があります。

コンパイル時オプションを使用する方法

nvccコンパイラのオプション--default-stream per-threadを使用します。

nvcc --default-stream per-thread -o my_program my_program.cu

マクロを使用する方法

ソースコード内でCUDA_API_PER_THREAD_DEFAULT_STREAMマクロを定義します。このマクロは、CUDAヘッダーファイルをインクルードする前に定義する必要があります。

#define CUDA_API_PER_THREAD_DEFAULT_STREAM
#include <cuda_runtime.h>

スレッドごとのデフォルトストリームの使用例

以下は、スレッドごとのデフォルトストリームを使用して非同期処理を行う例です。

例1: 基本的な使用例

#include <cuda_runtime.h>
#include <iostream>

__global__ void my_kernel(int *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = idx;
}

int main() {
    const int N = 1024;
    int *d_data;
    cudaMalloc(&d_data, N * sizeof(int));

    // スレッドごとのデフォルトストリームを使用してカーネルを起動
    my_kernel<<<N / 256, 256>>>(d_data);

    cudaDeviceSynchronize();

    cudaFree(d_data);
    return 0;
}

例2: 複数スレッドでの使用

スレッドごとのデフォルトストリームを使用して、複数のホストスレッドから非同期処理を行う例です。

#include <cuda_runtime.h>
#include <thread>
#include <iostream>

__global__ void my_kernel(int *data, int value) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    data[idx] = value;
}

void launch_kernel(int *d_data, int value) {
    // スレッドごとのデフォルトストリームを使用してカーネルを起動
    my_kernel<<<1024 / 256, 256>>>(d_data, value);
    cudaDeviceSynchronize();
}

int main() {
    const int N = 1024;
    int *d_data;
    cudaMalloc(&d_data, N * sizeof(int));

    // 複数のホストスレッドを作成し、各スレッドでカーネルを起動
    std::thread t1(launch_kernel, d_data, 1);
    std::thread t2(launch_kernel, d_data, 2);

    t1.join();
    t2.join();

    cudaFree(d_data);
    return 0;
}

注意点

  • 同期の管理: スレッドごとのデフォルトストリームは他のストリームと同期しないため、必要に応じてcudaDeviceSynchronizeやcudaStreamSynchronizeを使用して明示的に同期を取る必要があります。
  • パフォーマンスの確認: 非同期処理のパフォーマンスを確認するために、Nsight SystemsやNsight Computeなどのプロファイラを使用することが推奨されます。

スレッドごとのデフォルトストリームを適切に使用することで、CUDAプログラムの非同期処理を効率化し、GPUリソースの利用効率を向上させることができます。

Citations:
[1] https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html
[2] https://qiita.com/kari_tech/items/a3891f13f3acc3575c08
[3] http://www.sim.gsic.titech.ac.jp/Japanese/GPU/pdf/20131226717065-477-0-41.pdf
[4] https://docs.nvidia.com/cuda/archive/12.1.0/cuda-c-programming-guide/contents.html
[5] https://qiita.com/dandelion1124/items/d069007e2a60021c0d19
[6] https://hpcworld.jp/techcolumn/cuda-stream%E3%82%92%E7%94%A8%E3%81%84%E3%81%9F%E4%B8%A6%E8%A1%8C%E5%87%A6%E7%90%86/

cudaDeviceSynchronizeとcudaStreamSynchronizeは、CUDAプログラムにおける同期操作を行うための関数ですが、それぞれの役割と使用シナリオは異なります。以下にその違いと使用方法について詳しく説明します。

cudaDeviceSynchronize

概要

cudaDeviceSynchronizeは、GPU上で実行中のすべての操作が完了するまでホスト（CPU）を待機させる関数です。これにより、すべてのストリームで実行中のカーネルやメモリ転送が完了するまでホストコードの実行を停止します。

使用例

cudaDeviceSynchronize();

使用シナリオ

  • デバッグ: 非同期操作の完了を確認するために使用されます。これにより、エラーが発生した場所を特定しやすくなります。
  • タイミング測定: カーネルの実行時間を正確に測定するために、カーネル実行前後にcudaDeviceSynchronizeを挿入します。
  • ホストとデバイス間の同期: カーネル実行後にホスト側で結果を使用する場合、カーネルが完了したことを保証するために使用します。

注意点

cudaDeviceSynchronizeを頻繁に使用すると、ホストが待機する時間が増え、全体的なパフォーマンスが低下する可能性があります[1][2][3]。

cudaStreamSynchronize

概要

cudaStreamSynchronizeは、指定されたストリーム内のすべての操作が完了するまでホストを待機させる関数です。特定のストリームに対してのみ同期を行うため、他のストリームの操作には影響を与えません。

使用例

cudaStream_t stream;
cudaStreamCreate(&stream);

// 非同期カーネル実行
my_kernel<<<grid, block, 0, stream>>>(d_data);

// 特定のストリームの同期
cudaStreamSynchronize(stream);

cudaStreamDestroy(stream);

使用シナリオ

  • 特定のストリームの同期: 複数のストリームを使用している場合、特定のストリームの操作が完了したことを確認するために使用します。
  • 非同期メモリ転送の完了確認: 非同期メモリ転送が完了したことを確認するために使用します。

注意点

cudaStreamSynchronizeは特定のストリームに対してのみ同期を行うため、他のストリームの操作は並行して実行され続けます[3][6]。

まとめ

機能 cudaDeviceSynchronize cudaStreamSynchronize
同期対象 すべてのストリーム 指定されたストリーム
使用シナリオ デバッグ、タイミング測定、全体同期 特定ストリームの同期、非同期操作の完了確認
パフォーマンスへの影響 高（ホストがすべての操作を待機） 低（特定ストリームのみ待機）

cudaDeviceSynchronizeとcudaStreamSynchronizeは、それぞれ異なる用途で使用されるため、適切なシナリオで使い分けることが重要です。全体的な同期が必要な場合はcudaDeviceSynchronizeを使用し、特定のストリームのみを同期したい場合はcudaStreamSynchronizeを使用します。

今回のPGM utilities

NVIDIA RAPIDS cuDFのutilitiesディレクトリには、cuDFライブラリの様々な機能をサポートするユーティリティ関数が含まれています。各ファイルの主要な関数と機能を以下に解説します。

cuda_memcpy.cu

このファイルはCUDAメモリコピー操作に関連する関数を提供します。

  • cuda_memcpy: デバイス間またはデバイスとホスト間のメモリコピーを行います。
  • cuda_memcpy_async: 非同期的なメモリコピーを実行します。
  • cuda_memset: CUDAデバイスメモリを特定の値で初期化します。

default_stream.cpp

CUDAのデフォルトストリーム操作を管理します。

  • get_default_stream: デフォルトのCUDAストリームを取得します。
  • set_default_stream: 新しいデフォルトストリームを設定します。

host_memory.cpp

ホストメモリの割り当てと管理に関する関数を提供します。

  • get_host_memory_resource: ホストメモリリソースを取得します。
  • set_host_memory_resource: カスタムホストメモリリソースを設定します。

linked_column.cpp

リンクされたカラムの操作を扱います。

  • create_linked_column: 既存のカラムにリンクされた新しいカラムを作成します。
  • update_linked_column: リンクされたカラムの内容を更新します。

logger.cpp

cuDFのロギング機能を実装します。

  • set_logging_level: ログレベルを設定します。
  • log_message: 指定されたレベルでメッセージをログに記録します。

prefetch.cpp

データのプリフェッチ操作を行う関数を提供します。

  • prefetch_data: 指定されたデータをGPUメモリにプリフェッチします。

stacktrace.cpp

スタックトレース機能を実装します。

  • capture_stacktrace: 現在のスタックトレースをキャプチャします。
  • print_stacktrace: キャプチャされたスタックトレースを出力します。

stream_pool.cpp

CUDAストリームプールの管理を行います。

  • get_stream: ストリームプールからストリームを取得します。
  • return_stream: 使用済みのストリームをプールに返却します。

traits.cpp

型特性（traits）に関する機能を提供します。

  • is_numeric: 型が数値型かどうかを判定します。
  • is_floating_point: 型が浮動小数点型かどうかを判定します。

type_checks.cpp

型チェックに関する関数を実装します。

  • is_same_type: 2つの型が同じかどうかを確認します。
  • is_convertible: ある型から別の型への変換が可能かどうかをチェックします。

type_dispatcher.cpp

型ディスパッチ機能を提供します。

  • dispatch_type: 与えられた型に基づいて適切な関数を呼び出します。

これらのユーティリティ関数は、cuDFライブラリ全体で使用され、メモリ管理、ストリーム操作、型チェック、ロギングなどの基本的な機能を提供します。これにより、cuDFの主要な機能を効率的に実装し、パフォーマンスを最適化することができます。

Citations:
[1] https://github.com/rapidsai/cudf/tree/branch-24.10/cpp/src/utilities
[2] https://github.com/rapidsai/cudf/blob/branch-24.10/cpp/src/utilities/cuda_memcpy.cu
[3] https://github.com/rapidsai/cudf/blob/branch-24.10/cpp/src/u

CUDAのデフォルトストリーム（default stream）とは、CUDAプログラムで特定のストリームを指定しない場合に使用されるストリームのことです。デフォルトストリームには以下の2種類があります。

2
1
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
2
1

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?