GPU编程

my讲的GPU编程

大规模并行计算

Gpu讲稿

 

 

1.CPU、gpu区别

CPU：中央处理器

GPU（Graphic Processing Unit）：图形处理器

 

 

强大的算术运算单元

大的缓存（数据）

复杂的逻辑控制单元（分支预测降低延时）

 

 

 

基于大的吞吐量设计

很多的ALU、很少的cache、

很多线程需要访问同一个相同的数据，缓存会合并，再去访问DRAM，获取后cache转发这些数据给相应线程。

Control（控制单元）把多个访问合并成少的访问。

CPU擅长逻辑控制、串行运算。

Gpu擅长大规模并行计算

 

2.如何使用gpu。使用CUDA进行编程。

 

安装前的环境准备和检查 ：

检查是否支持CUDA的GPU，操作系统linux，C编译器

安装CUDA ：

https://developer.nvidia.com/cuda-downloads

 

Runfile安装

http://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html#runfile-nouveau

 

 

Windows：Microsoft Visual Studio2008

 

3.什么情况使用gpu

GPU's performance advantage comes from its high parallelism and high memory bandwidth. To decide if an operation can benefit from GPUs, three factors should be considered:

(1)whether the operation is compute-intensive;

(2)whether the operation accesses the same data multiple times;

(3)whether there are multiple consecutive GPU operations on the data.

If any of the three factors hold, the operation should be considered to be offloaded to GPUs.

什么类型的程序适合在GPU上运行？　　

1）计算密集型的程序。　

2）易于并行的程序。GPU其实是一种SIMD(Single Instruction Multiple Data)架构， 他有成百上千个核，每一个核在同一时间最好能做同样的事情。

 

 

1.Allocating and releasing GPU device memory,

2.Transferring data between the host memory and the GPU device memory,

3.Launching different GPU kernels.

 

4.Gpu编程

CUDA C程序和C程序没什么区别，只是多了一些以cuda开头的一些库函数和特殊声明的函数。

在GPU设备上执行的函数：核函数kernel

Kernel以一个网格grid的形式执行，每个grid由若干个线程块block组成，每个block由若干个线程thread组成。

 

块并行和线程并行：

线程并行：细粒度并行，调度效率高

块并行：粗粒度并行，每次调度都要重新分配资源。（如果资源只有一份，那么所有block都只能排成一队串行执行）

 

线程块与线程块之间无相关，有利于做更粗粒度的并行。

 

 

CUDA架构

架构分为两部分：host端和device端

Host 端是指在 CPU 上执行的部份，而 device 端则是在GPU上执行的部份。Device 端的程序又称为 kernel函数。

通常 host 端程序会将数据准备好后，复制到GPU的内存中，再由GPU执行 device 端程序，完成后再由 host 端程序将结果从GPU的内存中取回。

 

由于 CPU 存取gpu内存时只能透过 PCI Express 接口，因此速度较慢（PCI Express x16 的理论带宽是双向各 4GB/s），因此不能太常进行这类动作，以免降低效率。

 

 

 

每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个thread 则有共享的一份 share memory。此外，所有的 thread（包括不同 block 的 thread）都共享一份 global memory、constant memory、和 texture memory。

 

5.举例

 

先将数据从内存复制到显存，再写好运算的核函数，之后用CUDA调用核函数，完成GPU上的计算，之后将结果复制回内存，释放掉显存。

 

计算立方和：

单线程

 

#include <stdio.h>

#include <stdlib.h>

 

//CUDA RunTime API

#include <cuda_runtime.h>

#define DATA_SIZE 1048576

int data[DATA_SIZE];

 

void GenerateNumbers(int *number, int size)

{//产生大量0-9之间的随机数，返回number[size]

    for (int i = 0; i < size; i++) {

        number[i] = rand() % 10;

    }

}

 

bool InitCUDA()

{//CUDA 初始化

    int count;

    cudaGetDeviceCount(&count);//取得支持Cuda的装置的数目

    if (count == 0) {

        fprintf(stderr, "There is no device.\n");

        return false;

    }

    int i;

    for (i = 0; i < count; i++) {

        cudaDeviceProp prop;

        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {

            if (prop.major >= 1) {

                break;

            }

        }

    }

    if (i == count) {

        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");

        return false;

    }

    cudaSetDevice(i);

    return true;

}

 

