欢迎您访问 最编程 本站为您分享编程语言代码，编程技术文章！

热门搜索/Hot Search

您现在的位置是： 首页

减少优化（待更新）

最编程 2024-10-14 10:05:19
...

深入浅出GPU优化系列：reduce优化 - 知乎 (zhihu.com)https://zhuanlan.zhihu.com/p/426978026

reduce baseline

#include <cuda_runtime.h>
#include <iostream>
 
__global__ void reduceSumKernel(float *src, float *dest, int n) {
    extern __shared__ float lSum[];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_loc = threadIdx.x;
 
    // 线程内归约
    lSum[idx_loc] = (idx < n) ? src[idx] : 0;
    __syncthreads();
 
    // 线程块内归约
    for (int i = 1; i < blockDim.x; i = 2 * i) {
        if (idx_loc % (2 * i) == 0) {
            lSum[idx_loc] += lSum[idx_loc + i];
        }
        __syncthreads();
    }
 
    // 线程块间归约
    if (idx_loc == 0) {
        atomicAdd(dest, lSum[0]);
    }
}
int main()
{
    const int N = 1024 * 1024; // 数据大小
    const int blockSize = 256; // 线程块大小
    const int numBlocks = (N + blockSize - 1) / blockSize; // 线程块数量
 
    // 主机端数据
    float *src;
    float *dest;
    src = new float[N];
    dest = new float[1];
 
    // 初始化数据
    for (int i = 0; i < N; i++)
    {
        src[i] = 1.0f;
    }
    dest[0] = 0.0f;
 
    // 设备端内存分配
    float *d_src;
    float *d_dest;
    cudaMalloc(&d_src, N * sizeof(float));
    cudaMalloc(&d_dest, sizeof(float));
 
    // 数据传输到设备
    cudaMemcpy(d_src, src, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dest, dest, sizeof(float), cudaMemcpyHostToDevice);
 
    // 调用内核
    reduceSumKernel<<<numBlocks, blockSize, blockSize * sizeof(float)>>>(d_src, d_dest, N);
    cudaDeviceSynchronize();
 
    // 数据从设备传输回主机
    cudaMemcpy(dest, d_dest, sizeof(float), cudaMemcpyDeviceToHost);
 
    // 输出结果
    std::cout << "Sum: " << dest[0] << std::endl;
 
    // 验证结果
    float expectedSum = static_cast<float>(N);
    if (dest[0] == expectedSum)
    {
        std::cout << "Result is correct." << std::endl;
    }
    else
    {
        std::cout << "Result is incorrect." << std::endl;
    }
 
    // 释放内存
    delete[] src;
    delete[] dest;
    cudaFree(d_src);
    cudaFree(d_dest);
 
    return 0;
}

解决warp divergence

