如何为CUDA中的每个线程使用寄存器内存？

Question

我正在尝试翻转尺寸较大的数组。（例如 4096x8192）
起初，我尝试使用两个数组作为输入和输出，并且有效！
（我会说输入是原始的，输出是翻转数组）

但我认为如果每个线程都可以容纳输入元素，将会更容易且更高效。那我就只能用一个数组了！

你们能否分享您的知识或介绍任何有助于解决此问题的文档？

谢谢，这是我的代码。


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define ThreadPB 32 // optimal size
dim3 threadsPerBlock(ThreadPB, ThreadPB);

__global__ void initKernel(int *input, int nx, int ny)
{
    int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
    int idx_y = blockDim.y * blockIdx.y + threadIdx.y;
    int idx = idx_y * nx + idx_x;

    if (idx_x < nx && idx_y < ny) {
        input[idx] = idx_y;
    }

}


__global__ void flipKernel(int *output, int *input, int nx, int ny)
{
    int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
    int idx_y = blockDim.y * blockIdx.y + threadIdx.y;
    int idx = idx_y * nx + idx_x;
    
    // is it possible to use only one array?
    if (idx_x < nx && idx_y < ny) { 
        output[(ny - idx_y - 1) * nx + idx_x] = input[idx_y * nx + idx_x];
    }
}

int main()
{
    // time check
    cudaEvent_t start, stop, start_temp, stop_temp;
    cudaEvent_t start_temp2, stop_temp2;
    float elapsedTime, elapsedTime_temp, elapsedTime_temp2;
    cudaEventCreate(&start);        cudaEventCreate(&stop);
    cudaEventCreate(&start_temp);   cudaEventCreate(&stop_temp);
    cudaEventCreate(&start_temp2);  cudaEventCreate(&stop_temp2);

    const int num_x = 4096;
    const int num_y = 8192;
    const int arraySize = num_x * num_y;

    int *orig, *flip;
    orig = (int *)malloc(sizeof(int) * arraySize);
    flip = (int *)malloc(sizeof(int) * arraySize);

    int *dev_orig = 0;
    int *dev_flip = 0;

    cudaMalloc((void**)&dev_orig, arraySize * sizeof(int));
    cudaMalloc((void**)&dev_flip, arraySize * sizeof(int));
    cudaMemcpy(dev_orig, orig, arraySize * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_flip, flip, arraySize * sizeof(int), cudaMemcpyHostToDevice);



    dim3 blocksFlip((num_x + threadsPerBlock.x - 1) / threadsPerBlock.x, (num_y + threadsPerBlock.y - 1) / threadsPerBlock.y);
    initKernel << <blocksFlip, threadsPerBlock >> > (dev_orig, num_x, num_y);
    
    cudaEventRecord(start, 0);

    flipKernel << <blocksFlip, threadsPerBlock >> > (dev_flip, dev_orig, num_x, num_y);

    // time check end
    cudaEventRecord(stop, 0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsedTime, start, stop); printf("flip 1024x2048 처리 시간 = %f ms.\n", elapsedTime);


    cudaMemcpy(orig, dev_orig, arraySize * sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(flip, dev_flip, arraySize * sizeof(int), cudaMemcpyDeviceToHost);

    // check flip works
    printf("FLIP this array { 0, 1, 2, 3, 4 , 5, 6, 7, 8, 9...} \n= { %d, %d, %d, %d, %d, %d, %d, %d, %d, %d...}\n",
        flip[num_x * 0], flip[num_x * 1], flip[num_x * 2], flip[num_x * 3], flip[num_x * 4],
        flip[num_x * 5], flip[num_x * 6], flip[num_x * 7], flip[num_x * 8], flip[num_x * 9]);


    return 0;
}

Answer 1

对于数组中的偶数行，您应该可以这样做：

__global__ void flipKernel(int *input, int nx, int ny)
{
    int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
    int idx_y = blockDim.y * blockIdx.y + threadIdx.y;
    int idx = idx_y * nx + idx_x;
    
    if (idx_x < nx && idx_y < ny/2) { 
        int output_temp = input[(ny - idx_y - 1) * nx + idx_x];
        input[(ny - idx_y - 1) * nx + idx_x] = input[idx_y * nx + idx_x];
        input[idx_y * nx + idx_x] = output_temp;
    }
}

您只需要在 y 中的线程数减半（y 中的行数减半）的情况下启动该内核。每个线程正在更新矩阵中的两个值。

与其考虑“注册”之类的事情或想象 CUDA 是某种奇怪的语言，如果您具有 C 或 C++ 编程能力，我鼓励您考虑如何解决问题被设计为普通的 C 或 C++ 编程挑战。您的直觉通常在 CUDA 中非常有效。

上面套路的核心就是一个swap。您所说的“寄存器”只是 C 或 C++ 中的一个普通局部变量。在 C++ 中有一个 register 关键字，但它在 CUDA 中基本上没有任何作用，这里也不需要。

您可以处理奇数行，只需保持中间行不变，然后交换其余行即可。这只需要对索引计算稍作更改。

如何为CUDA中的每个线程使用寄存器内存？

How to use register memory for each thread in CUDA?

cuda