CUDA By Example笔记--Julia集合的并行计算

目录

1--linux报错汇总

1-1 使用 nvcc 命令编译报错

1-2--使用 CMake 编译源码报错

2--源码解读

---

1--linux报错汇总

1-1 使用 nvcc 命令编译报错

        使用 nvcc ./julia_gpu.cu -lglut -lGLU -lGL 运行时，显示 cannot find -lglut 的错误，定位 "gl_helper.h" 时（即下图所示），无法找到 “glut.h”；

         解决方法：在 linux 下安装 glut 的开源版本，即 freeglut ：

sudo apt-get install build-essential freeglut3 freeglut3-dev binutils-gold

        安装后，可以在路径 “/usr/include/GL” 找到头文件 glut.h；

        测试：重新编译 julia_gpu.cu 文件

nvcc ./julia_gpu.cu -lglut -lGLU -lGL

         运行可执行文件 a.out：

./a.out

1-2--使用 CMake 编译源码报错

        使用 CMake 编译提供的源码时，出现以下错误：

undefined reference to `glClearColor';

undefined reference to `glClear';

undefined reference to `glDrawPixels';

undefined reference to `glFlush';

        解决方法：在 CMakeLists.txt 中增加库的链接：

cmake_minimum_required(VERSION 3.14)

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

        测试：运行可执行文件 main：

2--源码解读

#include "./common/book.h"
#include "./common/cpu_bitmap.h"

#define DIM 1000

struct cuComplex {
    float   r;
    float   i;
    __device__ cuComplex( float a, float b ) : r(a), i(b) {}
    __device__ float magnitude2( void ) {
        return r * r + i * i;
    }
    __device__ cuComplex operator*(const cuComplex& a) {
        return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
    }
    __device__ cuComplex operator+(const cuComplex& a) {
        return cuComplex(r+a.r, i+a.i);
    }
};

// 判断点(x, y)是否属于 julia 集
__device__ int julia( int x, int y ) {
    const float scale = 1.5;
    float jx = scale * (float)(DIM/2 - x)/(DIM/2);
    float jy = scale * (float)(DIM/2 - y)/(DIM/2);

    cuComplex c(-0.8, 0.156);
    cuComplex a(jx, jy);

    int i = 0;
    for (i=0; i<200; i++) {
        a = a * a + c;
        if (a.magnitude2() > 1000)
            return 0;
    }

    return 1;
}

__global__ void kernel( unsigned char *ptr ) {
    
    // 在声明二维线程格时，线程格每一维的大小与图像每一维的大小相等
    // 则(0, 0)和(DIM-1, DIM-1)之间每个像素点(x, y)都能获得一个线程块
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x; // 行索引 y 乘以线程块宽度 gridDim.x，再加上列索引 x，得到每一个(x, y)的指针索引

    int juliaValue = julia( x, y ); // 判断点(x, y)是否属于julia集
    ptr[offset*4 + 0] = 255 * juliaValue;
    ptr[offset*4 + 1] = 0;
    ptr[offset*4 + 2] = 0;
    ptr[offset*4 + 3] = 255;
}

int main( void ) {

    CPUBitmap bitmap(DIM, DIM); // 创建位图图像
    unsigned char    *dev_bitmap; // 创建指针，申请GPU内存
    HANDLE_ERROR( cudaMalloc( (void**)&dev_bitmap, bitmap.image_size() ) );

    dim3    grid(DIM,DIM); // 用于指定启动的线程块数量，这里用到的是二维线程格
    kernel<<<grid,1>>>( dev_bitmap ); // 执行定义的核函数，计算julia集

    HANDLE_ERROR( cudaMemcpy( bitmap.get_ptr(), dev_bitmap,
                              bitmap.image_size(),
                              cudaMemcpyDeviceToHost ) ); // 将计算结果从 device 传递到 host
                              
    HANDLE_ERROR( cudaFree( dev_bitmap ) ); // 释放在 GPU 申请的内存
                              
    bitmap.display_and_exit(); // 可视化位图
}

        核心代码分析：

dim3    grid(DIM,DIM); // 用于指定启动的线程块数量，这里用到的是二维线程格

        代码创建了一个二维线程格，其维度大小为 (DIM, DIM)，即共有 DIM*DIM 个线程块，每一个线程块对应一个位图像素点(x, y)，用于计算其是否属于 julia 集；

__global__ void kernel( unsigned char *ptr ) {
    
    // 在声明二维线程格时，线程格每一维的大小与图像每一维的大小相等
    // 则(0, 0)和(DIM-1, DIM-1)之间每个像素点(x, y)都能获得一个线程块
    int x = blockIdx.x;
    int y = blockIdx.y;
    int offset = x + y * gridDim.x; // 行索引 y 乘以线程块宽度 gridDim.x，再加上列索引 x，得到每一个(x, y)的指针索引

    int juliaValue = julia( x, y ); // 判断点(x, y)是否属于julia集
    ptr[offset*4 + 0] = 255 * juliaValue;
    ptr[offset*4 + 1] = 0;
    ptr[offset*4 + 2] = 0;
    ptr[offset*4 + 3] = 255;
}

         使用__global__声明核函数，gridDim 是一个常数用于保存线程格每一维的大小；由于每一个线程块对应一个像素点，则可以通过 x + y * gridDim.x 计算每一个像素点相应的偏移量；

        kernel 函数是这个例子并行计算的核心代码，通过为每一个像素点分配一个线程块，通过并行计算快速判断每一个像素点是否属于 julia 集，无需在CPU中通过循环每次只能判断一个像素点 (x, y)；