CUDAの非同期memcpy中のデバイス同期

Question

CUDAで非同期memcpyホストをデバイスに実行し、直ちにカーネルを実行したいとします。非同期転送が完了したら、カーネルでどのようにテストできますか？

Answer 1

CUDA "ストリーム"を使用して非同期コピーとカーネルの起動をシーケンシングすると、非同期転送が完了した後でカーネルが確実に実行されます。次のコード例は示しています

#include <stdio.h> 

__global__ void kernel(const int *ptr) 
{ 
    printf("Hello, %d\n", *ptr); 
} 

int main() 
{ 
    int *h_ptr = 0; 

    // allocate pinned host memory with cudaMallocHost 
    // pinned memory is required for asynchronous copy 
    cudaMallocHost(&h_ptr, sizeof(int)); 

    // look for thirteen in the output 
    *h_ptr = 13; 

    // allocate device memory 
    int *d_ptr = 0; 
    cudaMalloc(&d_ptr, sizeof(int)); 

    // create a stream 
    cudaStream_t stream; 
    cudaStreamCreate(&stream); 

    // sequence the asynchronous copy on our stream 
    cudaMemcpyAsync(d_ptr, h_ptr, sizeof(int), cudaMemcpyHostToDevice, stream); 

    // sequence the kernel on our stream after the copy 
    // the kernel will execute after the copy has completed 
    kernel<<<1,1,0,stream>>>(d_ptr); 

    // clean up after ourselves 
    cudaStreamDestroy(stream); 
    cudaFree(d_ptr); 
    cudaFreeHost(h_ptr); 
} 

そして出力：私はいくつかの非同期状態かどうか（例えば完了としてカーネル内からテストするためにサポートされている任意の方法があります信じていません

$ nvcc -arch=sm_20 async.cu -run 
Hello, 13 

を非同期転送）が満たされています。 CUDAスレッドブロックは、他の実行スレッドとは完全に独立して実行されると想定されます。