CUDA判断素数的程序

http://blog.csdn.net/smsmn/article/details/6315160 

调用核函数：IsPrime<<<dimGrid,dimBlock>>>()

但是程序里使用的是： IsPrime<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM *  sizeof ( bool )>>>gpudata, result,time);

三个参数，第三个是共享内存的大小。

程序是这样做的，每个块block统一一个结果，因为有共享内存供同一个块共享。

result[bid]= false ;  每个块得出result bool值数组的一个位置。

而 bool  shared[];的大小呢就是每个块中线程的大小。

1：

每个具体的线程，即blockid的某个threadid的线程，所计算的是一个数或者好几个数，怎么说呢，总共有block_NUM*THREAD_NUM这么多的线程，假设要计算素数test，则必须对所有从2,3,4...到test-1这么多数做除法运算，总共应该是test-1个数，这些数就按顺序安排到 block_NUM*THREAD_NUM这些线程上运行，多出来的再次重复的安排也就是说i和 i += BLOCK_NUM * THREAD_NUM都应该是安排在同一个线程块的同一个线程id上运行。

2：

每个线程块中的shared[]的数据怎么综合到一个全局的result[]的一个值上去呢。我们知道result是通过 cudaMalloc(( void **) &result,  sizeof ( bool )*BLOCK_NUM);  分配的内存，所以是全局内存。

而且在什么时候综合这也是个问题，应该是在每个块里面的最后一个线程上综合。

关于关键字：_syncthreads()，共享内存只有在执行_syncthreads()后才能使共享变量对其他线程可见。

可以调用它在内核中指定同步点，块中所有线程都必须在这里等待处理。

而

1.   if(tid == 0)  
2.     {  
3.         for(i=0;i<THREAD_NUM;i++)  
4.         {  
5.             if(shared[i])  
6.             {  
7.                 result[bid]=true;  
8.             }  
9.   
10.         }  
11.     }  
12.     if(tid == 0)   
13.         time[bid + BLOCK_NUM] = clock();  
14.   
15.     time[bid + BLOCK_NUM] = clock();  

代码又使得我们不会同步后在块中每个线程里都做归一处理，只在tid=0的线程。

volatile是一个类型修饰符（type specifier）。它是被设计用来修饰被不同线程访问和修改的变量。

关于clock_t，程序中用来计算时间，time的大小是，2*BLOCK_NUM，因为每个block要记一个开始时间还要记一个结束时间。最后作者还计算了最早块开始时间，和块的最晚结束时间。

cudaMalloc分配的是全局内存