cuda编程入门（2）—— 浅谈核函数

​​cuda编程入门（1)​​

执行空间修饰符有三个 ​​__global__​​​, ​​__device__​​​, ​​__host__​​。

Function Execution Space Specifiers

参考官网文档 ​​Function Execution Space Specifiers​​

global

主要有以下几点：

1.Executed on the device, Callable from the host.
2.must specify its execution configuration
3.asynchronous, meaning it returns before the device has completed its execution.

device
主要有以下几点：

1.Executed on the device,
2.Callable from the device only.

// 两个向量加法kernel，grid和block均为一维
__device__ float* add_2(float* x, float * y, float* z, int n) {
    // 获取该线程的全局索引
    int index = threadIdx.x + blockIdx.x * blockDim.x;
    // 步长(线程总数)
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
    {
        z[i] = x[i] + y[i];
    }
    return z;
}
__global__ void add(float* x, float * y, float* z, int n)
{
    add_2(x, y, z, n);
}

#include <iostream>
int main()
{
    int N = 1 << 20; //(1048576)
    int nBytes = N * sizeof(float);
    // 申请host内存
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);

    // 初始化数据
    for (int i = 0; i < N; ++i)
    {
        x[i] = 10.0;
        y[i] = 20.0;
    }

    // 申请device内存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_x, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);

    // 将host数据拷贝到device
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    // 定义kernel的执行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // 执行kernel
    add << < gridSize, blockSize >> >(d_x, d_y, d_z, N);

    // 将device得到的结果拷贝到host
    cudaMemcpy((void*)z, (void*)d_z, nBytes, cudaMemcpyDeviceToHost);

    // 检查执行结果
    float maxError = 0.0;
    for (int i = 0; i < N; i++)
        maxError = fmax(maxError, fabs(z[i] - 30.0));
    std::cout << "最大误差: " << maxError << std::endl;

    // 释放device内存
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    // 释放host内存
    free(x);
    free(y);
    free(z);

    return 0;
}

host
主要有以下几点：

1.Executed on the host,
2.Callable from the host only.

Execution Configuration

​​execution-configuration​​

Any call to a global function must specify the execution configuration for that call. The execution configuration defines the dimension of the grid and blocks that will be used to execute the function on the device, as well as the associated stream (see CUDA Runtime for a description of streams).

执行配置主要四个参数

<<< Dg, Db, Ns, S >>>

即 grid的维度信息， block的维度信息，共享内存以及cuda_stream。详细信息参考官网信息。
grid， block设置好了之后，可以根据内置变量 ​​​gridDim​​​,​​blockDim​​​,​​blockIdx​​​,​​threadIdx​​,来计算线程号。

共享内存 Ns 以及 cuda_stream S 在后续博客中介绍。