__global__ static void sumOfSquares(int *num, int* result)

{// __global__ 函数 (GPU上执行) 计算立方和

    int sum = 0；

    int i;

    for (i = 0; i < DATA_SIZE; i++) {

        sum += num[i] * num[i] * num[i];

    }

    *result = sum;

}

 

int main()

{

    

if (!InitCUDA())

 {//CUDA 初始化

        return 0;

    }

 

    //生成随机数

    GenerateNumbers(data, DATA_SIZE);

 

    /*把数据复制到显卡内存中*/

    int* gpudata, *result;

 

    //cudaMalloc 取得一块device内存 ( 其中result用来存储计算结果 )

    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);

    cudaMalloc((void**)&result, sizeof(int));

 

    //cudaMemcpy 将产生的随机数复制到显卡内存中

    //cudaMemcpyHostToDevice - 从内存复制到显卡内存

    //cudaMemcpyDeviceToHost - 从显卡内存复制到内存

    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

 

// 在CUDA 中执行函数

// 语法：函数名称<<<block数目, thread数目, shared memory大小>>>(参数...);

    sumOfSquares <<<1, 1, 0 >>>(gpudata, result);

 

 

    /*把结果从device复制回host内存*/

 

    int sum;

 

    //cudaMemcpy 将结果从device中复制回host内存

    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);

 

    //Free

    cudaFree(gpudata);

    cudaFree(result);

 

    printf("GPUsum: %d \n", sum);

-------------------------------------------------------------------------------------------------------------

    sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {

        sum += data[i] * data[i] * data[i];

    }

    printf("CPUsum: %d \n", sum);

 

    return 0;

}

 

 

并行改进：

 

主机程序：设备环境初始化，数据传输等必备过程；

设备程序：只负责计算；

 

#define THREAD_NUM 256

 

sumOfSquares <<<1, THREAD_NUM, 0 >>>(gpudata, result);

 

__global__ static void sumOfSquares(int *num, int* result)

{

    const int tid = threadIdx.x;    //表示目前的 thread 是第几个 thread（由 0 开始计算）

    const int size = DATA_SIZE / THREAD_NUM;    //计算每个线程需要完成的量

    int sum = 0;

    int i;

    for (i = tid * size; i < (tid + 1) * size; i++) {

       sum += num[i] * num[i] * num[i];

    }

    result[tid] = sum;

}

 

在我们的例子中，有 256 个 threads，所以同时会有 256 个 sumOfSquares 函数在执行，但每一个的 threadIdx.x 是不一样的，分别会是 0 ~ 255。所以利用这个变量，我们就可以让每一个函数执行时，对整个数据的不同部份计算立方和。

 

 

 

 

 

 

 

 

Int maxThreadsPerBlock    在一个线程块中可以包含的最大线程数量

 

 

 

 

 

 

 

练习题：

热分布计算问题是一种简单的热能传播问题，对于一个矩形区域，分散存在有N*N个温度点 ( i , j )，每个时间步内，温度会向四周进行传播。假设边缘之外隔温，无需考虑至区域外的温度传播。则每时间步各平温度点的温度计算方式为：

其他计算情况以此类推，即相邻元素的温度平均值。

本实验规定初始情况为：边界设置固定初始值，V=max(x,y)，进行迭代计算至整体温度稳定后（即温度变化前后温度变化小于1e^-3），计算大于温度阈值为5.5的点的个数。

 

灰色为区域边界的初始值：

0	1	2	3	4	5	6	7	8	9	10	11
1	0	0	0	0	0	0	0	0	0	0	11
2	0	0	0	0	0	0	0	0	0	0	11
3	0	0	0	0	0	0	0	0	0	0	11
4	0	0	0	0	0	0	0	0	0	0	11
5	0	0	0	0	0	0	0	0	0	0	11
6	0	0	0	0	0	0	0	0	0	0	11
7	0	0	0	0	0	0	0	0	0	0	11
8	0	0	0	0	0	0	0	0	0	0	11
9	0	0	0	0	0	0	0	0	0	0	11
10	0	0	0	0	0	0	0	0	0	0	11
11	11	11	11	11	11	11	11	11	11	11	11