Cuda Streams的概述（六）-- 常见问题

常见问题

问题1:使用默认流

问题2：内存传输问题

问题3:隐式同步

问题4:主机端的限制

问题5:启动overhead限制

问题6:过度同步

问题7：PROFILER OVERHEAD

case 1-A

for(int i=0;i<repeat;i++) {

kernel<<<1,1,0,stream1>>>();

kernel<<<1,1>>>(); 
}

存在问题：有一个kernel在默认流里。

解决方案：每个kernel放在自己的流里。

for(int i=0;i<repeat;i++) { 
kernel<<<1,1,0,stream1>>>(); 
kernel<<<1,1,0,stream2>>>(); 
}

Case 1-B

for(int i=0;i<repeat;i++) {

kernel<<<1,1,0,stream1>>>(); 
cudaEventRecord(event1); 
kernel<<<1,1,0,stream2>>>(); 
cudaEventRecord(event2);

cudaEventSynchronize(event1); 
cudaEventSynchronize(event2);
}

存在问题：event导致的问题。

存在问题：没有流的cudaEventRecord进入默认流。

for(int i=0;i<repeat;i++) {
kernel<<<1,1,0,stream1>>>(); 
cudaEventRecord(event1); 
kernel<<<1,1,0,stream2>>>(); 
cudaEventRecord(event2);
}

解决方案：把record event放入非默认流。

for(int i=0;i<repeat;i++) {
kernel<<<1,1,0,stream1>>>(); 
cudaEventRecord(event1,stream1); 
kernel<<<1,1,0,stream2>>>(); 
cudaEventRecord(event2,stream2);

cudaEventSynchronize(event1); 
cudaEventSynchronize(event2);
}

问题1:使用默认流

• 存在现象：
  • 两个流不能重叠（在Cuda5.0中，默认流是stream2）
  • 查找cudaEventRecord(event) , cudaMemcpyAsync()
    • 如果没有指定流，就在默认流里。
  • 查找在启动默认流的Kernel
    • <<<a,b>>>
• 解决方法：
  • 把work移到非默认流里。
  • cudaEventRecord(event,stream), cudaMemcpyAsync(…,stream)
  • 可选择方法：申请其他流作为非阻塞流。

Case 2-A

for(int i=0;i<repeat;i++) { 
cudaMemcpy(d_ptr,h_ptr,bytes, cudaMemcpyHostToDevice); kernel<<<1,1,0,stream2>>>(); 
cudaDeviceSynchronize(); 
}

存在问题：内存的拷贝是同步的

for(int i=0;i<repeat;i++) { 
cudaMemcpyAsync(d_ptr,h_ptr,bytes, cudaMemcpyHostToDevice, stream1); kernel<<<1,1,0,stream2>>>(); 
cudaDeviceSynchronize(); }

解决方案：使用异步API

Case 2-B

for(int i=0;i<repeat;i++) { 
cudaMemcpyAsync(d_ptr,h_ptr,bytes, cudaMemcpyHostToDevice, stream1); kernel<<<1,1,0,stream2>>>(); 
cudaDeviceSynchronize(); }

存在问题：主机端无法使用Cuda5.5的可分页类型。

cudaHostRegister(h_ptr,bytes,0); 
for(int i=0;i<repeat;i++) {
cudaMemcpyAsync(d_ptr,h_ptr,bytes, cudaMemcpyHostToDevice, stream1);
kernel<<<1,1,0,stream2>>>();
cudaDeviceSynchronize(); 
} 
cudaHostUnregister(h_ptr);

解决方案：主机端pin内存使用cudaHostRegister或者cudaMallocHost。

问题2：内存传输问题

• 存在现象：
  • 内存拷贝没有重叠
  • 主机端花费过多时间在内存拷贝API上
  • 在Cuda5.5+，reports 可分页内存
• 解决方案：
  • 使用异步内存拷贝
  • 在主机端内存使用pinned内存。 （cudaMallocHost或cudaHostRegister）

Case 3

void launchwork(cudaStream_t stream) { 
int *mem; 
cudaMalloc(&mem,bytes); 
kernel<<<1,1,0,stream>>>(mem); 
cudaFree(mem); 
}
for(int i=0;i<repeat;i++) { 
launchwork(stream1); 
launchwork(stream2); 
}

存在问题：申请和释放在设备上是同步的。
主机端在申请和释放是会阻塞。

void launchwork(cudaStream_t stream, int *mem) { 
kernel<<<1,1,0,stream>>>(mem); 
}

for(int i=0;i<repeat;i++) { 
launchwork<1>(stream1,mem1); 
launchwork<2>(stream2,mem2); 
}

解决方案：复用cuda内存，同时将流和event加入参数。

问题3:隐式同步

• 存在现象：
  • 主机端无法go ahead
  • 主机端在某些api调用上停留过多时间。（cudaMalloc, cudaFree, cudaEventCreate, cudaEventDestroy, cudaStreamCreate, cudaStreamCreate, cudaHostRegister, cudaHostUnregister, cudaFuncSetCacheConfig）
• 解决方法
  • 复用内存和数据结构

Case 4

for(int i=0;i<repeat;i++) { 
hostwork(); 
kernel<<<1,1,0,stream1>>>(); 
hostwork(); 
kernel<<<1,1,0,stream2>>>(); 
}

存在问题：主机端限制了性能。

问题4:主机端的限制

• 存在现象：
  • 主机端在cuda APIs的外侧
  • 在主机和设备为空的地方，有大量的gaps在时间线上。
• 解决方案：
  • 把更多的工作移到GPU上。
  • 主机端使用多线程代码。

Case 5

for(int i=0;i<repeat;i++) {
kernel<<<1,1,0,stream1>>>();
kernel<<<1,1,0,stream2>>>(); 
}

存在问题：没有足够的work覆盖启动的overhead。

问题5:启动overhead限制

• 存在现象：
  • 主机端无法提前（Host does not get ahead）
  • kenel 少于 <30us
  • 连续kernel之间的时间 >10us
• 解决方案：
  • 设计运行时间更长的kernel
    • 将相邻的kernel融合在一起。
    • 批work放在一个kernel里面。
    • 解决更打的问题。

问题6:过度同步

• 存在现象：
  • 主机端无法提前。
  • 时间线的idle时间中有大量的gaps。
  • 主机端显示了同步的API调用。
• 解决方案：
  • 使用event限制同步的数量。
  • 使用cudaStreamWaitEvent组织主机端同步。
  • 使用cudaEventSynchronize

问题7：PROFILER OVERHEAD

• 存在现象：
  • 时间线中有大量的gaps，时间线同时显示了profiler的overhead
• 真实代码可能不存在同样的问题。
• 解决方案：
  • 避免cudaDeviceSynchronize()和cudaStreamSynchronize()
  • 尽量使用cudaEventRecord(event,stream)和cudaEventSynchronize(event)