cuda编程

最新推荐文章于 2024-07-05 20:32:07 发布

Rareay

最新推荐文章于 2024-07-05 20:32:07 发布

阅读量511

点赞数

分类专栏：技术文章标签： cuda linux 显卡深度学习

本文链接：https://blog.csdn.net/qq_33236581/article/details/101699288

版权

技术专栏收录该内容

25 篇文章 0 订阅

订阅专栏

1 底层结构

在下图中，Host表示CPU，Device表示GPU，Kernel的代码在divice上执行时，实际上是启动很多线程，1个Kernel所启动的所有线程称为1个Grid（网络），Grid又可以分为多个Block（线程块），Block又可以分为多个Thread（线程）；Block在Grid中有自己的索引号，Thread在Block中也有自己的索引号，这些索引号默认是三维的，如果只用到了其中两维，第三维只是默认为1而已；
在这里插入图片描述
GPU的结构图如下：

介绍图中的SP和SM：

SP：最基本的处理单元，streaming processor，也称为CUDA core；最后具体的指令和任务都是在SP上处理的；GPU进行并行计算，也就是很多个SP同时做处理；
SM：多个SP加上其他的一些资源组成一个streaming multiprocessor；也叫GPU大核，其他资源如：warp scheduler，register，shared memory等；SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源；CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。

软件意义上的线程及线程块与应将的对应关系如图：
在这里插入图片描述
图中：

每个线程由每个线程处理器（SP）执行
线程块由多核处理器（SM）执行
一个kernel其实由一个grid来执行，一个kernel一次只能在一个GPU上执行

block是软件概念，一个block只会由一个sm调度，程序员在开发时，通过设定block的属性，告诉GPU硬件，我有多少个线程，线程怎么组织；而具体怎么调度由sm的warps scheduler负责，block一旦被分配好SM，该block就会一直驻留在该SM中，直到执行结束；一个SM可以同时拥有多个blocks，但需要序列执行；

2 CUDA内存模型

CUDA中的内存模型分为以下几个层次：

每个线程都用自己的registers（寄存器）
每个线程都有自己的local memory（局部内存）
每个线程块内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源
每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用
每个grid都有自己的constant memory（常量内存）和texture memory（纹理内存），），不同线程块的线程都可使用

线程访问这几类存储器的速度是register > local memory >shared memory > global memory
下面这幅图表示就是这些内存在计算机架构中的所在层次:
在这里插入图片描述

3 CUDA编程模型

3.1 指定函数的运行环境（CPU/GPU）

通过关键字就可以表示某个程序在CPU上跑还是在GPU上跑！如下表所示，比如我们用__global__定义一个kernel函数，就是CPU上调用，GPU上执行，注意__global__函数的返回值必须设置为void：

修饰符	运行在	调用于
`__device__` float DeviceFunc()	device	device
`__global__` float KernelFunc()	device	host
`__host__` float HostFunc()	host	host

注：这里的device就是GPU，host是CPU；

3.2 CPU和GPU间的数据传输怎么写

1）在显存中分配、回收内存的函数接口：

cudaMalloc()// 在设备端分配global memory
cudaFree()// 释放存储空间

2）CPU的数据和GPU端数据做数据传输的函数接口是一样的，他们通过传递的函数实参（枚举类型）来表示传输方向：

cudaMemcpy(void dst, void src, size_t nbytes, enum cudaMemcpyKind direction)
//其中：
// 	enum cudaMemcpyKind:
// 		cudaMemcpyHostToDevice（CPU到GPU）
// 		cudaMemcpyDeviceToHost（GPU到CPU）
// 		cudaMemcpyDeviceToDevice（GPU到GPU）

3.2 创建线程组织模型

我们可以用dim3类来表示网格和线程块的组织方式，网格grid可以表示为一维和二维格式，线程块block可以表示为一维、二维和三维的数据格式。

dim3 DimGrid(100, 50);  //5000个线程块，维度是100*50
dim3 DimBlock(4, 8, 8);  //每个线层块内包含256个线程，线程块内的维度是4*8*8

3.3 线程号的映射

映射所用到的一些变量名参考下图，gridDim和blockDim分别表示其矩阵块的尺寸，这里只是二维的，如果是三维的自己脑补；
在这里插入图片描述

