CUDA编程模型介绍

YYRAN_ZZU

已于 2024-10-27 21:33:01 修改

阅读量1.7k

点赞数 31

CC 4.0 BY-SA版权

文章标签： gpu算力 c++ linux

于 2024-10-26 22:44:45 首次发布

本文链接：https://blog.csdn.net/YYRAN_ZZU/article/details/143259973

1.介绍

CUDA是一种通用的并行计算平台和编程模型，在C语言基础扩展，借助于cuda可以像编写C语言一样的实现并行算法，本文主要介绍基础的编程概念和程序入门的第一个cuda程序。

2.CUDA编程模型概述

2.1 CUDA编程结构

cuda编程模型使用C语言扩展成的代码在异构计算系统中执行应用程序，一个异构环境中包含多个CPU和GPU，下面介绍两个常用的专用概念词

主机：CPU及其内存
设备：GPU及其内存
设备(kernel)-指代码在GPU上运行的代码

如下图所示
编程模型
一般一个典型的CUDA程序实现的流程如下

数据从CPU内存拷贝到GPU内存
调用核函数(kernel)对存储在GPU内存中的数据进行操作
将数据从GPU内存送回到CPU内存

参考官网流程示意图如下
在这里插入图片描述

2.2 内存管理

CUDA模型假设系统由一个主机和设备组成的，各自拥有独立的内存。下面介绍一些函数针对GPU内存是的申请，释放和CPU之间传输的接口函数。

标准的C函数	CUDA C函数	标准的C函数	CUDA C函数
malloc	cudaMalloc	memset	cudaMemset
memcpy	cudaMemcpy	free	cudaFree

重点函数 cudaMemcpy 函数原型如下

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

cudaMemcpyKind 有如下类型：

cudaMemcpyHostToHost：从主机内存复制到主机内存（通常用于不同主机内存区域之间的复制，但在CUDA中较少使用，因为更多时候是与设备内存交互）
cudaMemcpyHostToDevice：从主机内存复制到设备内存（即GPU内存）
cudaMemcpyDeviceToHost：从设备内存复制到主机内存
cudaMemcpyDeviceToDevice：从设备内存复制到另一个设备内存（在同一GPU或不同GPU之间复制数据，需要相应的硬件和驱动支持）

函数 cudaGetErrorString 原型如下

const char* cudaGetErrorString(cudaError_t error); //类似C语言中的strerror函数

2.3 线程管理

当核函数在主机端启动，会执行移动到设备上，设备上会产生大量的线程。CUDA明确线程层次抽象的概念，两层的线程层次结构由线程块和线程块网格构成。如下图：
在这里插入图片描述
CUDA组织三维的网格和块。使用dim3类型变量，基于uint3定义的整数型向量，用来表示维度。

blockIdx:块索引
threadIdx:线程索引
gridDim:网格维度

下面是一个打印网格和块索引的与维度的例子：

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * Display the dimensionality of a thread block and grid from the host and
 * device.
 */

__global__ void checkIndex(void)
{
   
   
    printf("threadIdx:(%d, %d, %d) blockIdx:(%d, %d, %d) blockDim:(%d, %d, %d) gridDim:(%d, %d, %d) \n", \
																threadIdx.x, threadIdx.y, threadIdx.z, \
																blockIdx.x, blockIdx.y, blockIdx.z, \
																blockDim.x, blockDim.y, blockDim.z, \
																gridDim.x, gridDim.y, gridDim.z);
}

int main(int argc, char **argv)
{
   
   
    // define total data element
    int nElem = 6;

    // define grid and block structure
    dim3 block(3);
    dim3 grid((nElem + block.x - 1) / block.x);

    // check grid and block dimension from host side
    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);
    printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);

    // check grid and block dimension from device side
    checkIndex<<<grid, block>>>();

    // reset device before you leave
    CHECK(cudaDeviceReset());

    return(0);
}

打印如下，更加形象说明上图的含义

grid.x 2 grid.y 1 grid.z 1
block.x 3 block.y 1 block.z 1
threadIdx:(0, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1) 
threadIdx:(1, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1) 
threadIdx:(2, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1) 
threadIdx:(0, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1) 
threadIdx:(1, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1) 
threadIdx:(2, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)

对于一个给定的数据，确定网格和块尺寸的一般步骤为：

确定块的大小
在已知数据的大小和块大小的基础上计算网格维度

要确定块的尺寸，通常需要考虑如下两点：

内核的性能特性
GPU资源的限制

在后续博客对以上几个因素进行详细的介绍。

2.4 启动一个CUDA核函数

应该对C语言函数的调用语句很熟悉：function_name (argument list); 下面说下CUDA内核调用函数格式，是对C语言函数调用的延伸：kernel_name <<<grid, block>>>(argument list);如下图所示：

核函数调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端，可以用强制主机端程序等待所有核函数执行结束：cudaError_t cudaDeviceSynchronize(void);
或者使用隐式同步，当使用cudaMemcpy函数在主机和设备之间拷贝数据时，主机隐式同步，即主机端必须等待数据拷贝完成后才能继续执行程序。

