CUDA 学习记录 [1]

争取3个月的时间对cuda基础操作熟悉

数据并行首先需要给线程做分配

块划分分多个数据小块随机处理
周期划分固定轮流每个线程执行数据块

计算机架构:

常见的是单指令 → 多数据多指令 → 多数据

关键词: 延迟(操作时间) 带宽(单位时间处理数据量) 吞吐量(单位时间成功运算数量)

内存视角来说:

分布式内存 → 机房服务器集群内存通过网络连接

单主板多处理器 → 共享用一片内存寻址空间

CPU和GPU线程的区别：

CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。
CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而GPU核则是大量线程，最大幅度提高吞吐量

对于GPU resource API 调用有两种方式(不可以混合调用)

CUDA Driver → CPU 中最底层紧贴GPU
CUDA Runtime → higher level API

一个CUDA应用通常可以分解为两部分，

CPU 主机端代码
GPU 设备端代码

CUDA nvcc编译器会自动分离你代码里面的不同部分，如图中主机代码用C写成，使用本地的C语言编译器编译，设备端代码，也就是核函数，用CUDA C编写，通过nvcc编译，链接阶段，在内核程序调用或者明显的GPU设备操作时，添加运行时库。

NVCC ← LLVM 编译器

CUDA toolkit 包含 nvcc? libraries, developer tools

我们先开始看hello world

什么叫做在Kernel上编程?

我们首先要了解Kernel是什么 Kernel其实是function that executed in parallel fashion by multiple threads

thread.index

grid → block: {thread1, thread2, } Threads in a block work together and share memory

sm负责执行 thread blocks 在这里 further organized into wraps

(32 threads execute instructions in a single cycle)

Each SM has cores, registers, shared memory (Core number varies)

multiple SMs can execute different thread blocks simultaneously(→ direct parallelism)

large number of thread blocks will be distributed automatically?

Launch Kernel → need to define #of block/thread to be created

each thread execute independently, and what data it operates is controlled by thread.index

我们来看一个基础hello world

/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void)
{
  printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset();//if no this line ,it can not output hello world from gpu
  return 0;
}

一个更加直观的

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

__global__ void hello_world() {
printf("Hello from thread %d in block %d!\n", threadIdx.x, blockIdx.x);
}

int main() {
int numBlocks = 1;
int numThreadsPerBlock = 10;

// Launch kernel
hello_world<<<numBlocks, numThreadsPerBlock>>>();

// Wait for GPU to finish
cudaDeviceSynchronize();

return 0;
}

几个关键词:

global declare kernel function (run on GPU, called from CPU) 注意这里的从主机调用

device 从设备端调用

Kernel核函数编写有以下限制

只能访问设备内存
必须有void返回类型
不支持可变数量的参数
不支持静态变量
显示异步行为

<<<#Block, #Threads>>> 用来specify Kernel function 还可以规定3D block size

接受dim3类型 grid/block维度或int 这里是说每一个B对应的threads数

cudaDeviceSynchronize() → Blocks CPU thread until all threads on GPU finished 同步CPU/GPU

两个设备本身是异步的也就是CPU线程调用完kernel func 就会继续这个program

不管执行情况

ensure all launched kernels and memory operations are completed before program proceeds

一般CUDA程序分成下面这些步骤：

分配GPU内存
拷贝内存到设备
调用CUDA内核函数来执行计算
把计算完成数据拷贝回主机端
内存销毁

CUDA抽象了硬件实现：

线程组的层次结构
内存的层次结构
障碍同步

我们看怎么把简单操作转换成并行

但是这样是否以为这我们需要 size n个的线程量

void sumArraysOnHost(float *A, float *B, float *C, const int N) {
  for (int i = 0; i < N; i++)
    C[i] = A[i] + B[i];
}

__global__ void sumArraysOnGPU(float *A, float *B, float *C) {
  int i = threadIdx.x;
  C[i] = A[i] + B[i];
}

验证kernel func

