CUDA02_04GPU计算线程分配

本主题体会下GPU计算的结构设计对并发编程的天然支持；使用GPU处理并发任务，简洁方便。
本主题主要内容就是怎么计算线程分配
1. GPU C/C++的扩展语法
2. Grid的结构；
3. Block的结构；
最后给出了一个使用GPU处理图像的例子。

序言

GPU采用的是并行处理来提高速度，主要是因为：
- GPU高度并行；
- 多线程；
- 多核处理器；
- 极高的内存带宽。
GPU性能提升证据图
- 每秒的浮点运算次数数据。明显单精度浮点数运算很彪悍。
- 下图来自官方文档，截止时间2020年1月。
  
  浮点数运算次数提升与比较
GPU内存带宽提升证据图
- 是每秒的内存IO的理论峰值。
- 下图来自官方文档。
  
  内存带宽提升图
GPU的设计使用更具的计算单元而不是存储单元。
- 官方说法：GPU将更多的晶体管用于数据处理(ALU:算术逻辑单元Arithmetic Logical Unit)。
- 下图是CPU与GPU的存储与计算晶体管的比较。绿色部分是计算单元。
  
  CPU与GPU性能提升的设计比较

CPU就多核（multicore），GPU是很多核（manycore）
- CPU的多线程编程，在GPU上也是多线程编程，思想是一样的！但是GPU天生就是多线程。 CPU需要我们自己使用线程库处理。

GPU的核动态可分配性
- GPU是围绕流式多处理器（SMs）阵列构建的。
  - 每个SM分配多核块；
  - 每个块可以指定线程数（线程数会根据处理器的情况来充分利用GPU的处理器）
    - 比如：5个线程，一个处理器，处理器调用5次；
    - 比如：5个线程，5个处理器，处理器调用1次；
- 多线程程序被划分成独立执行的线程块，这样多处理器的GPU将比少处理器的GPU在更短的时间内自动执行程序。

SM流式多处理器的示意图

SM流式多处理器分配示意图
其中理论上一个SM对应一个函数，一个程序由多个SM（使用Grid描述）构成运算
- 但是一个程序的SM是可以根据需要自己定义的。

GPU计算的SM模型

GPU的C/C++扩展向量数据类型

基本类型的向量扩展

向量的最大长度为4，一次按照x，y，z，w访问

Type	Alignment
char1, uchar1	1
char2, uchar2	2
char3, uchar3	1
char4, uchar4	4
short1, ushort1	2
short2, ushort2	4
short3, ushort3	2
short4, ushort4	8
int1, uint1	4
int2, uint2	8
int3, uint3	4
int4, uint4	16
long1, ulong1	4 if sizeof(long) is equal to sizeof(int) 8, otherwise
long2, ulong2	8 if sizeof(long) is equal to sizeof(int), 16, otherwise
long3, ulong3	4 if sizeof(long) is equal to sizeof(int), 8, otherwise
long4, ulong4	16
longlong1, ulonglong1	8
longlong2, ulonglong2	16
longlong3, ulonglong3	8
longlong4, ulonglong4	16
float1	4
float2	8
float3	4
float4	16
double1	8
double2	16
double3	8
double4	16

dim3类型

这个类型太过常用，所以实际上是把uint3定义成dim3。
这个类型默认初始化为（1，1，1）

数据变量的初始化

系统提供了这些数据类型的构造函数完成初始化。
- int2 make_int2(int x, int y);
- int2 变量(int x, int y);

扩展数据类型使用例子

#include <stdio.h>
#include <cuda.h>

int main(int argc, const char **argv){
    dim3  pos;
//     pos.x = 10;
//     pos.y = 20;
//     pos.z = 30;
    printf("%d,%d,%d\n", pos.x, pos.y, pos.z);

    float4 v4 = make_float4(1.1f, 2.2f, 3.3f, 4.4f);
    printf("%f,%f,%f,%f\n", v4.x, v4.y, v4.z, v4.w);
    return 0;
}

// nvcc -o main.exe  -Xcompiler /source-charset:utf-8 c01_datatype.cu

提示：
- 因为全部采用UTF-8编码，所以在window下控制终端的编码需要切换为utf-8比较合适。

GPU扩展内置变量

