工作总结@2010.3.16

在CUDA中，一般的数据复制到的显卡内存的部分，称为global memory。这些内存时没有cache 的，而且，存取global memory所需要的时间（即latency）是非常长的，通常是数百个cycles. 如果我们的程序只有一个thread，所以每次它读取global memory的内容，就要等到实际读取到数据、累加到sum之后，才能进行下一步，这就是为什么如果采取一个线程搞的话表现会很差。

       由于global memory并没有cache，所以要避免开巨大的latency的方法，就是要利用大量的threads。假设现在又大量的threads在同时执行，那么当一个thread读取内存，开始等待结果的时候，GPU就可以立刻切换到下一个thread，并读取下一个内存位置。因此，理想上当thread的数目够多的时候，就可以完全把global memory的巨大latency隐藏起来了。

       想一下如果把计算平方和的程序并行换额？其中的一个比较好的方法是把数字分成若干组，把各组数字分别计算平方和后，最好再把每组的和加总起来就可以了，所以，一开始，我们可以把最后加总的动作，由CPU来进行。

       现有一kernel函数__global__ static void sumOfQuery(int*, int* clock_t*)

n  第一种方法“伪多线程法”

即在主函数中调用<<<1, 1, 0>>>sumOfQuery(…)

即采用一个block 一个线程 串行实现。

for (int i= 0, i < element_size; ++i)

       sum += elem[i]

n  第二种采用多线程法

即在主函数中调用<<<1, thread_size, 0>>>sumOfQuery

Sum += result[i] (i从0到thread_size) //host

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

存在缺点为内存延迟大，因为当第一个线程等到内存数据时，第二个线程就开始运行，所以数据没有成块。

n  第三种采用改进型多线程法

Size = DATA_SIZE / THREAD_NUM

for (i = tid; i < DATA_SIZE; i += THREAD_NUM){

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

       }

result[tid] = sum;

此种方法的采用的很大一部分原因在于显卡上的内存是DRAM，因此基于它的操作（最有效的存取方式）是以连续的方式存取。上一个程序中，虽然表面看起来是连续存储内存位置，但是在实际上是与想法违背的。

实验结果发现虽然latency有所减少，但没有取得明显效果，主要原因在于thread数目较小，即N个thread最多只能隐藏N个latency。所以，如果增加thread数目，就可以看到更好的效率，例如可以把thread的数量调的更高，但不能超过GPU所能执行的最大线程数目。

失之东隅收之桑榆，thread数目过多，同样会导致cpu端的工作也可能家走。

n  第四种是采用块的形式

即添加block块，在cuda中，thread是可以分组的，一个block中的thread，具有一个共享的share memory，也可以进行同步工作。不同block之间的thread则不行，在计算平方和这个项目中，则并不太需要进行thread的同步动作，因此我们可以使用多个block来进一步增加thread的数目。

for (I = bid * THREAD_NUM + tid; i < DATA_SIZE;

i += BLOCK_NUM * THREAD_NUM)

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

       result[bid * THREAD_NUM + tid] = sum;

n  第五种是采取块内同步算法与树状算法，正在看中