cuda核简单了解

#include <stdio.h>
#include <cuda_runtime.h>

/*
CUDA核函数
__global__ 核函数的前缀定义
    - 使用__global__修饰的函数，必须是void无返回值
    - __global__核函数修饰，必须是nvcc编译才有效，否则无效
    - __global__修饰的函数，使用name<<<grid, block, memory, stream>>>(params)启动核函数
        - 启动在host，但是执行在device

__device__ 修饰的函数，只能够在设备上执行，设备上调用(例如核函数内调用)
    - nvidia提供了很多很多内置设备函数，比如日常的cos、sin之类的
        - 在nvidia团队中，不同内置函数的api接口版本号，被称之为，计算能力

定义如下：
__device__，函数执行在设备上
__global__，函数执行在设备上，但是调用在host上。定义核函数的符号
__host__，函数执行在host上，调用也在host上
*/

// sigmoid 不能够使用 sigmoid<<<1, 3>>>这种启动他
// 你也不能够直接 sigmoid(0.1)
// 只能够在核函数内调用他
__device__ float sigmoid(float value){
    return 1 / (1 + exp(-value));
}

__global__ void compute(float* a, float* b, float* c){
    /*
    线程layout的概念，启动的线程会被设计为grid和block，如同提供的参数那样
    这个layout的概念是虚拟的，通过cuda驱动实现真实硬件映射，抽象了一个中间层（调度层）
       - 如果我们有4352 Core
       - 如果我们需要启动 5000 个线程
       - 抽象层他会把5000个线程安排到各个Core中去执行，根据情况来执行数次
       - 每次调度单位WarpSize，如果启动的线程不足，也会执行WarpSize，不过core是非激活状态而已

    你想告诉他启动多少个线程，通过gridDim和blockDim告诉他
    线程数 = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

    2147483647 是 有符号整数(int)的最大值
    65535      是 无符号短整数(int16, short)的最大值
    gridDim的最大值范围： (x,y,z): (2147483647, 65535, 65535)
    blockDim的最大值范围：(x,y,z): (1024, 1024, 64)

    他的定义，其实在 device_launch_parameters.h
    uint3 __device_builtin__ __STORAGE__ threadIdx;
    uint3 __device_builtin__ __STORAGE__ blockIdx;
    dim3 __device_builtin__ __STORAGE__ blockDim;
    dim3 __device_builtin__ __STORAGE__ gridDim;
    int __device_builtin__ __STORAGE__ warpSize;

    获取线程ID，进行数据操作
    数据索引，是通过blockIdx和threadIdx计算得到
    gridDim告诉你Grid的大小 ，blockDim告诉你block大小
    blockIdx告诉你所在Grid内的索引，  threadIdx告诉你所在block的索引
    把gridDim和blockDim设想为一个tensor
    gridDim的shape  = gridDim.z   x  gridDim.y  x  gridDim.x
    blockDim的shape = blockDim.z  x  blockDim.y x  blockDim.x
    最终的启动线程的shape维度是：gridDim.z   x  gridDim.y  x  gridDim.x  x  blockDim.z  x  blockDim.y x  blockDim.x

    如果启动的线程是个6维度tensor，那么索引，是不是也可以类似
    blockIdx.z, blockIdx.y, blockIdx.x, threadIdx.z, threadIdx.y, threadIdx.x

    仅仅是在这个场景需要把6个维度索引变为连续的内存单元索引
    如果我们有6个维度a, b, c, d, e, f，如果我们有6个位置索引u, v, w, x, y, z
    position = ((((u * b + v) * c + w) * d + x) * e + y) * f + z
    */

    // gridDim = 1 x 1 x 1
    // blockDim = 3 x 1 x 1
    // 
    int position = blockDim.x * blockIdx.x + threadIdx.x;
    c[position] = a[position] * sigmoid(b[position]);
}

 
int main(){

    // 定义3个数组，host指针
    const int num = 3;
    float a[num] = {1, 2, 3};
    float b[num] = {5, 7, 9};
    float c[num] = {0};

    // 定义3个设备指针，device指针
    size_t size_array = sizeof(c);
    float* device_a = nullptr;
    float* device_b = nullptr;
    float* device_c = nullptr;

    // 分配设备空间，大小是size_array，单位byte
    cudaMalloc(&device_a, size_array);
    cudaMalloc(&device_b, size_array);
    cudaMalloc(&device_c, size_array);

    // 把数据从host复制到device，其实就是主机复制到显卡
    // 复制的是a和b
    cudaMemcpy(device_a, a, size_array, cudaMemcpyHostToDevice);
    cudaMemcpy(device_b, b, size_array, cudaMemcpyHostToDevice);

    // 执行核函数，结果放到c上
    compute<<<dim3(1), dim3(3)>>>(device_a, device_b, device_c);

    // 把计算后的结果c复制回主机
    cudaMemcpy(c, device_c, size_array, cudaMemcpyDeviceToHost);

    // 查看主机上的c内容是多少
    for(int i = 0; i < num; ++i)
        printf("c[%d] = %f\n", i, c[i]);
    return 0;
}