CUDA 기록 및 이벤트가 작동하지 않을 때까지 기다려주십시오?

Question

내 설정의 골격을 따릅니다. 이렇게 실행하면 올바른 결과를 얻지 못합니다. 이것은 커널이 커널을 사용할 때 완료하지 않은 비동기 데이터 전송 때문일 가능성이 큽니다. 프리 프로세서 if-else 문을 사용하여 "failsafe"버전을 구현했습니다. else 부분을 번역 할 때 프로그램이 정상적으로 실행됩니다. 나는 그것을 얻지 않는다. 왜?

in1, out1 등은 단지 자리 표시 자입니다. 물론 그들은 for 루프의 반복마다 다른 컨테이너를 가리 킵니다. 비동기 전송이 발생할 수 있습니다. 그러나 반복에있어서 out1은 전송에 사용되고 하나는 커널에 의해 사용됩니다.

cudaStream_t streams[2]; 
    cudaEvent_t evCopied; 

    cudaStreamCreate(&streams[0]); // TRANSFER 
    cudaStreamCreate(&streams[1]); // KERNEL 

    cudaEventCreate(&evCopied); 

    // many iterations 
    for() { 

    // Here I want overlapping of transfers with previous kernel 
    cudaMemcpyAsync(out1, in1, size1, cudaMemcpyDefault, streams[0]); 
    cudaMemcpyAsync(out2, in2, size2, cudaMemcpyDefault, streams[0]); 
    cudaMemcpyAsync(out3, in3, size3, cudaMemcpyDefault, streams[0]); 

#if 1 
    // make sure host thread doesn't "run away" 
    cudaStreamSynchronize(streams[1]); 
    cudaEventRecord(evCopied , streams[0]); 
    cudaStreamWaitEvent(streams[1] , evCopied , 0); 
#else 
    // this gives the correct results 
    cudaStreamSynchronize(streams[0]); 
    cudaStreamSynchronize(streams[1]); 
#endif 

    kernel<<< grid , sh_mem , streams[1] >>>(out1,out2,out3); 

    } 

설정을 재정렬하라는 답변을 게시하지 마십시오. 그런 식으로 커널을 여러개의 커널로 나눠서 각각의 스트림으로 나눕니다.

Answer 1

적어도 두 개의 스트림을 동기화하기 위해 이벤트를 사용하는 것은 효과가 있습니다. 기본적으로 코드는 게시하지 않기로 선택했기 때문에 코드가 작동하지 않으며 악마는 항상 세부 정보에 표시됩니다.

그러나, 여기에 내가 당신이 시도하고있는하는 것과 유사한 방식으로 스트림 API를 사용하여 생각 완전한, 실행 가능한 예제가 제대로 작동합니다 :

#include <cstdio> 

typedef unsigned int uint; 

template<uint bsz> 
__global__ void kernel(uint * a, uint * b, uint * c, const uint N) 
{ 
    __shared__ volatile uint buf[bsz]; 
    uint tid = threadIdx.x + blockIdx.x * blockDim.x; 
    uint stride = blockDim.x * gridDim.x; 
    uint val = 0; 
    for(uint i=tid; i<N; i+=stride) { 
     val += a[i] + b[i]; 
    } 
    buf[threadIdx.x] = val; __syncthreads(); 

#pragma unroll 
    for(uint i=(threadIdx.x+warpSize); (threadIdx.x<warpSize)&&(i<bsz); i+=warpSize) 
     buf[threadIdx.x] += buf[i]; 

    if (threadIdx.x < 16) buf[threadIdx.x] += buf[threadIdx.x+16]; 
    if (threadIdx.x < 8) buf[threadIdx.x] += buf[threadIdx.x+8]; 
    if (threadIdx.x < 4) buf[threadIdx.x] += buf[threadIdx.x+4]; 
    if (threadIdx.x < 2) buf[threadIdx.x] += buf[threadIdx.x+2]; 
    if (threadIdx.x == 0) c[blockIdx.x] += buf[0] + buf[1]; 

} 

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } 
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) 
{ 
    if (code != cudaSuccess) 
    { 
     fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); 
     if (abort) exit(code); 
    } 
} 