// 情况1：grid划分成1维，block划分为1维。  
__device__ int getGlobalIdx_1D_1D() {  
    int threadId = blockIdx.x *blockDim.x + threadIdx.x;  
    return threadId;  
}  
  
// 情况2：grid划分成1维，block划分为2维。  
__device__ int getGlobalIdx_1D_2D() {  
    int threadId = blockIdx.x * blockDim.x * blockDim.y  
        + threadIdx.y * blockDim.x + threadIdx.x;  
    return threadId;   
}  
  
// 情况3：grid划分成1维，block划分为3维。  
__device__ int getGlobalIdx_1D_3D() {  
    int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z  
        + threadIdx.z * blockDim.y * blockDim.x  
        + threadIdx.y * blockDim.x + threadIdx.x;  
    return threadId;  
}  
  
// 情况4：grid划分成2维，block划分为1维。  
__device__ int getGlobalIdx_2D_1D() {  
    int blockId = blockIdx.y * gridDim.x + blockIdx.x;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
    return threadId;  
}  
  
// 情况5：grid划分成2维，block划分为2维。  
__device__ int getGlobalIdx_2D_2D() {  
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
        + (threadIdx.y * blockDim.x) + threadIdx.x;  
    return threadId;  
}  
  
// 情况6：grid划分成2维，block划分为3维。  
__device__ int getGlobalIdx_2D_3D() {  
    int blockId = blockIdx.x + blockIdx.y * gridDim.x;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
        + (threadIdx.z * (blockDim.x * blockDim.y))  
        + (threadIdx.y * blockDim.x) + threadIdx.x;  
    return threadId;  
}  
  
// 情况7：grid划分成3维，block划分为1维。  
__device__ int getGlobalIdx_3D_1D() {  
    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
        + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * blockDim.x + threadIdx.x;  
    return threadId;  
}  
  
// 情况8：grid划分成3维，block划分为2维。  
__device__ int getGlobalIdx_3D_2D() {  
    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
        + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y)  
        + (threadIdx.y * blockDim.x) + threadIdx.x;  
    return threadId;  
}  
  
// 情况9：grid划分成3维，block划分为3维。  
__device__ int getGlobalIdx_3D_3D() {  
    int blockId = blockIdx.x + blockIdx.y * gridDim.x  
        + gridDim.x * gridDim.y * blockIdx.z;  
    int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)  
        + (threadIdx.z * (blockDim.x * blockDim.y))  
        + (threadIdx.y * blockDim.x) + threadIdx.x;  
    return threadId;  
}

3.4 CUDA实例

代码所对应文件名后缀为.cu，写好代码后，用命令nvcc main.cu或命令nvcc main.cu -o main进行编译；

本程序可以用于读取显卡信息，因为在建立显卡线程时需要分配线程块的个数和线程块里面的线程个数

#include "device_launch_parameters.h"
#include <iostream>

int main()
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    for(int i=0;i<deviceCount;i++)
    {
        cudaDeviceProp devProp;
        cudaGetDeviceProperties(&devProp, i);
        std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
        std::cout << "设备全局内存总量： " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
        std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
        std::cout << "每个线程块的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
        std::cout << "每个线程块的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
        std::cout << "设备上一个线程块（Block）种可用的32位寄存器数量： " << devProp.regsPerBlock << std::endl;
        std::cout << "每个EM的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
        std::cout << "每个EM的最大线程束数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
        std::cout << "设备上多处理器的数量： " << devProp.multiProcessorCount << std::endl;
        std::cout << "======================================================" << std::endl;        
    }
    return 0;
}

本程序是在GPU上做矩阵乘法运算

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <sys/time.h> 
#include <stdio.h>
#include <math.h>

// 矩阵类
class matrix
{
public:
    int *data;
    int width;
    int highly;
    int size;
    void malloc_data()
    {
        size = width * highly;
        data = (int *)malloc(sizeof(int) * size);
    }
};

// GPU里面的矩阵乘法运算
__global__ void matrix_mul_gpu(int *M, int* N, int* P, int width)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    int W = gridDim.x * blockDim.x;
                
    int sum = 0;
    for(int k = 0; k < width; k++)
    {
        int a = M[j * width + k];
        int b = N[k * W + i];
        sum += a*b;
    }
    P[j * W + i] = sum;
}

