简单的 CUDA 内核优化答案

【问题标题】：Simple CUDA Kernel Optimization简单的 CUDA 内核优化
【发布时间】：2023-04-09 23:11:02
【问题描述】：

在加速应用程序的过程中，我有一个非常简单的内核，它执行如下所示的类型转换：

__global__ void UChar2FloatKernel(float *out, unsigned char *in, int nElem){
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(i<nElem)
        out[i] = (float) in[i];
}

全局内存访问是合并的，据我了解，使用共享内存也没有好处，因为没有多次读取同一内存。有没有人知道是否可以执行任何优化来加速这个内核。输入和输出数据已经在设备上，因此不需要主机到设备的内存复制。

【问题讨论】：

你是对的，使用共享内存不会给你带来任何麻烦，因为你仍然需要从全局内存中加载一次来自in的元素，然后将它一次写回@987654324 @。如果单个线程计算多个元素，也许您可以获得优势。但是你必须尝试一下。如果您还没有这样做，您可以阅读"Best Practises Guide"。也许有一些新的提示给你。
添加到 hubs 评论，尝试使用 float4、uchar4 等矢量数据类型，甚至让我们知道它在性能上是否有意义。
您无法进一步改进它。做这么简单的事情的内核是一种浪费。如果您正在转换的数据将用作另一个内核的输入，则在该内核执行转换。这将为内核增加一个小的额外计算，但会被读取字符而不是浮点数的 I/O 增益所隐藏。
谢谢。我会看一下指南，看看我是否遗漏了什么。 @DarkZeros 我正在做这个额外的步骤以保持移植版本与原始版本相似，但正如你所说的那样，这是一种浪费，有效的解决方案是在生产结束或消费之前进行。

标签： cuda gpu

【解决方案1】：

您可以通过const __restrict__ 限定符来装饰输入数组，该限定符通知编译器数据是只读的并且没有任何其他指针的别名。这样，编译器将检测到访问是统一的，并可以通过使用其中一个只读缓存（常量缓存或在计算能力 >=3.5 上称为纹理缓存的只读数据缓存）对其进行优化。

您还可以通过__restrict__ 限定符修饰输出数组，以建议编译器进行其他优化。

最后，DarkZeros 的推荐值得借鉴。

【讨论】：

【解决方案2】：

您最好编写代码的矢量化版本，立即将 float4 写入 out。如果 nElem 恰好是 4-multiple 的边界，这应该非常简单，否则，您可能需要注意残差。

【讨论】：

【解决方案3】：

您可以对类似代码执行的最大优化是使用常驻线程并增加每个线程执行的事务数。虽然 CUDA 块调度模型非常轻量级，但它不是免费的，并且启动大量包含仅执行单个内存加载和单个内存存储的线程的块将产生大量的块调度开销。因此，只需启动尽可能多的块，“填满”GPU 的所有 SM，并让每个线程做更多的工作。

第二个明显的优化是切换到 128 字节的内存事务处理负载，这应该会给您带来切实的带宽利用率增益。在 Fermi 或 Kepler GPU 上，这不会像在第一代和第二代硬件上那样大幅提升性能。

将其完全放入一个简单的基准测试中：

__global__ 
void UChar2FloatKernel(float *out, unsigned char *in, int nElem)
{
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(i<nElem)
        out[i] = (float) in[i];
}

__global__
void UChar2FloatKernel2(float  *out, 
                const unsigned char *in, 
            int nElem)
{
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;    
    for(; i<nElem; i+=gridDim.x*blockDim.x) {
        out[i] = (float) in[i];
    }
}

__global__
void UChar2FloatKernel3(float4  *out, 
                const uchar4 *in, 
            int nElem)
{
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;    
    for(; i<nElem; i+=gridDim.x*blockDim.x) {
        uchar4 ival = in[i]; // 32 bit load
        float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w);
        out[i] = oval; // 128 bit store
    }
}

int main(void)
{

    const int n = 2 << 20;
    unsigned char *a = new unsigned char[n];

    for(int i=0; i<n; i++) {
        a[i] = i%255;
    }

    unsigned char *a_;
    cudaMalloc((void **)&a_, sizeof(unsigned char) * size_t(n));
    float *b_;
    cudaMalloc((void **)&b_, sizeof(float) * size_t(n));
    cudaMemset(b_, 0, sizeof(float) * size_t(n)); // warmup

    for(int i=0; i<5; i++)
    {
        dim3 blocksize(512);
        dim3 griddize(n/512);
        UChar2FloatKernel<<<griddize, blocksize>>>(b_, a_, n);
    }

    for(int i=0; i<5; i++)
    {
        dim3 blocksize(512);
        dim3 griddize(8); // 4 blocks per SM
        UChar2FloatKernel2<<<griddize, blocksize>>>(b_, a_, n);
    }

    for(int i=0; i<5; i++)
    {
        dim3 blocksize(512);
        dim3 griddize(8); // 4 blocks per SM
        UChar2FloatKernel3<<<griddize, blocksize>>>((float4*)b_, (uchar4*)a_, n/4);
    }
    cudaDeviceReset();
    return 0;
}

在一个小的费米设备上给我这个：