int main(void) 
{ 
    const int nruns = 16, ntransfers = 3; 
    const int Nb = 32, Nt = 192, Nr = 3000, N = Nr * Nb * Nt; 
    const size_t szNb = Nb * sizeof(uint), szN = size_t(N) * sizeof(uint); 
    size_t sz[4] = { szN, szN, szNb, szNb }; 

    uint * d[ntransfers+1]; 
    for(int i=0; i<ntransfers+1; i++) 
     gpuErrchk(cudaMallocHost((void **)&d[i], sz[i])); 
    uint * a = d[0], * b = d[1], * c = d[2], * out = d[3]; 

    for(uint i=0; i<N; i++) { 
     a[i] = b[i] = 1; 
     if (i<Nb) c[i] = 0; 
    } 

    uint * _d[3]; 
    for(int i=0; i<ntransfers; i++) 
     gpuErrchk(cudaMalloc((void **)&_d[i], sz[i])); 
    uint * _a = _d[0], * _b = _d[1], * _c = _d[2]; 

    cudaStream_t stream[2]; 
    for (int i = 0; i < 2; i++) 
     gpuErrchk(cudaStreamCreate(&stream[i])); 

    cudaEvent_t sync_event; 
    gpuErrchk(cudaEventCreate(&sync_event)); 

    uint results[nruns]; 
    for(int j=0; j<nruns; j++) { 
     for(int i=0; i<ntransfers; i++) 
      gpuErrchk(cudaMemcpyAsync(_d[i], d[i], sz[i], cudaMemcpyHostToDevice, stream[0])); 

     gpuErrchk(cudaEventRecord(sync_event, stream[0])); 
     gpuErrchk(cudaStreamWaitEvent(stream[1], sync_event, 0)); 

     kernel<Nt><<<Nb, Nt, 0, stream[1]>>>(_a, _b, _c, N); 
     gpuErrchk(cudaPeekAtLastError()); 

     gpuErrchk(cudaMemcpyAsync(out, _c, szNb, cudaMemcpyDeviceToHost, stream[1])); 
     gpuErrchk(cudaStreamSynchronize(stream[1])); 

     results[j] = uint(0); 
     for(int i=0; i<Nb; i++) results[j]+= out[i]; 
    } 

    for(int j=0; j<nruns; j++) 
     fprintf(stdout, "%3d: ans = %u\n", j, results[j]); 

    gpuErrchk(cudaDeviceReset()); 
    return 0; 
} 

커널은 "융합이다 벡터 추가/감소 "는 말도 안되지만 커널 실행 이전에 세 입력의 마지막에 의존하여 정답을 산출합니다. 이는 입력 데이터 수의 두 배에 불과합니다. 예제에서와 같이 커널 실행과 비동기 입력 배열 복사는 서로 다른 스트림에 있으므로 복사와 실행이 중복 될 수 있습니다. 이 경우 모든 반복에서 첫 번째 두 개의 큰 입력을 복사하는 정당한 이유가 있습니다. 중요한 사본 인 마지막 사본이 완료되기 전에 지연을 도입하고 커널과 잘못 겹칠 가능성을 높입니다. 이것은 CUDA 메모리 모델이 실행중인 커널에 의해 액세스되는 메모리를 비동기 적으로 수정하는 것이 안전하다는 것을 보장하지 않기 때문에 잘못 될 수 있습니다. 수행하려는 작업이 인 경우 실패 할 것으로 예상하십시오. 그러나 실제 코드를 보지 않으면 더 말할 수 없습니다.

커널을 시작하기 전에 커널이 cudaStreamWaitEvent 두 스트림을 동기화하지 않고 올바른 결과를 내지 못한다는 것을 스스로 알 수 있습니다. 의사 코드와이 예제 간의 유일한 차이점은 실행 스트림에서 cudaStreamSynchronize의 위치입니다. 여기서 커널을 실행 한 다음에 커널을 배치하여 커널이 전송 전에 결과를 호스트로 다시 수집하는지 확인합니다. 그것은 중요한 차이 일 수 있지만, 실제 코드는 실제 코드 분석과 동등하지 않습니다. ...

내가 제안 할 수있는 것은 그것이 작동하는 방식에 대한 느낌을 얻기 위해이 예제로 노는 것입니다. 최신 버전의 Nsight for Windows에서 실행 스트림을 인위적으로 직렬화하지 않고 비동기 코드를 프로파일 링 할 가능성이 있음을 알고 있습니다. 이 예제 나 자신의 코드에서 문제를 해결할 수 없다면 문제를 진단하는 데 도움이 될 수 있습니다.