▶ 本章介绍了线程块并行，并给出两个例子：长向量加法和绘制julia集。
● 长向量加法，中规中矩的GPU加法，包含申请内存和显存，赋值，显存传入，计算，显存传出，处理结果，清理内存和显存。用到了 tid += gridDim.x; 使得线程块可以读取多个下标，计算长于线程块数量的向量（例子中向量长度为32768，线程块数量为1024）
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h" #define N (32 * 1024) __global__ void add(int *a, int *b, int *c)
{
int tid = blockIdx.x;
while (tid < N)
{
c[tid] = a[tid] + b[tid];
tid += gridDim.x;
}
return;
} int main(void)
{
int *a, *b, *c;
int *dev_a, *dev_b, *dev_c; // 申请内存和显存
a = (int*)malloc(N * sizeof(int));
b = (int*)malloc(N * sizeof(int));
c = (int*)malloc(N * sizeof(int));
cudaMalloc((void**)&dev_a, N * sizeof(int));
cudaMalloc((void**)&dev_b, N * sizeof(int));
cudaMalloc((void**)&dev_c, N * sizeof(int)); // 数组填充
for (int i = ; i < N; i++)
{
a[i] = i;
b[i] = * i;
} // 将内存中的a和b拷贝给显存中的dev_a和dev_b
cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice); // 调用核函数
add <<<, >>>(dev_a, dev_b, dev_c); // 将显存中的dev_c从显存拷贝回内存中的c
cudaMemcpy(c, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost); // 检验结果
bool success = true;
for (int i = ; i<N; i++)
{
if ((a[i] + b[i]) != c[i])
{
printf("Error at i==%d:\n\t%d + %d != %d\n", i, a[i], b[i], c[i]);
success = false;
break;
}
}
if (success)
printf("We did it!\n"); // 释放内存和显存
free(a);
free(b);
free(c);
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c); getchar();
return ;
}
● 绘制Julia集：
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "D:\Code\CUDA\book\common\book.h"
#include "D:\Code\CUDA\book\common\cpu_bitmap.h" #define DIM 1000 // 输出图形的边长
#define USE_CPU false // true使用CPU，false使用GPU struct DataBlock
{
unsigned char *dev_bitmap;
}; struct cuComplex// 自定义复数类型
{
float r;
float i;
__host__ __device__ cuComplex(float a, float b) : r(a), i(b) {}
__host__ __device__ float magnitude2(void)
{
return r * r + i * i;
}
__host__ __device__ cuComplex operator*(const cuComplex& a)
{
return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
}
__host__ __device__ cuComplex operator+(const cuComplex& a)
{
return cuComplex(r + a.r, i + a.i);
}
}; __host__ __device__ int julia(int x, int y)// 判定一个点是否属于julia集
{
const float scale = 1.5;
float jx = scale * (float)(DIM / - x) / (DIM / );
float jy = scale * (float)(DIM / - y) / (DIM / ); cuComplex c(-0.8, 0.056);// 事先规定的偏移值
cuComplex a(jx, jy); int i = ;
for (i = ; i<; i++)
{
a = a * a + c;
if (a.magnitude2() > )
return ;
}
return ;
} void kernelCPU(unsigned char *ptr)// CPU中按循环遍历每个点进行判定
{
for (int y = ; y<DIM; y++)
{
for (int x = ; x<DIM; x++)
{
int offset = x + y * DIM;
int juliaValue = julia(x, y);
ptr[offset * + ] = * juliaValue;
ptr[offset * + ] = ;
ptr[offset * + ] = ;
ptr[offset * + ] = ;
}
}
return;
} __global__ void kernelGPU(unsigned char *ptr)// GPU中每个线程块判定一个点
{
int x = blockIdx.x;
int y = blockIdx.y;
int offset = x + y * gridDim.x; int juliaValue = julia(x, y);
ptr[offset * + ] = ;
ptr[offset * + ] = * juliaValue;
ptr[offset * + ] = ;
ptr[offset * + ] = ;
return;
} int main(void)
{
#if USE_CPU
CPUBitmap bitmap(DIM, DIM);
unsigned char *ptr = bitmap.get_ptr(); kernelCPU(ptr);
#else
DataBlock data;
CPUBitmap bitmap(DIM, DIM, &data);
unsigned char *dev_bitmap;
cudaMalloc((void**)&dev_bitmap, bitmap.image_size());
data.dev_bitmap = dev_bitmap; kernelGPU << < dim3(DIM, DIM), >> > (dev_bitmap); cudaMemcpy(bitmap.get_ptr(), dev_bitmap, bitmap.image_size(), cudaMemcpyDeviceToHost);
cudaFree(dev_bitmap);
#endif bitmap.display_and_exit();
printf("\n\tfinished!");
getchar();
return ;
}
▶ 在定义结构体的时候定义结构的运算 __host__ __device__ cuComplex(float a, float b) : r(a), i(b) {} ，表明初始化结构实例时可以传入两个浮点参数 a 和 b，并将它们分别当做该实例的两个分量。类似的情形在结构dim3中也有体现。
//在vector_types.h中预定义的结构dim3
struct __device_builtin__ dim3
{
unsigned int x, y, z;
#if defined(__cplusplus)
__host__ __device__ dim3(unsigned int vx = , unsigned int vy = , unsigned int vz = ) : x(vx), y(vy), z(vz) {}
__host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
__host__ __device__ operator uint3(void) { uint3 t; t.x = x; t.y = y; t.z = z; return t; }
#endif /* __cplusplus */
}; //然后在CUDA的代码中可以使用
dim3 threads = dim3(, );
dim3 blocks = dim3(n / threads.x, );