这些扩展变量只能在GPU上运行的时候才能访问。
- 换句话说这些变量只能在核函数内可以直接访问。
这些内置变量实际是系统在工作的时候用来存储GPU调度参数的变量，并提供给用户处理数据与运算。

gridDim变量

类型
- dim3
理解SM（Grid，Block，Thread三者的关系）

Grid，Block，Thread的三级计算资源管理结构：SM管理方式

gridDim的功能
- 描述需要分配分类的block数量（见上图）
  - block的数量 = gridDim.x * gridDim.y * gridDim.z

blockDim变量

类型：
- dim3
功能：
- 用来描述每个block的线程数量，见上图。
- thread的数量 = blockDim.x * blockDim.y * blockDim.z

blockIdx变量

类型：
- uint3
功能：
- 用来记录核函数调用的时候，当前块的编号，应为函数同时被多个线程同时调用，为了避免数据脏，需要编号来维持数据的清晰。

threadIdx变量

类型：
- uint3
功能：
- 用来记录函数调用时的线程编号，如果结合到块编号，我们可以知道函数当前被第几次调用。

使用blockIdx与threadIdx变量，可以对不同线程提供核函数不同的数据运算，使得计算分配到不同的核，避免重复计算。
- 比如一副图像，一个像素的处理可以交给一个线程处理，为了避免一个像素处理两次，我们可以根据blockIdx与threadIdx来确定那个像素被处理，从而避免图像的处理不会重复，而且速度的提升就不是几倍的效率了，那是几百倍的提升。想想都觉得兴奋。

warpSize变量

类型
- int
功能：
- 多处理器创建、管理、调度和执行一组N个并行线程，称为warps。组成一个warp的各个线程在同一个程序地址一起启动，但是它们有自己的指令地址计数器和寄存器状态，因此可以独立地进行分支和执行。
- warpSize就是一个并行线程组的数量。

理解内置变量的代码

下面通过跟踪这些变量，理解调用核函数的线程分配的计算方式；
- 这种方式需要理解CUDA对GPU的线程管理与分配方式。
例子的逻辑介绍：
1. 准备了一个255大小的数据，一个线程处理一个元素（我们的处理逻辑是乘以3）；
2. 255 = 5 * 5 * 3 * 3
3. GPU分配方案：
  - 一个函数一个SM；一个SM就是Grid
  - 每个Grid的block为5 * 5 = 25个块。
  - 每个block的thread为3 * 3 = 9个线程。
4. 通过观察输出的block编号与thread编号可以体会CUDA对线程的管理与分别方式。
理解kernel函数
1. 我们调用函数传递的两个参数会赋值到内置变量：gridDim与blockDim。
2. CUDA在分配线程的时候:
  - 会按照gridDim得到一个索引，赋值给内置变量blockIdx;
  - 按照blockDim得到一个索引，赋值给内置变量threadIdx;
3. 根据gridDim与blockIdx计算block的一维索引就是高中数学中的故事了。
4. 使用block的一维索引，以及blockDim与threadIdx。可以计算出线程在所有线程中的编号。

// 定义核函数，用来实现并行计算
__global__ void record_var(int *p, int *bk_idx, int *th_idx){
    // 使用内置变量得到核函数被调用的编号。这个编号与线程调度，线程调度与gpu核有关。
    // 计算block的编号
    int block_id;
    int thread_id;
    block_id = blockIdx.y * gridDim.x + blockIdx.x;
    // 计算thread的编号  
    thread_id = block_id * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
    // 计算索引
    p[thread_id] *= 5;
    bk_idx[thread_id] = block_id;   // 记录block的id。一共25个block，分成5*5
    th_idx[thread_id] = thread_id;  // 记录thread的id，每个block一共4个线程，分成3*3
}

核函数的调用
- 我们先不关注数据，这个基本上是编程模式中的套路。重点关注其中的线程分配。

int main(int argc, const char **argv){
    // 准备cpu数据 (gridDim:5*5 + blockDim:3*3), 一共100个线程，每个线程计算一个元素
    int  *data_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    int  *bkid_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    int  *thid_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    
    ..... <省略若干代码>
    
    // GPU内存
    int  *data_gpu, *bkid_gpu, *thid_gpu; 
   
    ... <省略若干代码>
    
    // 调用核函数
    dim3 grid(5, 5);
    dim3 block(3, 3);

    record_var<<<grid, block>>>(data_gpu, bkid_gpu, thid_gpu);

    return 0;
}

