CUDA学习3

启动一个核函数：

kernel_functionName <<<grid block>>>();

执行配置可以指定线程在GPU上调度运行的方式，第一个是网格的维度，也就是启动块的数目第二个是块维度，也就是每个块中线程的数目，通过这样的方式，我们可以配置：

1. 内核中线程的数目

2. 内核中使用的线程布局

同意块中的线程之间是可以相互合作的，不同块之间的线程是不能合作的，比如说，有32个数据元素用于计算，你就可以这样设定：每个元素8各块，要启动4个块:

kernel_functionName<<<4,8>>>();

上面代码块的配置可以这样看：

图一：配置示例

数据在全局内存中是线性存储的，所以，可以用变量ThreadIdx.x和BlockIdx.x来进行以下操作：

1. 在网格中标示唯一的线程

2. 建立线程和数据元素之间的映射关系

核函数的调用和主机线程是异步的，调用完之后，立即将控制权返回给主机端，也可以强制主机端程序等待所有核函数执行结束：

cudaError_t cudaDeviceSynchronize(void);

一些CUDA运行的时候，API在主机和设备之间是隐式同步的，用cudaMemcpy在主机端和设备之间拷贝数据的时候，主机端隐式同步，即主机端程序必须等待数据拷贝完之后才能继续运行，哇，感觉优点像中断的味道。

所有CUDA核函数的启动都是异步的，CUDA内核调用完成之后，控制权立即返回给CPU

今天要学习的，主要是核函数的编写，什么是核函数嘞，核函数就是在设备端执行的代码，在和函数中，要为一个线程规定要进行的计算以及数据访问，当核函数被调用的时候，多个不同的CUDA线程并行执行同一个计算任务，下面的彪哥是函数类型限定符，指定一个函数在主机上实现还是在设备尝试先，以及是被主机调用还是被设备调用：

限定符	执行位置	调用	备注
__global__	设备	可以从主机调用，也可以从计算能力为3的设备中调用	必须有一个void返回类型
__device__	设备	仅能从设备端调用
__host__	主机	仅能从主机端调用	可以省略

__device__和__host__可以一起使用，这样的话函数可以同时在主机和设备端之间进行编译

CUDA核函数的限制（适用于所有核函数）：

1. 只能放翁设备内存

2. 必须有void返回类型

3. 不支持可变数量的参数

4. 不支持静态变量

5. 显示异步行为

举个例子：两个大小为N的向量A和B相加：

主机端：

void arrayAdd(float*a,float*b,float*c,int n){
    for(int i=0;i<n;i++){
    c[i]=a[i]+b[i]; 
    }
}

kernel:

__global__ void arrayAdd(float*a,float*b,float*c){
    int i=threadIdx.x;
    c[i]=a[i]+b[i];   
}
//然后对于两个长度为32的数组，在main函数中：
arrayAdd(A,B,C,N);
arrayAddOnGPU<<<1,32>>>(A,B,C)

其实啊，你可以将执行参数设置成<<<1,1>>>,在一个块上启动一个线程来执行，其实这就相当于你在GPU上进行了串行处理，这个可以帮你来验证你的运行结果正确与否。

错误处理也很简单，因为CUDA调用时异步的，所以有时候可能很难去确定一个错误到底时由于那一步程序搞错了，所以如果我们可以定义一个错误处理宏来封装所有的CUDA API调用，这就大大的简化了我们的错误检查过程：

#define CHECK(call){
    const cudaError_t error=call;
    if(error!=cudaSucess){
    printf("Error: %s:%d,"__FILE__,__LINE__);
    printf("code : %d,reason:%s\n,",error,cudaGetErrorString(error));
    exit(1);
    }
}

然后你就可以用这样的形式来使用这个宏定义：

CHECK(cudaMemcpy(d_c,gpuref,nbytes,cudaMemcpyHostToDevice));

来检查内存拷贝或者之前异步操作的错误，他会报告错误代码然后输出一系列消息之后，停止我们的程序。

也可以用：

CHECK(cudaDeviceSynchronize());

他会阻塞主机线程的运行直到设备端所有的请求都结束，同时保证最后的核函数启动部分不会出错。

