从0开始学习《GPU高性能运算之CUDA》——3

6 规约思想和同步概念

扩大点说，并行计算是有一种基本思想的，这个算法能解决很多很常规的问题，而且很实用，比如说累加和累积等——规约思想。对于基础的、重要的，我想有必要系统的学习。

我觉得有必要重新复制下之前写的这篇介绍：

http://www.cnblogs.com/viviman/archive/2012/11/21/2780286.html

并行程序的开发有其不同于单核程序的特殊性，算法是重中之重。根据不同业务设计出不同的并行算法，直接影响到程序的效率。因此，如何设计并行程序的算法，似乎成为并编程的最大难点。观其算法，包括cuda sdk的例子和网上的牛人，给出的一些例子，以矩阵和矢量处理为主，深入点的包括fft和julia等数学公式，再高级一点的算是图形处理方面的例子。学习这些算法的思想，免不了有自己的一点点总结。之前学习过omp编程，结合现在的cuda，我觉得要理解并行编程，首先理解划分和规约这两个概念。也许你的算法学的更加扎实。划分是《算法》里面的一个重要思想，将一个大的问题或任务，分解成小问题小任务，各个击破，最后归并结果；规约是《cuda**》书上介绍的一个入门的重要思想，规约算法(reduction)用来求连加、连乘、最值等，应用广泛。每次循环参加运算的线程减少一半。不管算法的思想如何花样，万变不离其中的一点--将一个大的任务分解成小的任务集合，分解原则是粒度合适尽量小、数据相关性尽量小。如此而已。因为，我们用GPU是为了加速，要加速必须提高执行任务的并行度！明白这个道理，那么我们将绞尽脑汁地去想方设法分析自己手上的任务，分解、分解、分解！这里拿规约来说事情，因为，规约这个东西，似乎可以拿来单做9*9乘法表来熟悉，熟悉了基础的口诀，那么99*99的难题也会迎刃而解。ex：矢量加法，要实现N=64*256长度的矢量的累加和。假设a+b计算一次耗时t。

cpu计算：显然单核的话需要64*256*t。我们容忍不了。

gpu计算：最初的设想，我们如果有个gpu能同时跑N/2个线程，我们这N/2个线程同时跑，那么不是只需要t时间就能将N个数相加编程N/2个数相加了吗？对的。这一轮我们用了t时间；接着的想法，我们不断的递归这个过程，能发现吗？第二轮，我们用N/2/2个线程同时跑，剩下N/2/2个数相加，这一轮我们同样用了t时间；一直这样想下去吧，最后一轮，我们用了1个线程跑，剩下1个数啦，这就是我们的结果！每一轮时间都为t，那么理想情况，我们用了多少轮这样的计算呢？计算次数=log(N)=6*8=48，对的，只用了48轮，也就是说，我们花了48*t的时间！

规约就是这样，很简单，很好用，我们且不管程序后期的优化，单从这个算法分析上来说，从时间复杂度N降到了logN，这在常规算法上，要提高成这样的效率，是不得了的，这是指数级别的效率提高！所以，你会发现，GPU有CPU无法取代的得天独厚的优势——处理单元真心多啊！

规约求和的核函数代码如下：__global__ void RowSum(float* A, float* B){ int bid = blockIdx.x; int tid = threadIdx.x;

__shared__ s_data[128]; //read data to shared memory s_data[tid] = A[bid*128 + tid]; __synctheads(); //sync

for(int i=64; i>0; i/=2){ if(tid<i) s_data[tid] = s_data[tid] + s_data[tid+i] ; __synctheads(); } if(tid==0) B[bid] = s_data[0];}

 

这个例子还让我学到另一个东西——同步！我先不说同步是什么，你听我说个故事：我们调遣了10个小组从南京去日本打仗，我们的约定是，10个组可以自己行动，所有组在第三天在上海机场会合，然后一起去日本。这件事情肯定是需要处理的，不能第1组到了上海就先去日本了，这些先到的组，唯一可以做的事情是——等待！这个先来后到的事情，需要统一管理的时候，必须同步一下，在上海这个地方，大家统一下步调，快的组等等慢的组，然后一起干接下去的旅程。

是不是很好理解，这就是同步在生活中的例子，应该这样说，计算机的所有机制和算法很多都是源于生活！结合起来，理解起来会简单一点。

在CUDA中，我们的同步机制用处大吗？又是如何用的呢？我告诉你，一个正常规模的工程中，一般来说数据都会有先来后到的关系，这一个计算结果可能是提供给另一个线程用的，这种依赖关系存在，会造成同步的应用。

__synctheads()这句话的作用是，这个block中的所有线程都执行到此的时候，都听下来，等所有都执行到这个地方的时候，再往下执行。

7 撬开编程的锁

锁是数据相关性里面肯定要用到的东西，很好，生活中也一样，没锁，家里不安全；GPU中没锁，数据会被“盗”。

对于存在竞争的数据，CUDA提供了原子操作函数——ATOM操作。

先亮出使用的例子：

__global__ void kernelfun()

{

__shared__ int i=0;

atomicAdd(&i, 1);

}

如果没有加互斥机制，则同一个half warp内的线程将对i的操作混淆林乱。

用原子操作函数，可以很简单的编写自己的锁，SKD中有给出的锁结构体如下：

#ifndef __LOCK_H__

#define __LOCK_H__

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include "atomic_functions.h"

struct Lock {

    int *mutex;

    Lock( void ) {

        HANDLE_ERROR( cudaMalloc( (void**)&mutex, sizeof(int) ) );

        HANDLE_ERROR( cudaMemset( mutex, 0, sizeof(int) ) );

    }

    ~Lock( void ) {

        cudaFree( mutex );

    }

    __device__ void lock( void ) {

        while( atomicCAS( mutex, 0, 1 ) != 0 );

    }

    __device__ void unlock( void ) {

        atomicExch( mutex, 0 );

    }

};

#endif

 

8 CUDA软件体系结构

 

9 利用好现有的资源

如果连开方运算都需要自己去编写程序实现，那么我相信程序员这个职业将会缩水，没有人愿意去干这种活。我想，程序员需要学会“偷懒”，现有的资源必须学会高效率的使用。当c++出现了STL库，c++程序员的开发效率可以说倍增，而且程序稳定性更高。

CUDA有提供给我们什么了吗？给了，其实给了很多。

先介绍几个库：CUFFT、CUBLAS、CUDPP。

这里我先不详细学习这些库里到底有哪些函数，但是，大方向是需要了解的，不然找都不知道去哪儿找。CUFFT是傅里叶变换的库，CUBLAS提供了基本的矩阵和向量运算，CUDPP提供了常用的并行排序、搜索等。

CUDA4.0以上，提供了一个类似STL的模板库，初步窥探，只是一个类似vector的模板类型。有map吗？map其实是一个散列表，可以用hashtable去实现这项机制。

SDK里面有很多例子，包括一些通用的基本操作，比如InitCUDA等，都可以固化成函数组件，供新程序的调用。

具体的一些可以固化的东西，我将在以后的学习中归纳总结，丰富自己的CUDA库！