c++ - CUDA - 在多个同步内核上处理单个像素缓冲区数据(数组)，这可能吗？

Question

目前我有一个像素缓冲区，我通过一次内核调用处理其中的数据:

dim3 threadsPerBlock(32, 32)
dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / threadsPerBlock.y);
kernel<<<blocks, threadsPerBlock>>>();

像素缓冲区包含尺寸为 screenWidth x screenHeight 的窗口中的所有像素。

我的想法是将窗口分成 2 或 4 个部分，同时处理像素数据。

这可以做到吗？如果可以，怎么做？

我对流的了解很少，但据我了解，两个流不能处理单个数据(例如我的 pixelBuffer)，或者我错了吗？

编辑:我的显卡具有计算能力 3.0

编辑 2:我使用 SDL 绘图，我有一个 GPU，我使用用户定义的数据数组:

ma​​in.cu

 Color vfb_linear[VFB_MAX_SIZE * VFB_MAX_SIZE]; // array on the Host
 Color vfb[VFB_MAX_SIZE][VFB_MAX_SIZE] // 2D array used for SDL
 extern "C" void callKernels(Color* dev_vfb);

int main()
{
    Color* dev_vfb; // pixel array used on the GPU
    // allocate memory for dev_vfb on the GPU
    cudaMalloc((void**)&dev_vfb, sizeof(Color) * RES_X * RES_Y);
    // memcpy HostToDevice
    cudaMemcpy(dev_vfb, vfb_linear, sizeof(Color) * RES_X * RES_Y, cudaMemcpyHostToDevice);

    callKernels(dev_vfb); // wrapper function that calls the kernels

    // memcpy DeviceToHost
    cudaMemcpy(vfb_linear, dev_vfb, sizeof(Color) * RES_X * RES_Y, cudaMemcpyDeviceToHost);

    // convert vfb_linear into 2D array so it can be handled by SDL
    convertDeviceToHostBuffer();    

    display(vfb); // render pixels on screen with SDL

}

cudaRenderer.cu

__global__ void kernel(Color* dev_vfb)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    if (offset < RES_X * RES_Y)
    {
        dev_vfb[offset] = getColorForPixel();
    }
}

extern "C" callKernels(Color* dev_vfb)
{
    dim3 threadsPerBlock(32, 32)
    dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / threadsPerBlock.y);
    kernel<<<blocks, threadsPerBlock>>>(dev_vfb);
}

显示内容(vfb):

void display(Color vfb[VFB_MAX_SIZE][VFB_MAX_SIZE])
{
    // screen is pointer to SDL_Surface
    int rs = screen->format->Rshift;
    int gs = screen->format->Gshift;
    int bs = screen->format->Bshift;

    for (int y = 0; y < screen->h; ++y)
    {
        Uint32* row = (Uint32*) ((Uint8*) screen->pixels + y * screen->pitch);
        for (int x = 0; x < screen->w; ++x)
            row[x] = vfb[y][x].toRGB32(rs, gs, bs);
    }
    SDL_Flip(screen);
}

这是我在项目中所做的一个简单示例。它是一个光线追踪器，也许 SDL 是与 CUDA 互操作的最差选择，但我不知道我是否有时间更改它。

Answer 1

没有什么可以阻止两个流处理一台设备的全局内存中的同一 block 数据。

正如我在评论中所说，我认为这不是让事情运行得更快的明智方法。但是，对您的代码的修改将是这样的(在浏览器中编码，未经测试):

__global__ void kernel(Color* dev_vfb, int slices)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    if (offset < (RES_X * RES_Y/slices)
    {
        dev_vfb[offset] = getColorForPixel();
    }
}

extern "C" callKernels(Color* dev_vfb)
{
    int num_slices=2;
    cudaStream_t streams[num_slices];
    for (int i = 0; i < num_slices; i++)
      cudaStreamCreate(&(streams[i]));
    dim3 threadsPerBlock(32, 32)
    dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / (num_slices*threadsPerBlock.y));
    for (int i = 0; i < num_slices; i++){
      int off = i * (screenWidth*screenHeight/num_slices);
      kernel<<<blocks, threadsPerBlock, 0, streams[i]>>>(dev_vfb+off, num_slices); }
}