#include <cuda_runtime.h>
#include <iostream>
//  现有问题
// 目前reduce0存在的最大问题就是warp divergent的问题。
// 对于一个block而言，它所有的thread都是执行同一条指令。如果存在if-else这样的分支情况的话，thread会执行所有的分支。
// 只是不满足条件的分支，所产生的结果不会记录下来。
// 解决方式
// 解决的方式也比较明了，就是尽可能地让所有线程走到同一个分支里面。
// 代码示意如下：
__global__ void reduce1(float *src, float *dest, int n) {
    extern __shared__ float lSum[];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_loc = threadIdx.x;
 
    // 线程内归约
    lSum[idx_loc] = (idx < n) ? src[idx] : 0;
    __syncthreads();
//  这段代码通过计算索引index = 2 * i * idx_loc并检查index < blockDim.x来避免warp divergence。
//  它之所以能避免发散，是因为所有线程执行的是相同的计算逻辑，即基于固定的数学运算来决定是否进行累加。
//  尽管有if (index < blockDim.x)条件判断，但这个判断是基于线程的固定索引和循环变量的数学运算，
//  而不是依赖于数据内容或运行时的条件，这意味着在任何给定的迭代中，对于同一个warp内的所有线程，这个条件要么都满足，要么都不满足。
//  因此，不会导致线程束内的线程执行不同的指令路径。
// 简而言之，由于所有线程根据相同的规则（基于它们的线程ID和固定算法）决定是否执行累加操作，这确保了在每个循环迭代中，warp内的线程要么一起执行累加，要么一起跳过该操作，从而避免了warp发散。
// 这种设计利用了SIMD执行模型的优势，确保了在关键的归约步骤中保持线程间的同步执行，提高了执行效率。
    // 线程块内归约
    for (int i = 1; i < blockDim.x; i = 2 * i) {
        // 计算当前线程需要累加的共享内存位置
        int index = 2 * i * idx_loc;
        // 检查索引是否在有效范围内，避免银行冲突
        if (index < blockDim.x) {
            // 执行两两累加
            lSum[index] += lSum[index + i];
        }
        // 确保所有线程完成其操作，避免数据竞争
        __syncthreads();
    }
 
    // 线程块间归约
    if (idx_loc == 0) {
        atomicAdd(dest, lSum[0]);
    }
}
int main()
{
    const int N = 1024 * 1024; // 数据大小
    const int blockSize = 256; // 线程块大小
    const int numBlocks = (N + blockSize - 1) / blockSize; // 线程块数量
 
    // 主机端数据
    float *src;
    float *dest;
    src = new float[N];
    dest = new float[1];
 
    // 初始化数据
    for (int i = 0; i < N; i++)
    {
        src[i] = 1.0f;
    }
    dest[0] = 0.0f;
 
    // 设备端内存分配
    float *d_src;
    float *d_dest;
    cudaMalloc(&d_src, N * sizeof(float));
    cudaMalloc(&d_dest, sizeof(float));
 
    // 数据传输到设备
    cudaMemcpy(d_src, src, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dest, dest, sizeof(float), cudaMemcpyHostToDevice);
 
    // 调用内核
    reduce1<<<numBlocks, blockSize, blockSize * sizeof(float)>>>(d_src, d_dest, N);
    cudaDeviceSynchronize();
 
    // 数据从设备传输回主机
    cudaMemcpy(dest, d_dest, sizeof(float), cudaMemcpyDeviceToHost);
 
    // 输出结果
    std::cout << "Sum: " << dest[0] << std::endl;
 
    // 验证结果
    float expectedSum = static_cast<float>(N);
    if (dest[0] == expectedSum)
    {
        std::cout << "Result is correct." << std::endl;
    }
    else
    {
        std::cout << "Result is incorrect." << std::endl;
    }
 
    // 释放内存
    delete[] src;
    delete[] dest;
    cudaFree(d_src);
    cudaFree(d_dest);
 
    return 0;
}

优化技巧2：解决bank冲突

#include <cuda_runtime.h>
#include <iostream>
// 现有问题
// reduce1的最大问题是bank冲突。
// 我们把目光聚焦在这个for循环中。并且只聚焦在0号warp。
// 在第一次迭代中，0号线程需要去load shared memory的0号地址以及1号地址的数，然后写回到0号地址。
// 而此时，这个warp中的16号线程，需要去load shared memory中的32号地址和33号地址。
// 可以发现，0号地址跟32号地址产生了2路的bank冲突。在第2次迭代中，0号线程需要去load shared memory中的0号地址和2号地址。
// 这个warp中的8号线程需要load shared memory中的32号地址以及34号地址，16号线程需要load shared memory中的64号地址和68号地址，24号线程需要load shared memory中的96号地址和100号地址。
// 又因为0、32、64、96号地址对应着同一个bank，所以此时产生了4路的bank冲突。
// 现在，可以继续算下去，8路bank冲突，16路bank冲突。由于bank冲突，所以reduce1性能受限。
__global__ void reduce2(float *src, float *dest, int n) {
    extern __shared__ float lSum[];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_loc = threadIdx.x;
 
    // 线程内归约
    lSum[idx_loc] = (idx < n) ? src[idx] : 0;
    __syncthreads();
    // 把目光继续看到这个for循环中，并且只分析0号warp。
    // 0号线程需要load shared memory的0号元素以及128号元素。
    // 1号线程需要load shared memory中的1号元素和129号元素。
    // 这一轮迭代中，在读取第一个数时，warp中的32个线程刚好load 一行shared memory数据。
    // 再分析第2轮迭代，0号线程load 0号元素和64号元素，1号线程load 1号元素和65号元素。
    // 咦，也是这样，每次load shared memory的一行。
    // 再来分析第3轮迭代，0号线程load 0号元素和32号元素，接下来不写了，总之，一个warp load shared memory的一行。
    // 没有bank冲突。到了4轮迭代，0号线程load 0号元素和16号元素。那16号线程呢，16号线程啥也不干，因为s=16，16-31号线程啥也不干，跳过去了。
    // 线程块内归约
    for (int i = blockDim.x/2; i >0; i >>=1) {
         // 如果当前线程 ID 小于 s，则进行累加操作
        if(idx_loc < i){
            lSum[idx_loc]+=lSum[idx_loc+i];
        }
         // 确保所有线程完成当前轮次的操作，避免数据竞争
        __syncthreads();
    }
 
    // 线程块间归约
    if (idx_loc == 0) {
        atomicAdd(dest, lSum[0]);
    }
}
int main()
{
    const int N = 1024 * 1024; // 数据大小
    const int blockSize = 256; // 线程块大小
    const int numBlocks = (N + blockSize - 1) / blockSize; // 线程块数量
 
    // 主机端数据
    float *src;
    float *dest;
    src = new float[N];
    dest = new float[1];
 
    // 初始化数据
    for (int i = 0; i < N; i++)
    {
        src[i] = 1.0f;
    }
    dest[0] = 0.0f;
 
    // 设备端内存分配
    float *d_src;
    float *d_dest;
    cudaMalloc(&d_src, N * sizeof(float));
    cudaMalloc(&d_dest, sizeof(float));
 
    // 数据传输到设备
    cudaMemcpy(d_src, src, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dest, dest, sizeof(float), cudaMemcpyHostToDevice);
 
    // 调用内核
    reduce2<<<numBlocks, blockSize, blockSize * sizeof(float)>>>(d_src, d_dest, N);
    cudaDeviceSynchronize();
 
    // 数据从设备传输回主机
    cudaMemcpy(dest, d_dest, sizeof(float), cudaMemcpyDeviceToHost);
 
    // 输出结果
    std::cout << "Sum: " << dest[0] << std::endl;
 
    // 验证结果
    float expectedSum = static_cast<float>(N);
    if (dest[0] == expectedSum)
    {
        std::cout << "Result is correct." << std::endl;
    }
    else
    {
        std::cout << "Result is incorrect." << std::endl;
    }
 
    // 释放内存
    delete[] src;
    delete[] dest;
    cudaFree(d_src);
    cudaFree(d_dest);
 
    return 0;
}

