Archive

Archive for the ‘Parallel Computing’ Category

CUDA深入浅出学习小结(1)

July 31st, 2010 No comments

平方和的并行计算
1. 直接分别使用CPU平方和, 一个Thread 一个Block的计算平方然后CPU计算平方和, 在10^7数据面前, CPU表示毫无压力(0cycle), 9600MGS惨不忍睹(多少cycles没记下)
2. 使用256个Threads后, GPU效率大大提升
3. 使用32blocks后, GPU再次提升效率
4. 此时已经有 32*256 = 8192 个线程在跑, 利用CPU计算和的时间被忽略造成不准确, 把计算和运算放在GPU中, 出现小幅性能下降
5. 计算sum可以并行化, 用树状图表示如下:

大意为第一次把1回到0上, 3加到2上, 因为所有线程都是要发起的, 对于判断哪些线程要执行加法很重要, 一个很简单想法是线程ID对offset的2倍取模看是否为0, 代码如下:

?View Code C
1
2
3
4
5
6
7
8
9

__syncthreads();  //make sure all threads are complete
extern __shared__ share[];
int offset=1
while( offset < THREADNUM ) //THREADNUM is the number of threads per block
{
    if ( threadIdx.x%(offset+offset)==0 )
        share[threadIdx.x] += share[threadIdx.x+offset];
    offset += offset
}

然而时间并没有减少, 发现取模运算太费神, 在深入浅出书上如此写

?View Code C
1
2
3
4
5
6
7
8
9
10

    __syncthreads(); 
    int offset = 1, mask = 1;
    while(offset < THREAD_NUM) { 
        if((tid & mask) == 0) {  // tid = threadIdx.x
            shared[tid] += shared[tid + offset]; 
        } 
        offset += offset; 
        mask = offset + mask; 
        __syncthreads(); 
    }

其中mask在里面的古怪作用正好使得結果如此…还是不太理解= =
6. 据说这方法会有问题..他说后面再说, 而采用另一种更简单方法, 大意是前半与后半相加不断递归…

?View Code C
1
2
3
4
5
6
7
8

    offset = THREAD_NUM >> 1; 
    while(offset > 0) { 
        if(tid < offset) {  //tid = threadIdx.x
            shared[tid] += shared[tid + offset]; 
        } 
        offset >>= 1; 
        __syncthreads(); 
    }

现在看求矩阵乘法了, cudaMemcpy 与 cudaMemcpy2D 区别还是没看懂… C语言2维数组与1维数组不是一样的…吗… int a[N*M] == int a[N][M] 那用cudaMemcpy 一下复制 N*M 不就得了…谁知道为何…

Categories: CUDA, Parallel Computing Tags: , ,