2.5 编写核函数

核函数是在设备端执行的代码。核函数调用时，书多不同的CUDA线程并行执行同一个计算任务。如下声明定义核函数：__global__ void kernel_name(argument list) 下表总结CUDA程序中的函数类型限定符。

限定符	执行	调用	备注
__global__	在设备端执行	可从主机端调用	必须有一个void返回类型
__device__	在设备段执行	仅能从设备端调用
__host__	在主机端执行	仅能从主机端调用	可以忽略

CUDA核函数的限制：

只能访问设备内存
必须具有void返回类型
不支持可变数量的参数
不支持静态变量
显示异步行为
下面举一个C语言和cuda语言，将两个大小为N的向量 A 和 B 相加，主机端C代码如下：

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

cuda核函数：

__global__ void vectorAddKernel(const float *A, const float *B, float *C, int N) {
   
     
    int i = blockIdx.x * blockDim.x + threadIdx.x;  
    if (i < N) {
   
     
        C[i] = A[i] + B[i];  
    }  
}

从上面两端代码可以判断内置的线程坐标编程替换了数组索引

2.6 编译和执行

现在举例一个完整基于GPU的向量加法例子

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>
void checkResult(float *hostRef, float *gpuRef, const int N)
{
   
   
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
   
   
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
   
   
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at current %d\n", hostRef[i],
                   gpuRef[i], i);
            break;
        }
    }
    if (match) printf("Arrays match.\n\n");
    return;
}
void initialData(float *ip, int size)
{
   
   
    // generate different seed for random number
    time_t t;
    srand((unsigned) time(&t));

    for (int i = 0; i < size; i++)
    {
   
   
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }

    return;
}
void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
   
   
    for (int idx = 0; idx < N; idx++)
        C[idx] = A[idx] + B[idx];
}
__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
   
   
    int i = threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}
int main(int argc, char **argv)
{
   
   
    printf("%s Starting...\n", argv[0]);
    // set up device
    int dev = 0;
    CHECK(cudaSetDevice(dev));
    // set up data size of vectors
    int nElem = 1 << 5;
    printf("Vector size %d\n", nElem);
    // malloc host memory
    size_t nBytes = nElem * sizeof(float);
    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float *)malloc(nBytes);
    h_B     = (float *)malloc(nBytes);
    hostRef = (float *)malloc(nBytes);
    gpuRef  = (float *)malloc(nBytes);
    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);
    // malloc device global memory
    float *d_A, *d_B, *d_C;
    CHECK(cudaMalloc((float**)&d_A, nBytes));
    CHECK(cudaMalloc((float**)&d_B, nBytes));
    CHECK(cudaMalloc((float**)&d_C, nBytes));
    // transfer data from host to device
    CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_C, gpuRef, nBytes, cudaMemcpyHostToDevice));
    // invoke kernel at host side
    dim3 block (nElem);
    dim3 grid  (1);

    sumArraysOnGPU<<<grid, block>>>(d_A, d_B, d_C, nElem);
    printf("Execution configure <<<%d, %d>>>\n", grid.x, block.x);
    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));
    // add vector at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);
    // check device results
    checkResult(hostRef, gpuRef, nElem);
    // free device global memory
    CHECK(cudaFree(d_A));
    CHECK(cudaFree(d_B));
    CHECK(cudaFree(d_C));
    // free host memory
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);
    CHECK(cudaDeviceReset());
    return(0);

Makefile文件

# Location of the CUDA Toolkit
CUDA_PATH ?= /usr/local/cuda
TARGET_ARCH ?= x86_64
# architecture
HOST_ARCH   := $(shell uname -m)
TARGET_ARCH ?= x86_64
TARGET_SIZE := 64
TARGET_OS ?= linux
HOST_COMPILER ?= g++
NVCC          := $(CUDA_PATH)/bin/nvcc -ccbin $(HOST_COMPILER)
# internal flags
NVCCFLAGS   := -m${TARGET_SIZE}
CCFLAGS     :=
LDFLAGS     :=
# Debug build flags
ifeq ($(dbg),1)
      NVCCFLAGS += -g -G
      BUILD_TYPE := debug
else
      BUILD_TYPE := release
endif

ALL_CCFLAGS :=
ALL_CCFLAGS += $(NVCCFLAGS)
ALL_CCFLAGS += $(EXTRA_NVCCFLAGS)
ALL_CCFLAGS += $(addprefix -Xcompiler ,$(CCFLAGS))
ALL_CCFLAGS += $(addprefix -Xcompiler ,$(EXTRA_CCFLAGS))
SAMPLE_ENABLED := 1
ALL_LDFLAGS :=
ALL_LDFLAGS += $(ALL_CCFLAGS)
ALL_LDFLAGS += $(addprefix -Xlinker ,$(LDFLAGS))
ALL_LDFLAGS += $(addprefix -Xlinker ,$(EXTRA_LDFLAGS))
# Common includes and paths for CUDA
INCLUDES  := -I./common
LIBRARIES :=
################################################################################