好了，核函数的启动工作就学的差不多了，这下来看看怎么样给核函数计时：

给cpu计时其实是很简单的，引入time.h文件之后，你可以按这篇博文来进行操作：C++中计算程序的运行时间_c++ 程序运行时间_peng1ei的博客-CSDN博客，核函数调用和主机端的程序时异步的，所以我们要用cudaDeviceSynchronize函数来等待所有的GPU线程运行结束才行。书上的cpuSecond（）用不了，不知道为啥，莫名其妙的，是不是说CUDA已经给淘汰了？

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<stdlib.h>
#include<string>
#include<time.h>
#include<stdio.h>
//#define CHECK(call)
//CHECK(call){
//    const cudaError_t error = call;
//    if (error != cudaSucess) {
//    printf("Error: %s:%d,"__FILE__, __LINE__);
//    printf("code : %d,reason:%s\n,", error, cudaGetErrorString(error));
//    exit(1);
//}
//}这个Chesk宏定义，msvc一直显示报错，咋搞都不对，就给注释了。
void initializeData(float* ip, int size) {
    //generate different seed for random number
    time_t t;
    srand((unsigned int)time(&t));
    for (int i = 0; i < size; i++) {
        ip[i] = float(rand() & 0xFF) / 10.0f;
    }
}
void sumArrayOnHost(float* a, float* b, float* c, const int n) {
    for (int  i = 0; i < n; i++)
    {
        c[i] = a[i] + b[i];
    }
}
__global__ void sumArrayOnGPU(float* a, float* b, float* c) {

    int i = threadIdx.x;
    c[i] = a[i] + b[i];

}
void checkresult(float* hostref, float* gpuref, const int n) {
    double eps = 1e-8;
    bool match = 1;
    for (int i = 0; i < n; i++) {
        if (abs(hostref[i] - gpuref[i]) > eps) {
            match = 0;
            printf("Array don't match");
            printf("host %5.2f gpu %5.2f at current %d\n", hostref[i], gpuref[i], i);
            break;
        }
        if (match) {
            printf("array match.\n\n");
        }
    }
}
int main() {
    printf("%s Starting...\n", __argv[0]);
    int dev = 0;
    cudaDeviceProp deviceprop;
    //CHECK(cudaGetDeviceProperties())
    int nElem = 1 << 24;
    printf("Vector size %d\n", nElem);
    size_t nBytes = nElem * sizeof(float);
    float* h_a, * h_b, * hostRef, * gpuRef;
    h_a = (float*)malloc(nBytes);
    h_b = (float*)malloc(nBytes);
    hostRef= (float*)malloc(nBytes);
    gpuRef= (float*)malloc(nBytes);
    //double istrat, ielaps;
    //istrat = cpuSecond();
    initializeData(h_a, nElem);
    initializeData(h_b, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);
    sumArrayOnHost(h_a, h_b, hostRef, nElem);
    float* d_a, * d_b, * d_c;
    cudaMalloc((float**)&d_a, nBytes);
    cudaMalloc((float**)&d_b, nBytes);
    cudaMalloc((float**)&d_c, nBytes);
    cudaMemcpy(d_a, h_a, nBytes, cudaMemcpyDeviceToDevice);
    cudaMemcpy(d_b, h_b, nBytes, cudaMemcpyDeviceToDevice);
    int ilen = 4;
    dim3 block(ilen);
    dim3 grid((nElem + block.x - 1) / block.x);
    sumArrayOnGPU << <grid, block >> > (d_a, d_b, d_c);
    cudaDeviceSynchronize();
    cudaMemcpy(gpuRef, d_c, nBytes, cudaMemcpyDeviceToHost);
    checkresult(hostRef, gpuRef, nElem);
     cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(gpuRef);
    free(hostRef);
    return 0;
}

按书上的代码，会报错，为啥嘞？维度太大了，改一下：

  sumArrayOnGPU << <16384，1024 >> > (d_a, d_b, d_c);

哎，怎么还报错了？？？

数组相加

第零位出错了，这就有些尴尬辽。。。grid和block的维度是可变的，我只能说到这里了，这个报错的问题，得再好好想想。