// 打印矩阵
void show_matrix(matrix * N)
{
    printf("--------------------\n");
    for (int i=0; i<N->highly; i++)
    {
        for (int j =0; j < N->width; j++)
        {
            printf("%d  ", N->data[i * N->width + j]);
        }
        printf("\n");
    }
}

// API
// 本程序对线程分配只有 1 个线程块，里面的线程是二维的，尺寸和运算结果的矩阵尺寸一样
int CudaMatrixMul(matrix *A, matrix *B, matrix *C)
{
    if (A->size == 0 || B->size == 0 || C->size == 0 || A->width != B->highly)
        return -1;

    // 给C分配控件
    C->width = B->width;
    C->highly = A->highly;
    C->malloc_data();

    // GPU上创建矩阵
    int *d_A, *d_B, *d_C;
    cudaMalloc((void**)&d_A, sizeof(int) * A->size);
    cudaMalloc((void**)&d_B, sizeof(int) * B->size);
    cudaMalloc((void**)&d_C, sizeof(int) * C->size);

    // 把CPU的矩阵赋值给GPU的矩阵
    cudaMemcpy(d_A, A->data, sizeof(int) * A->size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B->data, sizeof(int) * B->size, cudaMemcpyHostToDevice);

    // 创建CUDA线程
    dim3 threadPerBlock(C->width, C->highly);
    dim3 blockNumber(1, 1);
    matrix_mul_gpu << <blockNumber, threadPerBlock >> > (d_A, d_B, d_C, A->width);

    // 把GPU计算出的结果复制到CPU
    cudaMemcpy(C->data, d_C, sizeof(int) * C->size, cudaMemcpyDeviceToHost);

    // 释放GPU的矩阵
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    return 0;
}

int main()
{
    // CPU上创建矩阵
    int *d;
    matrix A, B, C;
    A.width = 2;
    A.highly = 3;
    A.malloc_data();
    d = A.data;
    *d++ = 1; *d++ = 2;
    *d++ = 3; *d++ = 4;
    *d++ = 5; *d++ = 6;

    B.width  = 4;
    B.highly = 2;
    B.malloc_data();
    d = B.data;
    *d++ = 1; *d++ = 2; *d++ = 3; *d++ = 4; 
    *d++ = 5; *d++ = 6; *d++ = 7; *d++ = 8; 

	// 调用上面的API
    CudaMatrixMul(&A, &B, &C);
	
	// 打印矩阵
    show_matrix(&A);
    show_matrix(&B);
    show_matrix(&C);

    free(A.data);
    free(B.data);
    free(C.data);
}

注：用 nvcc xxx.cu 即可编译.cu文件，和 gcc 编译一样；

4 混合编译.c/.cpp与.cu文件

参考

.cu文件是CUDA的kernel函数，需要nvcc编译器来编译；.c/.cpp是c代码，这里用gcc编译；假如现在有文件：

test1.cu
test2.c

则具体编译指令、编译脚本如下：

方法1：分别编译各个文件

nvcc -c test1.cu
gcc -c test2.c
gcc -o testc test1.o test2.o -lcudart -L/usr/local/cuda/lib64

方法2: 将cuda程序编译为静态库

nvcc -lib test1.cu -o libtestcu.a
gcc test2.c -ltestcu -L. -lcudart -L/usr/local/cuda/lib64

方法3：将CUDA程序弄成动态库（以makefile为例）

all: c

c: libtestcu.so
    gcc test2.c -ltestcu -L. -lcudart -L/usr/local/cuda/lib64 -o testc

libtestcu.so: test.cu
    nvcc -o libtestcu.so -shared -Xcompiler -fPIC test1.cu

方法4：基于CMake
暂无
参见这里
$E=mc^2$

Rareay

关注

0
点赞
踩
1

收藏

觉得还不错? 一键收藏
0
评论
cuda编程

目录1 底层结构2 CUDA内存模型3 CUDA编程模型3.1 指定函数的运行环境（CPU/GPU）3.2 CPU和GPU间的数据传输怎么写3.2 创建线程组织模型3.3 线程号的映射3.4 CUDA实例参考1 底层结构在下图中，Host表示CPU，Device表示GPU，Kernel的代码在divice上执行时，实际上是启动很多线程，1个Kernel所启动的所有线程称为1个Grid（网络）...
复制链接

扫一扫

专栏目录