如何解决CUDA C ++重叠SERIAL内核执行和数据传输
因此,本指南here显示了重叠内核执行和数据传输的一般方法。
{
"example1": {
"example2": {
"example3": 4
}
}
}
但是,内核是串行的。因此,它必须处理0-> 1000,然后处理1000-> 2000,...总之,在重叠数据传输时正确执行此内核的顺序为:
- 复制[a-> b]必须在内核[a-> b]之前发生
- 内核[a-> b]必须在内核[b-> c]之前发生,其中c> a,b
是否可以不使用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,stream[i]>>>(d_a,offset);
// edit: no deviceToHost copy
}
来执行此操作?如果没有,最快的方法是什么?
解决方法
因此每个内核都依赖(直到开始):
- 相关的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,cudaHostAllocDefault);
cudaMemset(dout,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,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,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字节/ us,为9.6GB / s。这基本上使Gen3链接在两个方向上都处于饱和状态,因此我的块处理大致是背对背的,并且我基本上将D-> H与H-> D记忆副本完全重叠。这里的内核是微不足道的,因此在配置文件中显示为条状。
对于您的情况,您表示不需要D-> H副本,但是它没有增加额外的复杂性,因此我选择显示它。如果您在循环外注释该行,则仍然会发生所需的行为(尽管这会影响以后的结果检查)。
对此方法的一种可能的批评是,cudaStreamSynchronize()
调用是必要的,这样我们就不会“超载”事件互锁,这意味着循环只会进行ns
次迭代超出设备上当前正在执行的操作。因此,不可能异步启动更多的工作。如果您想一次启动所有工作,然后继续在CPU上执行其他操作,则此方法将无法完全允许这样做(当流处理已从CPU到达ns
个迭代时,CPU会经过循环。最后一个)。
呈现该代码从概念上说明了一种方法。它不能保证没有缺陷,也不保证它适用于任何特定目的。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。