>nvcc -m32 -Xptxas="-v" -arch=sm_21 cast.cu
cast.cu
tmpxft_000014c4_00000000-5_cast.cudafe1.gpu
tmpxft_000014c4_00000000-10_cast.cudafe2.gpu
cast.cu
ptxas : info : 0 bytes gmem
ptxas : info : Compiling entry function '_Z18UChar2FloatKernel2PfPKhi' for 'sm_2
1'
ptxas : info : Function properties for _Z18UChar2FloatKernel2PfPKhi
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas : info : Used 5 registers, 44 bytes cmem[0]
ptxas : info : Compiling entry function '_Z18UChar2FloatKernel3P6float4PK6uchar4
i' for 'sm_21'
ptxas : info : Function properties for _Z18UChar2FloatKernel3P6float4PK6uchar4i
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas : info : Used 8 registers, 44 bytes cmem[0]
ptxas : info : Compiling entry function '_Z17UChar2FloatKernelPfPhi' for 'sm_21'

ptxas : info : Function properties for _Z17UChar2FloatKernelPfPhi
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas : info : Used 3 registers, 44 bytes cmem[0]
tmpxft_000014c4_00000000-5_cast.cudafe1.cpp
tmpxft_000014c4_00000000-15_cast.ii

>nvprof a.exe
======== NVPROF is profiling a.exe...
======== Command: a.exe
======== Profiling result:
 Time(%)      Time   Calls       Avg       Min       Max  Name
   40.20    6.61ms       5    1.32ms    1.32ms    1.32ms  UChar2FloatKernel(float*, unsigned char*, int)
   29.43    4.84ms       5  968.32us  966.53us  969.46us  UChar2FloatKernel2(float*, unsigned char const *, int)
   26.35    4.33ms       5  867.00us  866.26us  868.10us  UChar2FloatKernel3(float4*, uchar4 const *, int)
    4.02  661.34us       1  661.34us  661.34us  661.34us  [CUDA memset]

在后两个内核中，与 4096 个块相比，仅使用 8 个块提供了很大的加速，这证实了每个线程多个工作项是在这种内存限制、低指令数的情况下提高性能的最佳方法的想法内核。

【讨论】：

您对上面提到的线程调度开销有任何参考吗？我不太明白的是单个内存读/写如何影响线程调度。而且，线程调度只进行一次不是吗？
@GregKasapidis：我没有看到我在那个答案中提到线程调度开销的地方。
"虽然 CUDA 块调度模型非常轻量级，但它不是免费的，并且启动大量只执行单个内存加载和单个内存存储的线程会产生大量调度开销。 "
我建议再读一遍那句话中的所有单词。我没有提到 thread 调度。我提到了 block 调度。它们是不同的东西。

【解决方案4】：

这是该函数的一个 cpu 版本和 4 个 gpu 内核。 3 个内核来自 @talonmies 答案，我添加了仅使用矢量数据类型的 kernel2。

// cpu version for comparison
void UChar2Float(unsigned char *a, float *b, const int n){
    for(int i=0;i<n;i++)
        b[i] = (float)a[i];
}

__global__ void UChar2FloatKernel1(float *out, const unsigned char *in, int nElem){
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(i<nElem)     out[i] = (float) in[i];
}

__global__ void UChar2FloatKernel2(float4  *out, const uchar4 *in, int nElem){
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(i<nElem) {
        uchar4 ival = in[i]; // 32 bit load
        float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w);
        out[i] = oval; // 128 bit store
    }
}

__global__ void UChar2FloatKernel3(float  *out, const unsigned char *in, int nElem) {
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    for(; i<nElem; i+=gridDim.x*blockDim.x) 
    {
        out[i] = (float) in[i];
    }
}

__global__ void UChar2FloatKernel4(float4  *out, const uchar4 *in, int nElem) {
    unsigned int i = (blockIdx.x * blockDim.x) + threadIdx.x;
    for(; i<nElem; i+=gridDim.x*blockDim.x) 
    {
        uchar4 ival = in[i]; // 32 bit load
        float4 oval = make_float4(ival.x, ival.y, ival.z, ival.w);
        out[i] = oval; // 128 bit store
    }
}

在我的 Geforce GT 640 上，以下是计时结果：

simpleKernel (cpu):         0.101463 seconds.
simpleKernel 1 (gpu):       0.007845 seconds.
simpleKernel 2 (gpu):       0.004914 seconds.
simpleKernel 3 (gpu):       0.005461 seconds.
simpleKernel 4 (gpu):       0.005461 seconds.

所以我们可以看到只使用向量类型的 kernel2 是赢家。我已经对 (32 * 1024 * 768) 个元素进行了这些测试。 nvprof 输出也如下所示：

Time(%)      Time     Calls       Avg       Min       Max  Name
91.68%  442.45ms         4  110.61ms  107.43ms  119.51ms  [CUDA memcpy DtoH]
3.76%  18.125ms         1  18.125ms  18.125ms  18.125ms  [CUDA memcpy HtoD]
1.43%  6.8959ms         1  6.8959ms  6.8959ms  6.8959ms  UChar2FloatKernel1(float*, unsigned char const *, int)
1.10%  5.3315ms         1  5.3315ms  5.3315ms  5.3315ms  UChar2FloatKernel3(float*, unsigned char const *, int)
1.04%  5.0184ms         1  5.0184ms  5.0184ms  5.0184ms  UChar2FloatKernel4(float4*, uchar4 const *, int)
0.99%  4.7816ms         1  4.7816ms  4.7816ms  4.7816ms  UChar2FloatKernel2(float4*, uchar4 const *, int)

【讨论】：