《GPU高性能编程CUDA实战》第四章简单的线程块并行

▶ 本章介绍了线程块并行，并给出两个例子：长向量加法和绘制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, );

巴特西

《GPU高性能编程CUDA实战》第四章简单的线程块并行

最新文章

热门文章

巴特西

《GPU高性能编程CUDA实战》第四章 简单的线程块并行

最新文章

热门文章

《GPU高性能编程CUDA实战》第四章简单的线程块并行