优化技巧3：解决idle线程 

#include <cuda_runtime.h>
#include <iostream>

__global__ void reduce3(float *src, float *dest, int n) {
    extern __shared__ float lSum[];
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_loc = threadIdx.x;
    int s = blockIdx.x * (blockDim.x * 2) + threadIdx.x;
    // 添加索引范围检查
    if (s < n && s + blockDim.x < n) {
        lSum[idx_loc] = (idx < n)? src[s] + src[s + blockDim.x] : 0;
    } else {
        lSum[idx_loc] = 0;
    }
    __syncthreads();

    for (int i = blockDim.x / 2; i > 0; i >>= 1) {
        if (idx_loc < i) {
            lSum[idx_loc] += lSum[idx_loc + i];
        }
        __syncthreads();
    }

    if (idx_loc == 0) {
        atomicAdd(dest, lSum[0]);
    }
}

int main() {
    const int N = 1024 * 1024;
    const int blockSize = 256;
    // 修正线程块数量计算
    const int numBlocks = (N + blockSize - 1) / blockSize;

    float *src;
    float *dest;
    src = new float[N];
    dest = new float[1];

    for (int i = 0; i < N; i++) {
        src[i] = 1.0f;
    }
    dest[0] = 0.0f;

    float *d_src;
    float *d_dest;
    cudaMalloc(&d_src, N * sizeof(float));
    cudaMalloc(&d_dest, sizeof(float));

    cudaMemcpy(d_src, src, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_dest, dest, sizeof(float), cudaMemcpyHostToDevice);

    reduce3<<<numBlocks, blockSize, blockSize * sizeof(float)>>>(d_src, d_dest, N);
    cudaDeviceSynchronize();

    cudaMemcpy(dest, d_dest, sizeof(float), cudaMemcpyDeviceToHost);

    std::cout << "Sum: " << dest[0] << std::endl;

    // 定义一个小的误差阈值
    const float epsilon = 0.001f;
    float expectedSum = static_cast<float>(N);
    if (std::abs(dest[0] - expectedSum) < epsilon) {
        std::cout << "Result is correct." << std::endl;
    } else {
        std::cout << "Result is incorrect." << std::endl;
    }

    delete[] src;
    delete[] dest;
    cudaFree(d_src);
    cudaFree(d_dest);

    return 0;
}

上一篇： 入门级网络安全工程师的学习内容

下一篇： Vue-admin-box 后端管理框架 - 项目概览

推荐阅读