1--使用线程实现GPU上的矢量求和

        完整代码：

# include 
# define N (33 * 1024)__global__ void add(int* a, int* b, int* c){int tid = threadIdx.x + blockIdx.x * blockDim.x;while (tid < N){c[tid] = a[tid] + b[tid];tid += blockDim.x * gridDim.x;}
}int main(void){int a[N], b[N], c[N];int *dev_a, *dev_b, *dev_c;// 在 GPU 分配内存cudaMalloc( (void**)&dev_a, N * sizeof(int));cudaMalloc( (void**)&dev_b, N * sizeof(int));cudaMalloc( (void**)&dev_c, N * sizeof(int));// 在 cpu 上赋值for(int i = 0; i < N; i++){a[i] = i;b[i] = i * i;}// 将数据复制到GPUcudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);add<<<128, 128>>>(dev_a, dev_b, dev_c);cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);// 验证是否正确执行bool success = true;for(int i = 0; i < N; i++){if((a[i] + b[i]) != c[i]){printf("Error: %d + %d != %d\n", a[i], b[i], c[i]);success = false;}}if(success) printf("We did it! \n");// 释放内存cudaFree(dev_a);cudaFree(dev_b);cudaFree(dev_c);return 0;
}

        核心代码分析：

add<<<128, 128>>>(dev_a, dev_b, dev_c);

        第一个参数 128 表示创建 128 个线程块，第二个参数 128 表示每个线程块包含 128 个线程，即创建了 128 × 128 个并行线程；

int tid = threadIdx.x + blockIdx.x * blockDim.x;

        计算每一个线程在所有线程中对应的索引值，threadIdx.x表示线程的索引，blockIdx.x表示线程块的索引，blockDim.x表示一个线程块的x长度；

while (tid < N){c[tid] = a[tid] + b[tid];tid += blockDim.x * gridDim.x;
}

        由于设置的线程块数目为 128，每个线程块线程的数目也是 128，共有 128 × 128 个并行线程，因此一次只能处理 128 × 128 个数据；

        每次处理完 128 × 128 个数据后，索引值需要增加 128 × 128，即 tid += blockDim.x * gridDim.x；

2--使用线程实现波纹效果

        完整代码：

#include "cuda.h"
#include "./common/book.h"
#include "./common/cpu_anim.h"#define DIM 1024
#define PI 3.1415926535897932f__global__ void kernel( unsigned char *ptr, int ticks ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;// now calculate the value at that positionfloat fx = x - DIM/2;float fy = y - DIM/2;float d = sqrtf( fx * fx + fy * fy );unsigned char grey = (unsigned char)(128.0f + 127.0f *cos(d/10.0f - ticks/7.0f) /(d/10.0f + 1.0f));    ptr[offset*4 + 0] = grey;ptr[offset*4 + 1] = grey;ptr[offset*4 + 2] = grey;ptr[offset*4 + 3] = 255;
}struct DataBlock {unsigned char   *dev_bitmap;CPUAnimBitmap  *bitmap;
};void generate_frame( DataBlock *d, int ticks ) {dim3    blocks(DIM/16, DIM/16);dim3    threads(16, 16);kernel<<>>( d->dev_bitmap, ticks );HANDLE_ERROR( cudaMemcpy( d->bitmap->get_ptr(),d->dev_bitmap,d->bitmap->image_size(),cudaMemcpyDeviceToHost ) );
}void cleanup( DataBlock *d ) {HANDLE_ERROR( cudaFree( d->dev_bitmap ) ); 
}int main( void ) {DataBlock   data;CPUAnimBitmap  bitmap( DIM, DIM, &data );data.bitmap = &bitmap;HANDLE_ERROR( cudaMalloc( (void**)&data.dev_bitmap,bitmap.image_size() ) );bitmap.anim_and_exit( (void (*)(void*,int))generate_frame,(void (*)(void*))cleanup );
}

        CMakeLists.txt:

cmake_minimum_required(VERSION 3.14)project(Test1 LANGUAGES CUDA) # 添加支持CUDA语言
add_executable(main Test0307_2.cu)
target_link_libraries(main -lglut -lGLU -lGL)

        核心代码分析：

dim3    blocks(DIM/16, DIM/16);
dim3    threads(16, 16);
kernel<<>>( d->dev_bitmap, ticks );

        创建二维索引值的线程块和创建二维索引值的线程；

        每一个线程有唯一的索引 (x, y);

        在上述代码中，一副图像共有 DIM × DIM个像素，创建了 DIM/16 × DIM/16 个线程块，每一个线程块包含 16 × 16 个线程，一共有DIM × DIM个并行线程；因此每一个线程对应一个像素，每一个线程处理一个像素；

__global__ void kernel( unsigned char *ptr, int ticks ) {// map from threadIdx/BlockIdx to pixel positionint x = threadIdx.x + blockIdx.x * blockDim.x;int y = threadIdx.y + blockIdx.y * blockDim.y;int offset = x + y * blockDim.x * gridDim.x;// 生成波纹的代码float fx = x - DIM/2;float fy = y - DIM/2;float d = sqrtf( fx * fx + fy * fy );unsigned char grey = (unsigned char)(128.0f + 127.0f *cos(d/10.0f - ticks/7.0f) /(d/10.0f + 1.0f));    ptr[offset*4 + 0] = grey;ptr[offset*4 + 1] = grey;ptr[offset*4 + 2] = grey;ptr[offset*4 + 3] = 255;
}

每一个线程的索引 (x, y) 通过以下代码进行计算：

        int x = threadIdx.x + blockIdx.x * blockDim.x;
        int y = threadIdx.y + blockIdx.y * blockDim.y;

其偏移量通过以下代码进行计算：

        int offset = x + y * blockDim.x * gridDim.x;

未完待续！