colab + cuda

谷歌的colab免费提供GPU。但需要安装juputer的一个插件以支持NVCC的源码编译，在代码块中添加如下代码：

# add nvcc support for jupyter
!pip install git+git://github.com/depctg/nvcc4jupyter.git
%load_ext nvcc_plugin

!nvcc --version

一、从hello world开始：要在笔记本中运行代码，请在代码的开头添加%%cu扩展名。

%%cu
#include <stdio.h>

//kernel function
__global__ void helloFromGPU (void)
{
    printf("Hello World from GPU!\n");
}

int main()
{
  helloFromGPU <<<1, 1>>>();
  cudaDeviceSynchronize();
  return 0;
}

        __global__：

        __global__限定词声明一个函数作为一个存在的kernel。这样的一个函数是：

        1、在设备上执行的，

        2、仅可从主机调用。

        3、其调用形式为：helloFromGPU<<<1,10>>>();

        一个kernel是由一组线程执行，所有线程执行相同的代码。上面一行三对尖括号中的1和10 ，表示启动一个 grid 为 螺纹块 的内核。执行配置中的第一个参数1指定网格中线程块的数量，第二个参数10指定线程块中的线程数。

        有一点需要注意的是，printf的输出是在GPU内部执行的，你若想在控制台（网页上）收到该输出，你必须添加

cudaDeviceSynchronize();

二、矩阵加法: 单线程

%%cu
#include <stdio.h>

#define VECTOR_LENGTH 10000 
#define MAX_ERR 1e-4

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

int main()
{
    float *a, *b, *out;
    float *d_a, *d_b, *d_out; 

    //===================步骤1===================
    // Allocate memory on CPU
    a = (float*)malloc(sizeof(float) * VECTOR_LENGTH);
    b = (float*)malloc(sizeof(float) * VECTOR_LENGTH);
    out = (float*)malloc(sizeof(float) * VECTOR_LENGTH);
 
    // data initializtion
    for(int i = 0; i < VECTOR_LENGTH; i++)
    {
        a[i] = 3.0f;
        b[i] = 0.14f;
    }
    //===================步骤1===================
    // Allocate memory on GPU
    cudaMalloc((void**)&d_a, sizeof(float) * VECTOR_LENGTH);
    cudaMalloc((void**)&d_b, sizeof(float) * VECTOR_LENGTH);
    cudaMalloc((void**)&d_out, sizeof(float) * VECTOR_LENGTH);

    //===================步骤2===================
    // copy operator to GPU
    cudaMemcpy(d_a, a, sizeof(float) * VECTOR_LENGTH, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(float) * VECTOR_LENGTH, cudaMemcpyHostToDevice);
    
    //===================步骤3===================
    // GPU do the work, CPU waits
    vector_add<<<1,1>>>(d_out, d_a, d_b, VECTOR_LENGTH);
 
    //===================步骤4===================
    // Get results from the GPU
    cudaMemcpy(out, d_out, sizeof(float) * VECTOR_LENGTH, cudaMemcpyDeviceToHost);
    // Test the result
    for(int i = 0; i < VECTOR_LENGTH; i++)
    {
        printf("%f", out[i]);
    }
    //===================步骤5===================
    // Free the memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_out);
    free(a);
    free(b);
    free(out);
    
}

三、矩阵乘法: 多线程

__global__ void vector_add_fast(float *out, float *a, float *b, int n) {
    __shared__ float temp[10];
    int index = threadIdx.x; // index offset of this thread
    int stride = blockDim.x; // stride step of each iteration
    temp[threadIdx.x] = 0;
    for(int i = index; i < n; i += stride)
    {
        // load data from global memory to shared memory
        // temp[i] = a[i] * b[i];
        temp[threadIdx.x] = temp[threadIdx.x] + a[i] * b[i];
    }
    __syncthreads();
    if (threadIdx.x == 0)
    {
        float sum = 0;
        for (int i = 0; i < stride; i++)
        {
            sum += temp[i];
        }
        *out=sum;
     }
}