标签: cuda-streams

我有一些非常类似的代码:

int k, no_streams = 4;
cudaStream_t stream[no_streams];
for(k = 0; k < no_streams; k++) cudaStreamCreate(&stream[k]);

cudaMalloc(&g_in,  size1*no_streams);
cudaMalloc(&g_out, size2*no_streams);

for (k = 0; k < no_streams; k++)
  cudaMemcpyAsync(g_in+k*size1/sizeof(float), h_ptr_in[k], size1, cudaMemcpyHostToDevice, stream[k]);

for (k = 0; k < no_streams; k++)
  mykernel<<<dimGrid, dimBlock, 0, stream[k]>>>(g_in+k*size1/sizeof(float), g_out+k*size2/sizeof(float));

for (k = 0; k < no_streams; k++)
  cudaMemcpyAsync(h_ptr_out[k], g_out+k*size2/sizeof(float), size2, cudaMemcpyDeviceToHost, stream[k]);

cudaThreadSynchronize();

cudaFree(g_in);
cudaFree(g_out);

'h_ptr_in'和'h_ptr_out'是用cudaMallocHost分配的指针数组(没有标志).

问题是流不重叠.在可视化分析器中,我可以看到第一个流中的内核执行与第二个流中的副本(H2D)重叠,但没有其他重叠.

我可能没有资源来运行2个内核(我想我这样做)但至少内核执行和副本应该重叠,对吧？如果我把所有3(复制H2D,内核执行,复制D2H)放在同一个for循环中,它们之间没有任何重叠......

请帮忙,这可能导致什么？

我正在跑步:

Ubuntu 10.04 x64

设备:"GeForce GTX 460"(CUDA驱动程序版本:3.20,CUDA运行时版本:3.20,CUDA能力主要/次要版本号:2.1,并发复制和执行:是,并发内核执行:是)

对于我的CUDA开发,我使用的是具有16个内核的机器,以及1个带有16个SM的GTX 580 GPU.对于我正在做的工作,我计划启动16个主机线程(每个核心1个),每个线程启动1个内核,每个线程包含1个块和1024个线程.我的目标是在16个SM上并行运行16个内核.这可能/可行吗？

我试图尽可能多地阅读关于独立上下文的内容,但似乎没有太多可用的信息.据我了解,每个主机线程都可以拥有自己的GPU上下文.但是,如果我使用独立的上下文,我不确定内核是否会并行运行.

我可以将所有16个主机线程中的所有数据读入一个巨型结构,并将其传递给GPU以启动一个内核.但是,复制太多会降低应用程序的速度.

这个问题与使用cuda流来运行许多内核有关

在CUDA中,有许多同步命令cudaStreamSynchronize,CudaDeviceSynchronize,cudaThreadSynchronize以及cudaStreamQuery来检查流是否为空.

我注意到在使用分析器时,这些同步命令会给程序带来很大的延迟.我想知道是否有人知道减少这种延迟的任何方法,当然除了使用尽可能少的同步命令.

还有任何数字来判断最有效的同步方法.考虑在应用程序中使用的3个流,如果我使用2个cudaStreamSyncs或者只使用一个cudaDeviceSync,那么其中两个需要完成才能启动第四个流,这将导致更少的损失？

我有一个 CUDA 流，有人递给我 - 一个cudaStream_t值。CUDA运行时 API似乎没有指示我如何获取与该流关联的设备的索引。

现在，我知道这cudaStream_t只是一个指向驱动程序级流结构的指针，但我不太愿意深入研究驱动程序。有一个惯用的方法来做到这一点吗？或者有什么好的理由不想这样做？

编辑：这个问题的另一个方面是流是否确实与设备相关联，在这种方式中，CUDA 驱动程序本身可以根据给定的指向结构确定设备的标识。

使用CUDA Dynamic Parallelism时,我们遇到了性能问题.目前,CDP的表现比传统方法慢至少3倍.我们制作了最简单的可重现代码来显示这个问题,即将数组的所有元素的值增加+1.即

