cuda异步内存复制

在cuda handbook中发现了一个利用cuda event来进行的异步的内存的复制，觉得比较巧妙，但是美中不足的是书中代码里存在一个bug，不过思想还是很不错的

利用两个buffer来充当异步复制的缓冲区

1
void *g_hostBuffers[2];

定义需要的两个event,这里定义了之后就立马record，以便将来函数内第一次Sync时不会出错。

1
2
3
4
5
cudaEventCreate( &g_events[0] );
cudaEventCreate( &g_events[1] );
// record events so they are signaled on first synchronize
cudaEventRecord( g_events[0], 0 );
cudaEventRecord( g_events[1], 0 );

之后的内存分配和数据初始化部分我们就省略掉。接下来是最重要的函数部分

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
void
chMemcpyHtoD( void *device, const void *host, size_t N ) 
{
    cudaError_t status;
    char *dst = (char *) device;
    const char *src = (const char *) host;
    int stagingIndex = 0;
    while ( N ) {
        size_t thisCopySize = min( N, STAGING_BUFFER_SIZE );
        cudaEventSynchronize( g_events[stagingIndex] );
        memcpy( g_hostBuffers[stagingIndex], src, thisCopySize ); 
        cudaMemcpyAsync( dst, g_hostBuffers[stagingIndex],thisCopySize, cudaMemcpyHostToDevice, NULL );
        cudaEventRecord( g_events[stagingIndex], NULL );
        dst += thisCopySize;
        src += thisCopySize;
        N -= thisCopySize;
        stagingIndex = 1 - stagingIndex;
    }
    return;
}

这里我们可以看到在第11行我们把src里的数据拷贝到缓冲区中，之后12行执行一个异步的复制操作。
这个复制操作由于是异步的，所以很快就会返回，执行下一步记录的操作，这一轮的复制操作就结束了，之后更新index，把1变成0,0变成1。
当我们执行到下一轮的cudaEventSynchronize语句时，我们会等待index所指的事件结束，由于事件是在复制之后被记录的，事件结束也意味着异步复制的结束。这一步是保证当前index的下一个复制时前一个复制已经结束，这样就可以保证两个时间分别进行两个异步的内存复制。

这里上述的代码是正确的，原书中的代码第13行是错误的，原书中g_event的索引是错误的，为g_events[1-stagingIndex]，这样会导致整个流程都在串行，因为刚刚记录的g_event[0]在下一轮就立马执行了cudaEventSynchronize。

实验也证明了这一点，原书中的代码在实验室的TITAN V 显卡下的拷贝性能为2987.32 MB/s，而修改后上述的正确的代码的性能为4052.28 MB/s