CUDA C++ 重叠 SERIAL 内核执行和数据传输
CUDA C++ overlapping SERIAL kernel execution and data transfer
所以本指南here展示了重叠内核执行和数据传输的一般方法。
cudaStream_t streams[nStreams];
for (int i = 0; i < nStreams; ++i) {
cudaStreamCreate(&streams[i]);
int offset = ...;
cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
// edit: no deviceToHost copy
}
但是,内核是串行的。所以它必须处理0->1000,然后1000->2000,...简而言之,重叠数据传输时正确执行此内核的顺序是:
- copy[a->b] 必须发生在 kernel[a->b]
之前
- kernel [a->b] 必须发生在 kernel[b->c] 之前,其中 c > a, b
是否可以在不使用 cudaDeviceSynchronize()
的情况下执行此操作?如果没有,最快的方法是什么?
因此每个内核都依赖于(不能开始直到):
- 关联的 H->D 拷贝完成
- 之前的内核执行完成
普通流语义无法处理这种情况(2 个独立的依赖项,来自 2 个独立的流),因此我们需要在其中放置一个额外的互锁。我们可以使用一组事件和cudaStreamWaitEvent()
来处理它。
对于最一般的情况(不知道块的总数)我会推荐这样的东西:
$ cat t1783.cu
#include <iostream>
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
template <typename T>
__global__ void process(const T * __restrict__ in, const T * __restrict__ prev, T * __restrict__ out, size_t ds){
for (size_t i = threadIdx.x+blockDim.x*blockIdx.x; i < ds; i += gridDim.x*blockDim.x){
out[i] = in[i] + prev[i];
}
}
const int nTPB = 256;
typedef int mt;
const int chunk_size = 1048576;
const int data_size = 10*1048576;
const int ns = 3;
int main(){
mt *din, *dout, *hin, *hout;
cudaStream_t str[ns];
cudaEvent_t evt[ns];
for (int i = 0; i < ns; i++) {
cudaStreamCreate(str+i);
cudaEventCreate( evt+i);}
cudaMalloc(&din, sizeof(mt)*data_size);
cudaMalloc(&dout, sizeof(mt)*data_size);
cudaHostAlloc(&hin, sizeof(mt)*data_size, cudaHostAllocDefault);
cudaHostAlloc(&hout, sizeof(mt)*data_size, cudaHostAllocDefault);
cudaMemset(dout, 0, sizeof(mt)*chunk_size); // for first loop iteration
for (int i = 0; i < data_size; i++) hin[i] = 1;
cudaEventRecord(evt[ns-1], str[ns-1]); // this event will immediately "complete"
unsigned long long dt = dtime_usec(0);
for (int i = 0; i < (data_size/chunk_size); i++){
cudaStreamSynchronize(str[i%ns]); // so we can reuse event safely
cudaMemcpyAsync(din+i*chunk_size, hin+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyHostToDevice, str[i%ns]);
cudaStreamWaitEvent(str[i%ns], evt[(i>0)?(i-1)%ns:ns-1], 0);
process<<<(chunk_size+nTPB-1)/nTPB, nTPB, 0, str[i%ns]>>>(din+i*chunk_size, dout+((i>0)?(i-1)*chunk_size:0), dout+i*chunk_size, chunk_size);
cudaEventRecord(evt[i%ns]);
cudaMemcpyAsync(hout+i*chunk_size, dout+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyDeviceToHost, str[i%ns]);
}
cudaDeviceSynchronize();
dt = dtime_usec(dt);
for (int i = 0; i < data_size; i++) if (hout[i] != (i/chunk_size)+1) {std::cout << "error at index: " << i << " was: " << hout[i] << " should be: " << (i/chunk_size)+1 << std::endl; return 0;}
std::cout << "elapsed time: " << dt << " microseconds" << std::endl;
}
$ nvcc -o t1783 t1783.cu
$ ./t1783
elapsed time: 4366 microseconds
此处的良好做法是使用探查器来验证预期的重叠场景。但是,我们可以根据经过的时间测量走捷径。
循环将总共 40MB 的数据传输到设备,并返回 40MB。经过的时间是 4366us。这给出了 (40*1048576)/4366 或 9606 bytes/us 每个方向的平均吞吐量,即 9.6GB/s。这基本上是在两个方向上使 Gen3 link 饱和,因此我的块处理大约是 back-to-back,并且我基本上完全重叠了 D->H 和 H->D 内存拷贝。这里的内核是微不足道的,所以它在配置文件中显示为碎片。
对于您的情况,您表示不需要 D->H 副本,但它不会增加额外的复杂性,所以我选择展示它。如果您在循环外注释该行,仍会出现所需的行为(尽管这会影响稍后的结果检查)。
对这种方法的一个可能的批评是 cudaStreamSynchronize()
调用是必要的,因此我们不会“超出”事件互锁,这意味着循环将只进行到 ns
号超出当前在设备上执行的迭代次数。因此不可能异步启动比这更多的工作。如果您想一次启动所有工作并继续在 CPU 上做其他事情,此方法将不会完全允许(当流处理达到时 CPU 将继续通过循环ns
迭代从最后一个)。
提供代码是为了从概念上说明一种方法。它不保证没有缺陷,我也不声称它适合任何特定用途。
所以本指南here展示了重叠内核执行和数据传输的一般方法。
cudaStream_t streams[nStreams];
for (int i = 0; i < nStreams; ++i) {
cudaStreamCreate(&streams[i]);
int offset = ...;
cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i]);
kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
// edit: no deviceToHost copy
}
但是,内核是串行的。所以它必须处理0->1000,然后1000->2000,...简而言之,重叠数据传输时正确执行此内核的顺序是:
- copy[a->b] 必须发生在 kernel[a->b] 之前
- kernel [a->b] 必须发生在 kernel[b->c] 之前,其中 c > a, b
是否可以在不使用 cudaDeviceSynchronize()
的情况下执行此操作?如果没有,最快的方法是什么?
因此每个内核都依赖于(不能开始直到):
- 关联的 H->D 拷贝完成
- 之前的内核执行完成
普通流语义无法处理这种情况(2 个独立的依赖项,来自 2 个独立的流),因此我们需要在其中放置一个额外的互锁。我们可以使用一组事件和cudaStreamWaitEvent()
来处理它。
对于最一般的情况(不知道块的总数)我会推荐这样的东西:
$ cat t1783.cu
#include <iostream>
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
template <typename T>
__global__ void process(const T * __restrict__ in, const T * __restrict__ prev, T * __restrict__ out, size_t ds){
for (size_t i = threadIdx.x+blockDim.x*blockIdx.x; i < ds; i += gridDim.x*blockDim.x){
out[i] = in[i] + prev[i];
}
}
const int nTPB = 256;
typedef int mt;
const int chunk_size = 1048576;
const int data_size = 10*1048576;
const int ns = 3;
int main(){
mt *din, *dout, *hin, *hout;
cudaStream_t str[ns];
cudaEvent_t evt[ns];
for (int i = 0; i < ns; i++) {
cudaStreamCreate(str+i);
cudaEventCreate( evt+i);}
cudaMalloc(&din, sizeof(mt)*data_size);
cudaMalloc(&dout, sizeof(mt)*data_size);
cudaHostAlloc(&hin, sizeof(mt)*data_size, cudaHostAllocDefault);
cudaHostAlloc(&hout, sizeof(mt)*data_size, cudaHostAllocDefault);
cudaMemset(dout, 0, sizeof(mt)*chunk_size); // for first loop iteration
for (int i = 0; i < data_size; i++) hin[i] = 1;
cudaEventRecord(evt[ns-1], str[ns-1]); // this event will immediately "complete"
unsigned long long dt = dtime_usec(0);
for (int i = 0; i < (data_size/chunk_size); i++){
cudaStreamSynchronize(str[i%ns]); // so we can reuse event safely
cudaMemcpyAsync(din+i*chunk_size, hin+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyHostToDevice, str[i%ns]);
cudaStreamWaitEvent(str[i%ns], evt[(i>0)?(i-1)%ns:ns-1], 0);
process<<<(chunk_size+nTPB-1)/nTPB, nTPB, 0, str[i%ns]>>>(din+i*chunk_size, dout+((i>0)?(i-1)*chunk_size:0), dout+i*chunk_size, chunk_size);
cudaEventRecord(evt[i%ns]);
cudaMemcpyAsync(hout+i*chunk_size, dout+i*chunk_size, sizeof(mt)*chunk_size, cudaMemcpyDeviceToHost, str[i%ns]);
}
cudaDeviceSynchronize();
dt = dtime_usec(dt);
for (int i = 0; i < data_size; i++) if (hout[i] != (i/chunk_size)+1) {std::cout << "error at index: " << i << " was: " << hout[i] << " should be: " << (i/chunk_size)+1 << std::endl; return 0;}
std::cout << "elapsed time: " << dt << " microseconds" << std::endl;
}
$ nvcc -o t1783 t1783.cu
$ ./t1783
elapsed time: 4366 microseconds
此处的良好做法是使用探查器来验证预期的重叠场景。但是,我们可以根据经过的时间测量走捷径。
循环将总共 40MB 的数据传输到设备,并返回 40MB。经过的时间是 4366us。这给出了 (40*1048576)/4366 或 9606 bytes/us 每个方向的平均吞吐量,即 9.6GB/s。这基本上是在两个方向上使 Gen3 link 饱和,因此我的块处理大约是 back-to-back,并且我基本上完全重叠了 D->H 和 H->D 内存拷贝。这里的内核是微不足道的,所以它在配置文件中显示为碎片。
对于您的情况,您表示不需要 D->H 副本,但它不会增加额外的复杂性,所以我选择展示它。如果您在循环外注释该行,仍会出现所需的行为(尽管这会影响稍后的结果检查)。
对这种方法的一个可能的批评是 cudaStreamSynchronize()
调用是必要的,因此我们不会“超出”事件互锁,这意味着循环将只进行到 ns
号超出当前在设备上执行的迭代次数。因此不可能异步启动比这更多的工作。如果您想一次启动所有工作并继续在 CPU 上做其他事情,此方法将不会完全允许(当流处理达到时 CPU 将继续通过循环ns
迭代从最后一个)。
提供代码是为了从概念上说明一种方法。它不保证没有缺陷,我也不声称它适合任何特定用途。