【CUDA 】第2章 CUDA编程模型——2.1概述

最新推荐文章于 2024-09-26 17:12:37 发布

进击的一米六

最新推荐文章于 2024-09-26 17:12:37 发布

阅读量376

点赞数 6

分类专栏： # CUDA C编程权威指南文章标签：学习 c语言 c++

本文链接：https://blog.csdn.net/weixin_44231807/article/details/142377698

版权

CUDA C编程权威指南专栏收录该内容

2 篇文章 0 订阅

订阅专栏

CUDA C编程笔记

第二章 CUDA编程模型
- 2.1 概述

第二章 CUDA编程模型

2.1 概述

编程模型主要内容：1、GPU线程组织层次结构 2、GPU内存访问层次结构

2.1.1 编程结构

异步：GPU上运算可以和通信重叠
一旦内核被调用，控制权还给主机，GPU在计算核函数的同时，主机可以执行其它函数

2.1.2 内存管理函数

cudaMemcpy（目的，源，多少字节，方向）：主机和设备的数据传输
同步执行：函数返回和传输完成前，主机阻塞
cudaGetErrorString（调用函数返回的错误代码）：查询错误代码，转为可读的信息

抽象出GPU内存层次结构：全局（类似CPU系统内存）和共享（类似CPU缓存，但能由核函数控制）

下面是CPU端两个数组相加的代码：

#include <stdlib.h>
#include <string.h>
#include <time.h>
//浮点数数组加法

//数组加函数
void sumArrayOnHost(float *A,float *B,float *C,const int size){
    for(int i = 0; i < size; i++){
        C[i] = A[i] + B[i];
    }
}
//填数组函数-随机生成
void initialData(float *ip, int size){
    time_t t;//time_t是时间数据类型
    srand((unsigned int) time(&t));//srand用当前时间内存地址生成随机数种子，rand生成随机数

    for(int i = 0; i<size; i++){
        ip[i] = (float) (rand() & 0xFF)/10.0f;//rand生成随机数，与0xFF（0-255）按位与，得到一个0-255的整数，把这个数强转为float类型，然后除浮点数10.0，得到范围在0-25.5的随机浮点数
    }
}
//主函数
int main(int argc,int **argv){//argc表示命令行参数的个数（空白符分隔），argv存储了所有命令行参数
    int nElem;
    size_t nBytes = nElem * sizeof(float);//size_t是无符号整数，表示任何对象能表示的最大长度，多用于数组下标和内存管理
    float *h_A, *h_B, *h_C;
    h_A = (float *)malloc(nBytes);
    h_B = (float *)malloc(nBytes);
    h_C = (float *)malloc(nBytes);

    initialData(h_A, nElem);
    initialData(h_B, nElem);

    sumArrayOnHost(h_A, h_B, h_C, nElem);

    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

在此基础上修改代码编程来在GPU上进行数组加法：

2.1.3 线程管理

一个核函数——启动——>产生的所有线程【网格】（共享全局内存）——>多个块
在这里插入图片描述
块索引：blockIdx
线程索引：threadIdx
网格维度：girdDim
块的维度：blockDim
都是dim3类型的变量，都xyz三个维度，未指定元素初始化为1

###代码2-2遇到的问题

$ nvcc -arch=sm_86 2-2.1checkDimension.cu -o hello
2-2.1checkDimension.cu(5): error: identifier "z" is undefined

2-2.1checkDimension.cu(5): error: identifier "y" is undefined

2-2.1checkDimension.cu(5): warning #1290-D: a class type that is not trivially copyable passed through ellipsis

2-2.1checkDimension.cu(5): warning #181-D: argument is incompatible with corresponding format string conversion

2-2.1checkDimension.cu(5): warning #181-D: argument is incompatible with corresponding format string conversion

2-2.1checkDimension.cu(5): warning #225-D: the format string ends before this argument

分析：dim3 block (3);block在x的维度为3，没有显示声明的yz维度应该均为默认值1，但是这里报错，说明没有默认为1
原因：代码错误，把blockIdx.z的点打成了逗号，

正确代码2-2

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

__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){
    int nElem = 6;
    dim3 block (3);//block在x的维度为3
    dim3 grid ((nElem + block.x - 1)/block.x);//????为什么两处都是block.x不是blockdim

    //主机端检查block和grid维度
    printf("grid.x %d  grid.y %d  grid.z %d\n",grid.x,grid.y,grid.z);//只设置了x维度，没有设置yz维度报错第5行对yz的访问，为什么没有默认初始化为1
    printf("block.x %d  block.y %d  block.z %d\n",block.x,block.y,block.z);

    //设备端检查block和grid维度
    checkIndex<<<grid,block>>> ();

    //离开前重置设备
    cudaDeviceReset();
    return (0);//括号有啥用？？？
}

