CUDA学习之旅2009.11.10---《深入浅出CUDA》学习2

      今天继续学习这篇文章，日子比以前感觉好点了啊，用校园网就得知道什么是高峰期，我现在就是在网速可以的时候看看张纪中新拍的《倚天屠龙记》，网速慢的时候呢，一边听着轻轻的音乐，一边学习这篇文章，也算潇洒啊。

还是老规矩啊，先提几个今天没有看明白的地方吧。

１.sumOfSquares<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int)>>>(gpudata, result, time)和sumOfSquares<<<BLOCK_NUM, THREAD_NUM, sizeof(int)*THREAD_NUM >>>(gpudata, result, time)有区别吗？？我看其他一些函式里面定义空间大小时用的是后者。

２.还有就是第一个树状加法的程序，没有看明白啊。

int offset = 1, mask = 1;

__syncthreads();
    while(offset < THREAD_NUM) {
        if((tid & mask) == 0) {
            shared[tid] += shared[tid + offset];
        }
        offset += offset;
        mask = offset + mask;
        __syncthreads();
    }

其实原理是懂得，但是就是从这个程序出发看不懂，似乎无法实现。当然了，后面哪一个改进型的树状加法看明白了。同时对于将整个树状加法扩展开能够提高效率有点不明白，是否是因为ＧＰＵ在进行循环是低效率的啊？？

３.关于随机数的产生

首先是利用srand()初始化随机数发生器，但是有一点不明白就是，差生随机数时，rand(0)和rand(1)有什么区别啊？？

a[i * lda + j] = (float) rand() / RAND_MAX + (float) rand() / (RAND_MAX * RAND_MAX); //在这里的RAND_MAX又是何物？即使是要产生0 ~ 1之间的随机数，第一项(float) rand() / RAND_MAX就可以了啊？

4.最后一个，也是最疑惑的，一关于cuda的一个函式。

cudaMemcpy2D(ac, sizeof(float) * n, a, sizeof(float) * lda,sizeof(float) * n, n, cudaMemcpyHostToDevice);

原文的解释是：因为我们的矩阵乘法函式可以指定 pitch（即 lda、ldb、和 ldc），所以如果用一般的 cudaMemcpy 函式来复制内存的话，会需要每个 row 都分开复制，那会需要呼叫很多次 cudaMemcpy 函式，会使效率变得很差。因此，在这里我们用了一个新的 cudaMemcpy2D 函式，它是用来复制二维数组，可以指定数组的 pitch。这样就可以透过一次函数调用就可以了。疑问是里面的各个变量间是什么关系啊，又是怎么个复制法啊？？

当然了，除这么些个疑问，也还是有很多收获的。

1.blockIdx.x 和 threadIdx.x 都是是 CUDA 内建的变量，它表示的是目前的 block和thread 编号，当然我也猜测blockDim.x指的是block在x方向上的维度，也是cuda内建的变量，它有x、y、z三个方向上的维度。

2.为了同步同一个block里的thread，shared memory变量要用_shared_声明。注意只有同一个block里的thread才能共享，因此当需要多个block计算时，记得要同步的thread都在同一个block中。

3.虽然有一个树状加法看不明白，但是还是学会了另一种变通法子，这个以前在准备考研时看过数据结构，还好有一点印象啊。

4.__syncthreads() 是一个 CUDA 的内部函数，表示 block 中所有的 thread 都要同步到这个点，才能继续执行。在我们的例子中，由于之后要把所有 thread 计算的结果进行加总，所以我们需要确定每个 thread 都已经把结果写到 shared[tid] 里面了。注意，当出现嵌套循环时，要使用该函式一定要注意，容易报错。

5.基本会了利用cuda编程计算矩阵乘法，虽然还没有经过优化，效率、误差还有限，但是至少能够完整看完并且明白一个程序了。也对GPU编程的一般性步骤有了一个了解。

我看得很细，有好多小问题大家可能都觉得不是个问题，嘿嘿，刚想入门，还望海涵啊。