a[0,0,0,0,0,0,0,.....,0] --> kernel +1 --> a[1,1,1,1,1,1,1,1,1]

这个简单示例的目的只是为了查看CDP是否可以像其他CDP一样执行,或者是否存在严重的开销.

代码在这里:

#include <stdio.h>
#include <cuda.h>
#define BLOCKSIZE 512

__global__ void kernel_parent(int *a, int n, int N);
__global__ void kernel_simple(int *a, int n, int N, int offset);


// N is the total array size
// n is the worksize for a kernel (one third of N)
__global__ void kernel_parent(int *a, int n, int N){
    cudaStream_t s1, s2;
    cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
    cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);

    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if(tid == …

从K20开始,不同的流变为完全并发(用于在边缘上并发).

但是我的程序需要旧的方式.或者我需要做很多同步来解决依赖问题.

是否可以将流管理切换为旧方式？

我正在寻找一种方法来摆脱闲置代码中的主机线程中的繁忙等待（不复制该代码，它仅表示我的问题，它具有许多基本错误）：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     while (true) {
         if (cudaStreamQuery(streams[sid])) == cudaSuccess) { //BUSY WAITING !!!!
             cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
             kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
             break;
         }
         sid = ++sid % S_N;
     }

}

有没有一种方法可以使主机线程空闲并以某种方式等待某个流完成，然后准备并运行另一个流？

编辑：我在代码中添加了while（true），以强调忙等待。现在，我执行所有流，并检查其中哪个流完成以运行另一个新流。cudaStreamSynchronize等待特定的流完成，但是我想等待首先完成工作的任何流。

EDIT2：我摆脱了以休闲方式的繁忙等待：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
} …

我想使用流来并行化在单独的设备数据阵列上工作的内核的执行.数据在设备上分配并从先前的内核填充.

我写了以下程序,显示到目前为止我无法达到目标.实际上,两个非默认流上的内核在它们各自的流中顺序执行.

在具有最新Debian linux版本的2台Intel机器上观察到相同的行为.其中一款配备了CUDA 4.2的Tesla C2075,另一款配备了带有CUDA 5.0的Geforce 460GT.Visual Profiler显示4.2和5.0 CUDA版本中的顺序执行.

这是代码:

#include <iostream>
#include <stdio.h>
#include <ctime>

#include <curand.h>

using namespace std;

// compile and run this way:
// nvcc cuStreamsBasics.cu  -arch=sm_20  -o testCuStream   -lcuda -lcufft -lcurand 
// testCuStream  1024 512 512


/* -------------------------------------------------------------------------- */
//  "useful" macros
/* -------------------------------------------------------------------------- */


#define MSG_ASSERT( CONDITION, MSG )                    \
  if (! (CONDITION))                            \
    {                                   \
    std::cerr << std::endl << "Dynamic assertion `" #CONDITION "` failed in " << __FILE__ \
          << " …

在 CUDA 中，driver_types.h我们有：

typedef __device_builtin__ struct CUstream_st *cudaStream_t;

我们在很多地方都有cuda_runtime.h默认初始化的流参数。例如：

template<class T>
    static __inline__ __host__ cudaError_t cudaLaunchKernel(
    const T *func,
    dim3 gridDim,
    dim3 blockDim,
    void **args,
    size_t sharedMem = 0,
    cudaStream_t stream = 0
)

假设默认流有多安全(cudaStream) nullptr？

从官方的NVIDIA Multi-Process Server 文档看，我不清楚它如何与CUDA流交互。

这是一个例子：

应用0：将内核发布到逻辑流0；

App 1：将内核发布到（自己的）逻辑流0。

在这种情况下，

1）MPS是否/如何“劫持”这些CUDA调用？对于每个应用程序，它是否完全了解使用了哪些流以及哪些流中包含哪些内核？

2）MPS是否创建自己的2个流，并将相应的内核放入正确的流？还是MPS可能通过流以外的机制启用内核并发？

如果有帮助，我对MPS在Volta上的工作方式很感兴趣，但是也希望了解有关旧体系结构的信息。