###代码2-3遇到的问题

2-3.1defineGridBlock.cu(12): error: "block" has already been declared in the current scope
2-3.1defineGridBlock.cu(14): error: "grid" has already been declared in the current scope

原因：第一次定义block和grid的值，第二次开始修改它们的值时，应该用修改值的语句，而不是定义的语句

正确代码2-3

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

int main(int argc,char **argv){
    int nElem=1024;

    dim3 block(1024);
    dim3 grid((nElem+block.x-1)/block.x);
    printf("grid.x %d block.x %d \n",grid.x,block.x);

    //重置块数
    //dim3 block(512);block以及被声明过了，不能再用定义的语句，用修改值的语句
    block.x = 512;
    //dim3 grid((nElem+block.x-1)/block.x);grid同理，也被声明过了
    grid.x = (nElem+block.x-1)/block.x;
    printf("grid.x %d block.x %d \n",grid.x,block.x);

    block.x = 256;
    grid.x = (nElem+block.x-1)/block.x;
    printf("grid.x %d block.x %d \n",grid.x,block.x);

    block.x = 128;
    grid.x = (nElem+block.x-1)/block.x;
    printf("grid.x %d block.x %d \n",grid.x,block.x);

    //离开前重置设备
    cudaDeviceReset();
    return (0);
}

2.1.4 启动CUDA核函数

同一个块内的线程能协作，不同块的线程不能协作
kernal_name<<<一个网格几个块，一个块几个线程>>> (参数列表)；
显式同步：cudaDeviceSynchronize(void)强制同步，让主机端等所有核函数执行结束。
隐式同步：cudaMemcpy(目的，源，数量，方向)，让主机端等待数据拷贝完后才能继续执行
核函数：在设备端执行的代码，__global__修饰，返回类型为void

2.1.8 执行和编译

###代码2-4遇到的问题

2-4.1sumArraysOnGPU-small-case.cu(6): error: identifier "cudaERROR_t" is undefined

解决：把宏定义错误处理那段注释掉，大概是缺少头文件

正确代码2-4

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

//处理错误，用宏封装错误调用
// #define CHECK(call){
//     const cudaERROR_t error = call;
//     if(error != cudaSuccess){
//         printf("Error:%s:%d,",__FILE__,__LINE__);
//         printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));
//         exit(1);
//     }
// }

//验证核函数，判断是否正确运行
//CPU和GPU浮点数组匹配比较
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");
}

void initialData(float *ip,int size){
    time_t t;
    srand((unsigned) time(&t));

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

//主机上数组加
void sumArraysOhHost(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){
    int i = threadIdx.x;
    C[i] = A[i] + B[i]; 
}

int main(int argc,char **argv){
    printf("%s Starting...\n",argv[0]);

    //建立设备
    int dev = 0;
    cudaSetDevice(dev);

    //建立容器的数据尺寸
    int nElem = 32;//向量大小32
    printf("Vector size %d\n",nElem);

    //分配主机内存
    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);

    //在主机端初始化数据
    initialData(h_A,nElem);
    initialData(h_B,nElem);

    memset(hostRef,0,nBytes);
    memset(gpuRef,0,nBytes);

    //分配设备global内存
    float *d_A,*d_B,*d_C;
    cudaMalloc((float**)&d_A,nBytes);//第一个参数要求为void**
    cudaMalloc((float**)&d_B,nBytes);
    cudaMalloc((float**)&d_C,nBytes);

    //从主机向设备传输数据
    cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice);
    cudaMemcpy(d_B,h_B,nBytes,cudaMemcpyHostToDevice);

    //主机端激活函数
    //这里所有都放入一个块内，一个块32个元素
    dim3 block(nElem);
    dim3 grid(nElem/block.x);//？？？和之前的不一样，之前的是(nElem+block.x-1)/block.x;
    // //版本2：重新配置块的大小，导致代码2-4中核函数也修改int i=blockIdx.x;
    // dim3 block(1);
    // dim3 grid(nElem);

    sumArraysOnGPU<<<grid,block>>>(d_A,d_B,d_C);
    printf("Execution configuration <<<%d,%d>>>\n",grid.x,block.x);

    //把gpu结果复制回主机端
    cudaMemcpy(gpuRef,d_C,nBytes,cudaMemcpyDeviceToHost);

    //主机端计算结果
    sumArraysOhHost(h_A,h_B,hostRef,nElem);

    //检查结果
    checkResult(hostRef,gpuRef,nElem);

    //释放设备global 存储
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    //释放主机内存
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return(0);
}