【CUDA】基本知识点

1. 在GPU设备上执行的函数通常称为核函数（Kernel）
2. __global__编译器将告诉编译器：函数应该编译为在设备而不是主机上运行。（将主机代码发送到一个编译器，将设备代码发送到另一个编译器。
3. 程序员一定不能在主机代码中对cudaMalloc()返回的指针进行解引用。主机代码绝对不可以使用这个指针来读取或者写入内存。可以将cudaMalloxc()分配的指针传递给在设备上运行的函数、可以在设备代码中使用cudaMalloc()分配的指针进行内存读/写操作、可以将cudaMalloc()分配的指针传递给在主机上执行的函数、不能在主机代码中使用cudaMalloc()分配的指针进行内存读/写操作。
4. 由3可知不能使用标准C的free()函数来释放cudaMallioc分配的内存，需要调用cudaFree()，访问设备内存的两种最常见方法——在设备代码中使用设备指针以及调用cudaMemcpy()。
5. 总的来说，主机指针只能访问主机代码中的内存，而设备指针也只能访问设备代码中的内存。

6. __syncthreads();                          //对线程中的线程进行同步

   偏移：                                        //偏移的意思更多是指数组的下标

   threadIdx.x

   blockIdx.x

   blockDim.x

   gridDim.x

         //线程格中有线程块，线程块中有线程，gridDim和blockDim分别是线程格和线程块的常数，表明矩阵的维度属性。

        //在多CPU或多核版本中，用while循环对数据进行迭代的数量不是1，而是CPU的数量，在GPU中也可以使用相同的办法，迭代的数量改为并行线程的数量。

        //int offset =x + y * blockDim.x * gridDim.x              //线程块和线程二维分布，讲线程的ID进行一维线性化。

        //将CUDA C的关键字__shared__添加到变量声明中，将使这个变量驻存在此进程块的共享内存中，此线程块的每个线程都共享这个变量。能够实现一个线程块上的线程通过这个变量进行通信和协作。

  

#include "../common/book.h"

#define imin(a,b) (a<b?a:b)

const int N=33*1024;            //总线程数
const int threadsPerBlocks=256;  //值大于每个线程块内的线程数

__global__ void dot(float *a,float *b,float *c){
    __shared__ float cache[threadsPerBlocks];  //共享内存
    int tid= threadIdx.x + blockIdx.x * blockDim.x;  //一维数组的线性化
    int cacheIndex = threadIdx.x;
    float temp = 0;                  //temp用来暂时存储计算结果
    while(tid<N){
        temp+=a[tid] * b[tid];
        tid+=blockDim.x * gridDim.x;   //并行线程的数量
    }
    cache[cacheIndex] = temp   //设置cache相应位置上的值
    //对线程块中的线程进行同步
    __syncthreads();
    int i=blockDim.x / 2;     //进行归约运算
    while(i!=0){
        if(cacheIndex<i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();      //同步线程操作
        i /= 2;
    }
    if(cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}