現在のプロジェクトでは、信号処理と視覚化にGPUを使用しています。私はすでにストリームを使用して非同期操作を可能にしています。信号はフレームで処理され、各フレームのストリームの処理ステップは次のとおりです。
現在、手順は単一のGPUで行われていますが、私のマシンにはマルチGPUカード（GeForce GTX 690）が搭載されており、2つのデバイス間で操作を分散したいと思います。基本的に、操作1、2、3、および4を単一の非同期ストリームとして実行しながら、デバイスAでステップ1と2を実行し、デバイスBでステップ3と4を実行したいと思います。望ましい結果は、次のようなストリーミングレイアウトです。
Device A Stream a 1 2 1 2 ...
Stream b 1 2 ...
Device B Stream a 3 4 3 4 ...
Stream b 3 4 ...
これどうやってするの？
ストリームが作成されたデバイスに関連付けられているため、以前の試みは正しくありませんでした。ですから、タイトルにあるあなたの質問に対する最も直接的な答えは「できない」だと思います。単一のストリームを作成して、そこから複数のGPUにコマンドを発行することはできません。ここから：
Stream and Event Behavior
A kernel launch or memory copy will fail if it is issued to a stream that is not associated to the current device
ただし、調査中に、イベントは2つの異なるデバイスで2つのストリームを同期するための推奨される方法であることに気付きました。
cudaStreamWaitEvent()入力ストリームと入力イベントが異なるデバイスに関連付けられている場合でも成功します。したがって、cudaStreamWaitEvent（）を使用して、複数のデバイスを相互に同期させることができます。
そのため、これを説明するために次のコードを作成しました。
#include <stdio.h>
#define SIZE 32
#define K1VAL 5
#define K3VAL 3
#define cudaCheckErrors(msg) \
do { \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) { \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
} \
} while (0)
__global__ void kernel1(int *frame, int size){
int idx = threadIdx.x + (blockDim.x * blockIdx.x);
if (idx == 0){
int *a = new int[10000]; // just to make this kernel take a while
for (int i = 0; i<10000; i++)
a[i] = 0;
for (int i = 0; i < size; i++)
frame[i] += K1VAL;
}
}
__global__ void kernel3(int *frame, int size){
int idx = threadIdx.x + (blockDim.x * blockIdx.x);
if (idx == 0)
for (int i = 0; i < size; i++)
frame[i] -= K3VAL;
}
void set_device(int dev){
int ldev;
cudaSetDevice(dev);
cudaGetDevice(&ldev);
cudaCheckErrors("set device error");
if (ldev != dev){
printf("set device mismatch error\n");
exit(1);
}
}
int main(){
int A=0;
int B=1;
int framesize = SIZE*sizeof(int);
int *h_frame;
int *d_frame_aA, *d_frame_bB;
int numdev = 0;
cudaGetDeviceCount(&numdev);
cudaCheckErrors("can't determine number of devices");
if (numdev < 2){
printf("not enough devices!\n");
return 1;
}
set_device(A);
cudaMalloc((void **) &d_frame_aA, framesize); // stream_a
cudaMemset(d_frame_aA, 0, framesize);
set_device(B);
cudaMalloc((void **) &d_frame_bB, framesize); // stream_b
cudaMemset(d_frame_bB, 0, framesize);
cudaHostAlloc((void **) &h_frame, framesize, cudaHostAllocDefault);
cudaCheckErrors("allocations failure");
set_device(A);
cudaStream_t stream_a, stream_b;
cudaStreamCreate(&stream_a);
cudaEvent_t absync;
cudaEventCreate(&absync);
set_device(B);
cudaStreamCreate(&stream_b);
cudaCheckErrors("stream creation failure");
for (int i = 0; i < SIZE; i++)
h_frame[i] = 0;
set_device(A);
cudaDeviceEnablePeerAccess(B, 0);
set_device(B);
cudaDeviceEnablePeerAccess(A, 0);
cudaCheckErrors("enable peer access fail");
set_device(A);
cudaMemcpyAsync(d_frame_aA, h_frame, framesize, cudaMemcpyHostToDevice, stream_a);
kernel1<<<1,1,0, stream_a>>>(d_frame_aA, SIZE);
cudaCheckErrors("kernel1 fail");
cudaMemcpyPeerAsync(d_frame_bB, B, d_frame_aA, A, framesize, stream_a );
cudaCheckErrors("memcpypeer fail");
cudaEventRecord(absync, stream_a);
set_device(B);
// comment out the next line to see the failure
cudaStreamWaitEvent(stream_b, absync, 0);
kernel3<<<1,1,0, stream_b>>>(d_frame_bB, SIZE);
cudaCheckErrors("main sequence fail");
// cudaCheckErrors("main sequence failure");
cudaMemcpy(h_frame, d_frame_bB, framesize, cudaMemcpyDeviceToHost);
cudaCheckErrors("results_a memcpy fail");
for (int i = 0; i < SIZE; i++)
if (h_frame[i] != (K1VAL - K3VAL)) {
printf("results error\n");
return 1;
}
printf("success\n");
return 0;
}
コードをそのまま実行すると、successメッセージが表示されます。ストリームb（デバイスB）がストリームa（デバイスA）で待機するように強制する行をコメントアウトすると、results errorメッセージが表示されます。したがって、これは、あるデバイスのストリームを別のデバイスのストリームに同期する方法を示しています。それが役に立てば幸い。最初のラウンドで混乱してすみません。
cudaStreamWaitEvent()別のデバイスに属する CUDA イベントに待機を挿入できるため、GPU 間同期を有効にします。
したがって、プロデューサーとコンシューマーの間の GPU 間同期に必要なのは、2 つの GPU のそれぞれにいくつかのイベント (少なくとも 2 つ) を割り当て、プロデューサーcudaEventRecord()とコンシューマーcudaStreamWaitEvent()を同じイベントに配置することです。 cudaStreamWaitEvent()現在のデバイスのコマンド バッファにコマンドを挿入し、指定されたイベントが記録されるまで実行を中断します。
cudaStreamWaitEvent()以下は、この方法を使用してピアツーピア memcpy が実装されているコード フラグメントを示しています。ポンプがプライミングされると、プロデューサとコンシューマの両方が同時に PCIe 転送を実行し、それぞれが 2 つのステージング バッファ (ポータブル固定メモリに割り当てられている) の 1 つに転送されます。
cudaError_t
chMemcpyPeerToPeer(
void *_dst, int dstDevice,
const void *_src, int srcDevice,
size_t N )
{
cudaError_t status;
char *dst = (char *) _dst;
const char *src = (const char *) _src;
int stagingIndex = 0;
while ( N ) {
size_t thisCopySize = min( N, STAGING_BUFFER_SIZE );
CUDART_CHECK( cudaSetDevice( srcDevice ) );
CUDART_CHECK( cudaStreamWaitEvent( NULL, g_events[dstDevice][stagingIndex], 0 ) );
CUDART_CHECK( cudaMemcpyAsync( g_hostBuffers[stagingIndex], src, thisCopySize,
cudaMemcpyDeviceToHost, NULL ) );
CUDART_CHECK( cudaEventRecord( g_events[srcDevice][stagingIndex] ) );
CUDART_CHECK( cudaSetDevice( dstDevice ) );
CUDART_CHECK( cudaStreamWaitEvent( NULL, g_events[srcDevice][stagingIndex], 0 ) );
CUDART_CHECK( cudaMemcpyAsync( dst, g_hostBuffers[stagingIndex], thisCopySize,
cudaMemcpyHostToDevice, NULL ) );
CUDART_CHECK( cudaEventRecord( g_events[dstDevice][stagingIndex] ) );
dst += thisCopySize;
src += thisCopySize;
N -= thisCopySize;
stagingIndex = 1 - stagingIndex;
}
// Wait until both devices are done
CUDART_CHECK( cudaSetDevice( srcDevice ) );
CUDART_CHECK( cudaDeviceSynchronize() );
CUDART_CHECK( cudaSetDevice( dstDevice ) );
CUDART_CHECK( cudaDeviceSynchronize() );
Error:
return status;
}
https://github.com/ArchaeaSoftware/cudahandbook/blob/master/memory/peer2peerMemcpy.cuの完全なソース コード