完整代码如下：
- 因为使用了C++语法，所以要注意C++强类型语言导致的类型检测与要求。
- 这里可以使用new运算符分配空间，大家可以试试。

#include <iostream>
#include <cuda.h>
/*
 * 记录内置变量
 */

// 定义核函数，用来实现并行计算
__global__ void record_var(int *p, int *bk_idx, int *th_idx){
    // 使用内置变量得到核函数被调用的编号。这个编号与线程调度，线程调度与gpu核有关。
    // 计算block的编号
    int block_id;
    int thread_id;
    block_id = blockIdx.y * gridDim.x + blockIdx.x;
    // 计算thread的编号  
    thread_id = block_id * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;
    // 计算索引
    p[thread_id] *= 5;
    bk_idx[thread_id] = block_id;   // 记录block的id。一共25个block，分成5*5
    th_idx[thread_id] = thread_id;  // 记录thread的id，每个block一共4个线程，分成3*3
}

int main(int argc, const char **argv){
    // 准备cpu数据 (gridDim:5*5 + blockDim:3*3), 一共100个线程，每个线程计算一个元素
    int  *data_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    int  *bkid_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    int  *thid_cpu = (int*)malloc(5 * 5 * 3 * 3 * sizeof(int));
    // 初始化为3
    for(int i = 0; i < 5 * 5 * 3 * 3; i++){
        data_cpu[i] = 3;
        bkid_cpu[i] = -1;
        thid_cpu[i] = -1;
    }

    // GPU内存
    int  *data_gpu, *bkid_gpu, *thid_gpu; 
    cudaMalloc((void**)&data_gpu, 5 * 5 * 3 * 3 * sizeof(int));
    cudaMalloc((void**)&bkid_gpu, 5 * 5 * 3 * 3 * sizeof(int));
    cudaMalloc((void**)&thid_gpu, 5 * 5 * 3 * 3 * sizeof(int));
    // 拷贝数据
    cudaMemcpy((void*)data_gpu, (void*)data_cpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy((void*)bkid_gpu, (void*)bkid_cpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy((void*)thid_gpu, (void*)thid_cpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyHostToDevice);

    // 调用核函数
    dim3 grid(5, 5);
    dim3 block(3, 3);

    record_var<<<grid, block>>>(data_gpu, bkid_gpu, thid_gpu);

    // 计算完毕，拷贝数据到host
    cudaMemcpy((void*)data_cpu, (void*)data_gpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy((void*)bkid_cpu, (void*)bkid_gpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy((void*)thid_cpu, (void*)thid_gpu, 5 * 5 * 3 * 3 * sizeof(int), cudaMemcpyDeviceToHost);

    // 输出数据观察结果
    for(int i = 0; i < 5 * 5 * 3 * 3; i++){
        std::cout << data_cpu[i] << ", ";
    }
    std::cout <<std::endl;
    for(int i = 0; i < 5 * 5 * 3 * 3; i++){
        std::cout << bkid_cpu[i] << ", ";
    }
    std::cout <<std::endl;
    for(int i = 0; i < 5 * 5 * 3 * 3; i++){
        std::cout << thid_cpu[i] << ", ";
    }
    std::cout <<std::endl;

    // 养成释放内存的习惯
    cudaFree(data_gpu);
    cudaFree(bkid_gpu);
    cudaFree(thid_gpu);

    free(data_cpu);
    free(bkid_cpu);
    free(thid_cpu);
    return 0;
}

数据的计算结果

15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15, 15,

块索引跟踪的结果

0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 2, 2, 2, 2, 2, 2, 2, 2, 2, 3, 3, 3, 3, 3, 3, 3, 3, 3, 4, 4, 4, 4, 4, 4, 4, 4, 4, 5, 5, 5, 5, 5, 5, 5, 5, 5, 6, 6, 6, 6, 6, 6, 6, 6, 6, 7, 7, 7, 7, 7, 7, 7, 7, 7, 8, 8, 8, 8, 8, 8, 8, 8, 8, 9, 9, 9, 9, 9, 9, 9, 9, 9, 10, 10, 10, 10, 10, 10, 10, 10, 10, 11, 11, 11, 11, 11, 11, 11, 11, 11, 12, 12, 12, 12, 12, 12, 12, 12, 12, 13, 13, 13, 13, 13, 13, 13, 13, 13, 14, 14, 14, 14, 14, 14, 14, 14, 14, 15, 15, 15, 15, 15, 15, 15, 15, 15, 16, 16, 16, 16, 16, 16, 16, 16, 16, 17, 17, 17, 17, 17, 17, 17, 17, 17, 18, 18, 18, 18, 18, 18, 18, 18, 18, 19, 19, 19, 19, 19, 19, 19, 19, 19, 20, 20, 20, 20, 20, 20, 20, 20, 20, 21, 21, 21, 21, 21, 21, 21, 21, 21, 22, 22, 22, 22, 22, 22, 22, 22, 22, 23, 23, 23, 23, 23, 23, 23, 23, 23, 24, 24, 24, 24, 24, 24, 24, 24, 24,

线程索引跟踪的结果


0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123, 124, 125, 126, 127, 128, 129, 130, 131, 132, 133, 134, 135, 136, 137, 138, 139, 140, 141, 142, 143, 144, 145, 146, 147, 148, 149, 150, 151, 152, 153, 154, 155, 156, 157, 158, 159, 160, 161, 162, 163, 164, 165, 166, 167, 168, 169, 170, 171, 172, 173, 174, 175, 176, 177, 178, 179, 180, 181, 182, 183, 184, 185, 186, 187, 188, 189, 190, 191, 192, 193, 194, 195, 196, 197, 198, 199, 200, 201, 202, 203, 204, 205, 206, 207, 208, 209, 210, 211, 212, 213, 214, 215, 216, 217, 218, 219, 220, 221, 222, 223, 224,

结论：
- 看见这些线程，以及GPU的核设计结构，有一种蚂蚁搬家的感觉。

GPU核数与线程数

每个GPU的Kernel是与CPU的核数有关的。

获取CPU的核数

int main(int argc, char const *argv[]){
    // 为了代码的简洁性，省略了异常处理
    int kernel_count;
    cudaDeviceGetAttribute (&kernel_count, cudaDevAttrMultiProcessorCount, 0);
    printf("GPU设备上的核数量：%d\n", kernel_count);
    return 0;
}

获取每个核关联的GPU Core。
- GPU Core与GPU的计算能力版本有关。下面是计算能力版本的对应关系。

    sSMtoCores nGpuArchCoresPerSM[] = {
        {0x30, 192},
        {0x32, 192},
        {0x35, 192},
        {0x37, 192},
        {0x50, 128},
        {0x52, 128},
        {0x53, 128},
        {0x60,  64},
        {0x61, 128},    // 本人机器的兼容版本
        {0x62, 128},
        {0x70,  64},
        {0x72,  64},
        {0x75,  64},
        {-1, -1}
    };

获取GPU core支持数代码：

#include <iostream>
#include <cuda.h>

inline int _ConvertSMVer2CoresDRV(int major, int minor) {
    typedef struct {
        int SM;     // 0xMm (hexidecimal notation), M = SM Major version, and m = SM minor version
        int Cores;
    } sSMtoCores;

    sSMtoCores nGpuArchCoresPerSM[] = {
        {0x30, 192},
        {0x32, 192},
        {0x35, 192},
        {0x37, 192},
        {0x50, 128},
        {0x52, 128},
        {0x53, 128},
        {0x60,  64},
        {0x61, 128},
        {0x62, 128},
        {0x70,  64},
        {0x72,  64},
        {0x75,  64},
        {-1, -1}
    };

    int index = 0;

    while (nGpuArchCoresPerSM[index].SM != -1) {
        if (nGpuArchCoresPerSM[index].SM == ((major << 4) + minor)) {
            return nGpuArchCoresPerSM[index].Cores;
        }

        index++;
    }
    printf("MapSMtoCores for SM %d.%d is undefined.  Default to use %d Cores/SM\n",major, minor, nGpuArchCoresPerSM[index - 1].Cores);
    return nGpuArchCoresPerSM[index - 1].Cores;
}

int main(int argc, char const *argv[]){
    cuInit(0);   
    CUdevice device;
    cuDeviceGet (&device, 0);

    int kernel_count;
    cudaDeviceGetAttribute (&kernel_count, cudaDevAttrMultiProcessorCount, 0);
    std::cout<< "设备上的核数量：" <<  kernel_count << std::endl;

    int major, minor;
    cuDeviceGetAttribute(&major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, device);
    cuDeviceGetAttribute(&minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, device);

    // GPU Core数量
    int  gpu_core = _ConvertSMVer2CoresDRV(major, minor);
    std::cout<< "GPU Core / SM:" << gpu_core << std::endl;
    std::cout<< "GPU Core Total:" << gpu_core * kernel_count << std::endl;

    return 0;
}

// nvcc -o main.exe  -l cuda -Xcompiler /source-charset:utf-8 c03_gpu_core.cu

本机运行结果

C:\01works\02cuda\c04c_extension>main
设备上的核数量：10
GPU Core / SM:128
GPU Core Total:1280

GPU的grid与block的支持的限制
- grid没有限制；
- 但是每个块的线程数是有限制的，最大1024.
  - 因为块的所有线程都应该位于同一个处理器核心上，并且必须共享该核心的有限内存资源。在当前GPU上，一个线程块最多可以包含1024个线程。

GPU计算的应用例子

有了上面的理解，我们可以发现GPU对图像的运算，或者类似矩阵的运算，具备比较好的线程分配实现，从而实现性能提升。
下面使用GPU实现图像的颜色通道交换。

GPU的计算能力的特征

-官方地址
- https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities
计算能力特征截图

计算能力特征
计算能力的技术参数
- 只截图我们关注的部分。
  
  计算能力技术参数
gridDim的要求

gridDim主要定义块数：
1. 维度最多3位；
2. 第1维度最大为MAX_INT = 2,147,483,647 = $2 ^ {31} - 1$ ；
3. 低2，3维最大为65535 = MAX_SHORT_INT = $2 ^{16} - 1$ ;

blockDim的要求

blockDim主要定义块数：
1. 维度最多3位；
2. 第1,2维度最大为 1024
3. 低3维最大为 64

thread线程数 / block限制:
1. 每个块的最大线程数： $\le 1024$

核心代码

1.函数定义

- 注意：其中线程编号的计算是与调用的时候的线程分配结构有关系的。


__global__ void shift_color_channels(uchar4 *data){
    // 计算索引
    int  idx = blockIdx.x * blockDim.x + threadIdx.x;   // 因为是1维，获取x就是索引。
    // 处理像素
    unsigned char red   = data[idx].x;
    unsigned char green = data[idx].y;
    unsigned char blue  = data[idx].z;
    // unsigned char alpha = pixel.w;

    data[idx].x = blue;
    data[idx].y = red;
    data[idx].z = green;
    // alpha不动
}

函数调用：
- 注意因为行超过了1024，所以折半后的数据作为block的线程数。这类没有使用多维结构了！喜欢多维的可以自己设计下。
- 如果超过1024，则CUDA不工作。
  - x，y不能超过1024，z不能超过64，而且 $x \times y \times z$ 也不能超过1024。

    ////////////////////////////////GPU处理调用
    //为了简单，按照图像的行列分配线程，行用来定义块，列用来定义线程, 每个线程处理4个char，并交换像素通道
    dim3 grid(header.height * 2);    // C/C++的除法对整数是整除。
    dim3 block(header.width / 2);    // 注意：每个block做多1024个线程。因为块的所有线程都应该位于同一个处理器核心上


    shift_color_channels<<<grid, block>>>(img_gpu);

    ////////////////////////////////

完整代码

代码

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>
// 结构体定义
#pragma pack(1)
struct img_header{
    // 文件头
    char                  magic[2];                  // 魔法字
    unsigned int          file_size;                 // 文件大小
    unsigned char         reserve1[4];               // 跳4字节
    unsigned int          data_off;                  // 数据区开始位置
    // 信息头
    unsigned char         reserve2[4];               // 跳4字节
    int                   width;                     // 图像宽度
    int                   height;                    // 图像高度
    unsigned char         reserve3[2];               // 跳2字节
    unsigned short int    bit_count;                 // 图像位数1，4，8，16，24，32
    unsigned char         reserve4[24];              // 跳24字节
};


// 偷懒写一个匿名全局类

// 全局数据
struct img_header  header;
uchar4             *img;                             // 使用gpu的扩展类型
uchar4             *img_gpu; 

// 输入/输出文件名
const char *in_filename  = "gpu.bmp";
const char *out_filename = "gpu_out.bmp";

// 打开图像
void read_bmp();                                     // 无参数，采用全局成员
// 保存图像
void save_bmp();
// GPU运算的数据
void move_to_device();  
// 取得数据
void move_to_host();                            
// GPU处理
__global__ void shift_color_channels(uchar4 *data);
// 内存释放
void free_mem();

int main(int argc, const char **argv){
    read_bmp();
    move_to_device();
    ////////////////////////////////GPU处理调用
    //为了简单，按照图像的行列分配线程，行用来定义块，列用来定义线程, 每个线程处理4个char，并交换像素通道
    dim3 grid(header.height * 2);    // C/C++的除法对整数是整除。
    dim3 block(header.width / 2);    // 注意：每个block做多1024个线程。因为块的所有线程都应该位于同一个处理器核心上


    shift_color_channels<<<grid, block>>>(img_gpu);

    ////////////////////////////////
    move_to_host();
    save_bmp();
    free_mem();
    return 0;
}

__global__ void shift_color_channels(uchar4 *data){
    // 计算索引
    int  idx = blockIdx.x * blockDim.x + threadIdx.x;   // 因为是1维，获取x就是索引。
    // 处理像素
    unsigned char red   = data[idx].x;
    unsigned char green = data[idx].y;
    unsigned char blue  = data[idx].z;
    // unsigned char alpha = pixel.w;

    data[idx].x = blue;
    data[idx].y = red;
    data[idx].z = green;
    // alpha不动
}
void move_to_host(){
    cudaMemcpy((void*)img, (void*)img_gpu, header.height * header.width * sizeof(uchar4), cudaMemcpyDeviceToHost);
}
void move_to_device(){
    // 分配GPU内存
    cudaMalloc((void**)&img_gpu, header.height * header.width * sizeof(uchar4));   // 返回指针，则参数就需要二重指针。
    // 拷贝数据
    cudaMemcpy((void*)img_gpu, (void*)img,  header.height * header.width * sizeof(uchar4), cudaMemcpyHostToDevice);

}
void read_bmp(){ 
    /* 读取头，分配内存，读取数据，这里数据采用了一维数组，使用的时候，需要转换处理下。*/
    FILE *file = fopen(in_filename, "rb");
    // 读取头
    size_t n_bytes = fread(&header, 1, 54, file); 
    
    // 计算读取的大大小，并分配空间，并读取。
    header.height = header.height >= 0? header.height : -header.height;
    img = (uchar4 *)malloc(header.height * header.width * sizeof(uchar4));
    n_bytes = fread(img, sizeof(uchar4), header.height * header.width, file);  // 因为是4倍数对齐的，所以可以直接读取

    fclose(file); // 关闭文件
    
}
void save_bmp(){
    /* 使用与读取一样的头信息保存图像 */
    FILE *file = fopen(out_filename, "wb");
    // 写头
    header.height = -header.height;
    size_t n_bytes = fwrite(&header, 1, 54, file);
    header.height = -header.height;
    // 写图像数据
    n_bytes = fwrite(img, sizeof(uchar4), header.height * header.width, file);
    // 关闭文件
    fclose(file);
}
void free_mem(){
    /* 释放Host与Device内存 */
    free(img); // 直接释放（不需要指定大小，malloc系列函数有内部变量管理分配的内存）
    cudaFree(img_gpu);
}

// nvcc -o main.exe  -Xcompiler /source-charset:utf-8 c04_gpu_app.cu

编译
- 上面的代码保存为c04_gpu_app.cu文件名。

main:
    @nvcc -o main.exe  -Xcompiler /source-charset:utf-8 c04_gpu_app.cu
clean:
    @del *.exe *.exp *.lib  gpu_out.bmp 2>Nul

效果截图

GPU运算的应用例子

附录

总结
- 使用GPU处理多任务，比起线程，真是轻松不少，尤其手工调度计算任务，数据脏的同步处理等。

人面猴
序言：七十年代末，一起剥皮案震惊了整个滨河市，随后出现的几起案子，更是在滨河造成了极大的恐慌，老刑警刘岩，带你破解...
沈念sama阅读 204,793评论 6赞 478
死咒
序言：滨河连续发生了三起死亡事件，死亡现场离奇诡异，居然都是意外死亡，警方通过查阅死者的电脑和手机，发现死者居然都...
沈念sama阅读 87,567评论 2赞 381
救了他两次的神仙让他今天三更去死
文/潘晓璐我一进店门，熙熙楼的掌柜王于贵愁眉苦脸地迎上来，“玉大人，你说我怎么就摊上这事。” “怎么了？”我有些...
开封第一讲书人阅读 151,342评论 0赞 338
道士缉凶录：失踪的卖姜人
文/不坏的土叔我叫张陵，是天一观的道长。经常有香客问我，道长，这世上最难降的妖魔是什么？我笑而不...
开封第一讲书人阅读 54,825评论 1赞 277
港岛之恋（遗憾婚礼）
正文为了忘掉前任，我火速办了婚礼，结果婚礼上，老公的妹妹穿的比我还像新娘。我一直安慰自己，他们只是感情好，可当我...
茶点故事阅读 63,814评论 5赞 368
恶毒庶女顶嫁案：这布局不是一般人想出来的
文/花漫我一把揭开白布。她就那样静静地躺着，像睡着了一般。火红的嫁衣衬着肌肤如雪。梳的纹丝不乱的头发上，一...
开封第一讲书人阅读 48,680评论 1赞 281
城市分裂传说
那天，我揣着相机与录音，去河边找鬼。笑死，一个胖子当着我的面吹牛，可吹牛的内容都是我干的。我是一名探鬼主播，决...
沈念sama阅读 38,033评论 3赞 399
双鸳鸯连环套：你想象不到人心有多黑
文/苍兰香墨我猛地睁开眼，长吁一口气：“原来是场噩梦啊……” “哼！你这毒妇竟也来了？” 一声冷哼从身侧响起，我...
开封第一讲书人阅读 36,687评论 0赞 258
万荣杀人案实录
序言：老挝万荣一对情侣失踪，失踪者是张志新（化名）和其女友刘颖，没想到半个月后，有当地人在树林里发现了一具尸体，经...
沈念sama阅读 42,175评论 1赞 300
护林员之死
正文独居荒郊野岭守林人离奇死亡，尸身上长有42处带血的脓包…… 初始之章·张勋以下内容为张勋视角年9月15日...
茶点故事阅读 35,668评论 2赞 321
白月光启示录
正文我和宋清朗相恋三年，在试婚纱的时候发现自己被绿了。大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
茶点故事阅读 37,775评论 1赞 332
活死人
序言：一个原本活蹦乱跳的男人离奇死亡，死状恐怖，灵堂内的尸体忽然破棺而出，到底是诈尸还是另有隐情，我是刑警宁泽，带...
沈念sama阅读 33,419评论 4赞 321
日本核电站爆炸内幕
正文年R本政府宣布，位于F岛的核电站，受9级特大地震影响，放射性物质发生泄漏。R本人自食恶果不足惜，却给世界环境...
茶点故事阅读 39,020评论 3赞 307
男人毒药：我在死后第九天来索命
文/蒙蒙一、第九天我趴在偏房一处隐蔽的房顶上张望。院中可真热闹，春花似锦、人声如沸。这庄子的主人今日做“春日...
开封第一讲书人阅读 29,978评论 0赞 19
一桩弑父案，背后竟有这般阴谋
文/苍兰香墨我抬头看了看天上的太阳。三九已至，却和暖如春，着一层夹袄步出监牢的瞬间，已是汗流浃背。一阵脚步声响...
开封第一讲书人阅读 31,206评论 1赞 260
情欲美人皮
我被黑心中介骗来泰国打工，没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留，地道东北人。一个月前我还...
沈念sama阅读 45,092评论 2赞 351
代替公主和亲
正文我出身青楼，却偏偏与公主长得像，于是被迫代替她去往敌国和亲。传闻我的和亲对象是个残疾皇子，可洞房花烛夜当晚...
茶点故事阅读 42,510评论 2赞 343

CUDA02_04GPU计算线程分配

序言

GPU计算的SM模型

GPU的C/C++扩展向量数据类型

基本类型的向量扩展

dim3类型

数据变量的初始化

扩展数据类型使用例子

GPU扩展内置变量

gridDim变量

blockDim变量

blockIdx变量

threadIdx变量

warpSize变量

理解内置变量的代码

GPU核数与线程数

GPU计算的应用例子

GPU的计算能力的特征

核心代码

完整代码

附录

推荐阅读更多精彩内容