/*
* https://github.com/Tony-Tan/CUDA_Freshman
* 3_sum_arrays
*/
#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"

void sumArrays(float * a,float * b,float * res,const int size)
{
  for(int i=0;i<size;i+=4)
  {
    res[i]=a[i]+b[i];
    res[i+1]=a[i+1]+b[i+1];
    res[i+2]=a[i+2]+b[i+2];
    res[i+3]=a[i+3]+b[i+3];
  }
}
__global__ void sumArraysGPU(float*a,float*b,float*res)
{
  int i=threadIdx.x;
  res[i]=a[i]+b[i];
}
int main(int argc,char **argv)
{
  int dev = 0;
  cudaSetDevice(dev);

  int nElem=32;
  printf("Vector size:%d\n",nElem);
  int nByte=sizeof(float)*nElem;
  float *a_h=(float*)malloc(nByte);
  float *b_h=(float*)malloc(nByte);
  float *res_h=(float*)malloc(nByte);
  float *res_from_gpu_h=(float*)malloc(nByte);
  memset(res_h,0,nByte);
  memset(res_from_gpu_h,0,nByte);

  float *a_d,*b_d,*res_d;
  CHECK(cudaMalloc((float**)&a_d,nByte));
  CHECK(cudaMalloc((float**)&b_d,nByte));
  CHECK(cudaMalloc((float**)&res_d,nByte));

  initialData(a_h,nElem);
  initialData(b_h,nElem);

  CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));

  dim3 block(nElem);
  dim3 grid(nElem/block.x);
  sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);
  printf("Execution configuration<<<%d,%d>>>\n",block.x,grid.x);

  CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
  sumArrays(a_h,b_h,res_h,nElem);

  checkResult(res_h,res_from_gpu_h,nElem);
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);

  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);

  return 0;
}

以上使用了几个基础的内存管理 api

cudaMalloc

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

float* d_array; !!申明一个设备指针
cudaMalloc((void**)&d_array, 100 * sizeof(float));  // 分配400字节
!!!使设备指针的值 对应新分配的设备内存地址

#注意这里的void ** 可以被省略 如下一个例子可见

只有在malloc中需要 pointer 因为需要

cudaMemcpy

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

cudaMemcpyHostToHost     // 极少使用
cudaMemcpyHostToDevice   // CPU→GPU
cudaMemcpyDeviceToHost   // GPU→CPU 
cudaMemcpyDeviceToDevice // GPU→GPU

// CPU数据准备
float h_data[100] = {...}; 

// 分配设备内存
float* d_data;
cudaMalloc(&d_data, sizeof(h_data));

// 主机到设备拷贝
cudaMemcpy(d_data, h_data, sizeof(h_data), cudaMemcpyHostToDevice);

// 设备到主机回传
float result[100];
cudaMemcpy(result, d_data, sizeof(result), cudaMemcpyDeviceToHost);

#注意这里的顺序 a to b 也就是 b在前接受 a在后面 所以HostToDevice Device在前
#而且注意 这里先 cuda malloc 分配后才copy 同等size memory

cudaMemset → Device memory 初始化

cudaError_t cudaMemset(void* devPtr, int value, size_t count);

int* d_flags;
cudaMalloc(&d_flags, 1000 * sizeof(int));

// 将所有4字节int设为0x3f（实际值为0x3f3f3f3f）
cudaMemset(d_flags, 0x3f, 1000 * sizeof(int)); 

cudaMemset(d_flags, 0, 1000*sizeof(int))

cudaFree

只能释放通过 cudaMalloc/cudaMallocManaged 分配的内存

释放后应将指针设为 NULL 防止野指针

cudaError_t cudaFree(void* devPtr);

float* d_data = NULL;
cudaMalloc(&d_data, 1024);

// 使用设备内存...

cudaFree(d_data);
d_data = NULL;  // 重要：避免悬垂指针

参考:

https://face2ai.com/CUDA-F-2-0-CUDA编程模型概述1/