cuda编程基础

1.基础知识整理

一个典型的CUDA程序是按这样的步骤执行的：

1. 把数据从CPU内存拷贝到GPU内存。
2. 调用核函数对存储在GPU内存中的数据进行操作的。
3. 将数据从GPU内存传送回CPU内存。

   一般CPU一个核只支持一到两个硬件线程，而GPU往往在硬件层面上就支持同时成百上千个并发线程。不过这也要求我们在GPU编程中更加高效地管理这些线程，以达到更高的运行效率。在CUDA编程中，线程是通过线程网格（Grid）、线程块（Block）、线程束（Warp）、线程（Thread）这几个层次进行管理的.

  第二点，为了达到更高的效率，在CUDA编程中我们需要格外关注内存的使用。与CPU编程不同，GPU中的各级缓存以及各种内存是可以软件控制的，在编程时我们可以手动指定变量存储的位置。具体而言，这些内存包括寄存器、共享内存、常量内存、全局内存等。这就造成了CUDA编程中有很多内存使用的小技巧，比如我们要尽量使用寄存器，尽量将数据声明为局部变量。而当存在着数据的重复利用时，可以把数据存放在共享内存里。而对于全局内存，我们需要注意用一种合理的方式来进行数据的合并访问，以尽量减少设备对内存子系统再次发出访问操作的次数。

线程管理:

内存管理: 

其中寄存器(Registers)是GPU上运行速度最快的内存空间，通常其带宽为8TB/s左右，延迟为1个时钟周期。核函数中声明的一个没有其他修饰符的自变量，通常就存储在寄存器中。最快速也最受偏爱的存储器就是设备中的寄存器，属于具有重要价值有极度缺乏的资源。

接下来是共享内存(shared memory)，共享内存是GPU上可受用户控制的一级缓存。共享内存类似于CPU的缓存，不过与CPU的缓存不同，GPU的共享内存可以有CUDA内核直接编程控制。由于共享内存是片上内存，所以与全局内存相比，它具有更高的带宽与更低的延迟，通常其带宽为1.5TB/s左右，延迟为1～32个时钟周期。对于共享内存的使用，主要考虑数据的重用性。当存在着数据的重复利用时，使用共享内存是比较合适的。如果数据不被重用，则直接将数据从全局内存或常量内存读入寄存器即可。

全局内存(global memory)是GPU中最大、延迟最高并且最常使用的内存。全局内存类似于CPU的系统内存。在编程中对全局内存访问的优化以最大化程度提高全局内存的数据吞吐量是十分重要的。

1.向量相加的cuda实现

(按对应位置相加)

#include <iostream>
#include <cuda_runtime.h>

#define length 10000


//__global__ int foo(int a){}表示一个内核函数，
//是一组由GPU执行的并行计算任务，以foo<<>>(a)的形式或者driver API的形式调用。
//目前__global__函数必须由CPU调用，并将并行计算任务发射到GPU的任务调用单元。随着GPU可编程能力的进一步提高，未来可能可以由GPU调用。
__global__ void  vector_add(int *a,int *b,int *& c)
{
    for(int i=0;i<length;i++)
    {
        c[i]=a[i]+b[i];
    }
}




int main()
{
    int *a , *b, *c,*a_G , *b_G, *c_G;
    a=new int[length];
    b=new int[length];
    c=new int[length];

    for (int i =0;i<length;i++)
    {
        a[i]=i+1;
        b[i]=i+2;
    }


    //gpu分配空间
    //cudaMalloc
    //函数原型:      cudaError_t cudaMalloc (void **devPtr, size_t  size );
    //此函数返回值是CUDA中定义的一个错误代码。
    cudaMalloc((void**)&a_G, sizeof(int) * length);
    cudaMalloc((void**)&b_G, sizeof(int) * length);
    cudaMalloc((void**)&c_G, sizeof(int) * length);


    //CPU变量给GPU赋值
    cudaMemcpy(a_G,a,sizeof(int )*length,cudaMemcpyHostToDevice );
    cudaMemcpy(b_G,b,sizeof(int )*length,cudaMemcpyHostToDevice );




    //函数名称<<<block 数目, thread 数目, shared memory 大小>>>(參数...);
    vector_add<<<1,1>>>(a_G,b_G,c_G);


    //Gpu结果取到cpu
    cudaMemcpy(c,c_G, sizeof(int )* length,cudaMemcpyDeviceToHost);

    for (int i=0;i<length;i++)
    {
        cout<<c[i]<<endl;
    }


    return 0;
}

核函数:

__global__ int foo(int a){}表示一个内核函数，
是一组由GPU执行的并行计算任务，以foo<<>>(a)的形式或者driver API的形式调用。
目前__global__函数必须由CPU调用，并将并行计算任务发射到GPU的任务调用单元。随着GPU可编程能力的进一步提高，未来可能可以由GPU调用。

cudaMalloc函数:

函数原型:

cudaError_t cudaMalloc (void ** devPtr, size_t size );

返回值是cuda中定义的错误代码.

第一个变量:指向指针变量的指针,参数传递的是存储在cpu内存中的指针变量的地址，cudaMalloc在执行完成后，向这个地址中写入了一个GPU的地址值

第二个变量:大小.

cudaMemcpy函数:

cudaMemcpy用于在主机（Host）和设备（Device）之间往返的传递数据，用法如下：

主机到设备：cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice)

设备到主机：cudaMemcpy(h_A,d_A,nBytes,cudaMemcpyDeviceToHost)
拷贝数据时cpu主进程是锁死的.

2.向量相加的并行

cuda变成只单线程与在CPU下差不多,没有任何意义,因此要进行并行才能体现GPU的优势.

＞＞＞＞＞＞＞＞＞＞未完待续

3内置变量