CUDA编程基础

文章目录

  • 1、GPU介绍
  • 2、CUDA程序进行编译
  • 3、CUDA线程模型
    • 3.1、一维网格一维线程块
    • 3.2、二维网格二维线程块
    • 3.3、三维网格三维线程块
    • 3.3、不同组合形式
  • 4、nvcc编译流程
  • 5、CUDA程序基本架构
  • 6、错误检测函数
    • 6.1、运行时API错误代码
    • 6.2、检查核函数
  • 7、CUDA记时
    • 7.1、记时代码
    • 7.2、核函数记时实例
    • 7.3、nvprof性能刨析
    • 7.4、运行时API查询GPU信息
    • 7.5、查询GPU计算核心数量
  • 8、组织线程模型
    • 8.1、一维网格一维线程块计算二维矩阵加法
    • 8.2、二维网格一维线程块计算二维矩阵加法
    • 8.3、二维网格二维线程块计算二维矩阵加法
  • 9、内存结构

1、GPU介绍

参考链接

GPU 意为图形处理器，也常被称为显卡，GPU最早主要是进行图形处理的。如今深度学习大火，GPU高效的并行计算能力充分被发掘，GPU在AI应用上大放异彩。GPU拥有更多的运算核心，其特别适合数据并行的计算密集型任务，如大型矩阵运算，与GPU对应的一个概念是CPU，但CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务，CPU更擅长数据缓存和流程控制。

1、GPU不能单独进行工作，GPU相当于CPU的协处理器，由CPU进行调度。CPU+GPU组成异构计算架构，CPU的特点是更加擅长逻辑处理，而对大量数据的运算就不是那么擅长了，GPU恰好相反，GPU可以并行处理大量数据运算。
在这里插入图片描述

2、CUDA运行时API

CUDA提供两层API接口，CUDA驱动(driver)API和CUDA运行时(runtime)API；
两种API调用性能几乎无差异，课程使用操作对用户更加友好Runtime API；

在这里插入图片描述

3、第一个CUDA程序

#include <stdio.h>__global__ void hello_from_gpu()
{printf("Hello World from the the GPU\n");
}int main(void)
{hello_from_gpu<<<4, 4>>>();cudaDeviceSynchronize();return 0;
}

2、CUDA程序进行编译

通过nvidia-smi 查看当前显卡信息

在这里插入图片描述

使用nvcc对cuda代码进行编译
nvcc test1.cu -o test1

// 1、
// 核函数 在GPU上进行并执行
// 注意：限定词__global__
// 返回值必须是void// 两种都是正确的
// 形式1:__global__ void
// 形式2:__global__ void__global__ void hello_from_gpu()
{const int bid = blockIdx.x;const int tid = threadIdx.x;const int id = threadIdx.x + blockIdx.x * blockDim.x;printf("hello world from the GPU block:%d and thread:%d,global id:%d\n",bid,tid,id);
}int main()
{// 指定线程模型// 第一个指的是线程块的个数，第二个指的每个线程块线程的数量hello_from_gpu<<<4,4>>>();// 因为GPU是CPU的协调处理器，所以需要处理主机与设备直接的同步cudaDeviceSynchronize();return 0;
}

在这里插入图片描述

3、CUDA线程模型

3.1、一维网格一维线程块

当一个核函数在主机中启动时，他所有的线程构成了一个网格（grid）包含多个线程块（block）包含多个线程，线程是GPU中最小单位
<<<4,4>>> 表示gird中线程块的个数，block中线程数

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

#include <stdio.h>__global__ void hello_from_gpu()
{const int bid = blockIdx.x;const int tid = threadIdx.x;const int id = threadIdx.x + blockIdx.x * blockDim.x; printf("Hello World from block %d and thread %d, global id %d\n", bid, tid, id);
}int main(void)
{printf("Hello World from CPU!\n");hello_from_gpu<<<2, 2>>>();cudaDeviceSynchronize();return 0;
}

3.2、二维网格二维线程块

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

3.3、三维网格三维线程块

在这里插入图片描述

3.3、不同组合形式

在这里插入图片描述

在这里插入图片描述
在这里插入图片描述

4、nvcc编译流程

nvcc编译流程：需要注意的是，GPU的真实架构能力需要大于虚拟架构能力。

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

5、CUDA程序基本架构

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

在这里插入图片描述
在这里插入图片描述
使用GPU进行矩阵计算

#include <stdio.h>
int setGPU()
{int idevcount = 0;cudaError_t error = cudaGetDeviceCount(&idevcount);if(error != cudaSuccess || error == 0){printf("No found GPU\n");exit(-1);}else{printf("The count of GPU is :%d.\n",idevcount);}// 设置执行GPUint idev = 0;error = cudaSetDevice(idev);if(error != cudaSuccess){printf("fail set device 0 GPU\n");exit(-1);}else{printf("set GPU 0 for computing\n");}return 0;
}// 初始化函数
void initdata(float *addr,int element)
{for(int i = 0;i<element;i++){addr[i] = (float)(rand() & 0xFF) / 10.f;}return;
}// 使用设备函数
__device__ float add(float a, float b)
{return a+b;
}
// 核函数
__global__ void addFromGPU(float *a, float *b, float *c, const int n)
{const int bid = blockIdx.x;const int tid = threadIdx.x;const int id = tid + bid * blockDim.x;// 当513的时候，32*17 =544个线程，所以需要限制一下
//    c[id] = a[id] + b[id];if(id>n) return;c[id] = add(a[id]+b[id]);
}int main1()
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int ielement = 513; // 设置元素个数size_t stBytescount = ielement * sizeof(float); // 字节数// 分配主机内存，并初始化float *fphost_a,*fphost_b,*fphost_c;fphost_a = (float*)malloc(stBytescount);fphost_b = (float*)malloc(stBytescount);fphost_c = (float*)malloc(stBytescount);if(fphost_a != NULL && fphost_b != NULL && fphost_c != NULL){memset(fphost_a,0x00,stBytescount);memset(fphost_b,0x00,stBytescount);memset(fphost_c,0x00,stBytescount);}else{printf("fail to allocate host memory\n");exit(-1);}// 分配设备内存float *fpdevice_a,*fpdevice_b,*fpdevice_c;cudaMalloc((float**)&fpdevice_a,stBytescount);cudaMalloc((float**)&fpdevice_b,stBytescount);cudaMalloc((float**)&fpdevice_c,stBytescount);if(fpdevice_a != NULL && fpdevice_b != NULL && fpdevice_c != NULL){cudaMemset(fpdevice_a,0,stBytescount);cudaMemset(fpdevice_b,0,stBytescount);cudaMemset(fpdevice_c,0,stBytescount);}else{printf("fail to allocate device memory\n");free(fphost_a);free(fphost_b);free(fphost_c);exit(-1);}// 初始化随即种子srand(666);initdata(fphost_a,ielement);initdata(fphost_b,ielement);// 数据从主机中拷贝到设备中cudaMemcpy(fpdevice_a,fphost_a,stBytescount,cudaMemcpyHostToDevice);cudaMemcpy(fpdevice_b,fphost_b,stBytescount,cudaMemcpyHostToDevice);cudaMemcpy(fpdevice_c,fphost_c,stBytescount,cudaMemcpyHostToDevice);// 掉用核函数在设备中进行计算dim3 block(32);dim3 grid(ielement/32);// 掉用核函数addFromGPU<<<grid,block>>>(fpdevice_a,fpdevice_b,fpdevice_c,ielement);cudaDeviceSynchronize();// 将计算的到的数据从设备拷贝到主机中(隐式同步)cudaMemcpy(fphost_c,fpdevice_c,stBytescount,cudaMemcpyDeviceToHost);for(int i=0;i<10;i++){printf("idx=%2d\tmatrix_a:%.2f\tmatrix_b:%.2f\tresult=%.2f\n",fphost_a[i],fphost_b[i],fphost_c[i]);}// 释放内存free(fphost_a);free(fphost_b);free(fphost_c);cudaFree(fpdevice_a);cudaFree(fpdevice_b);cudaFree(fpdevice_c);cudaDeviceReset();return 0;
}

6、错误检测函数

6.1、运行时API错误代码

CUDA运行时API大多支持返回错误代码，返回值类型为cudaError_t，前面的例子我们也已经讲解过，CUDA运行时API成功执行，返回的错误代码为cudaSuccess，运行时API返回的执行状态值是枚举变量。

在这里插入图片描述
在这里插入图片描述

#pragma once
#include <stdlib.h>
#include <stdio.h>cudaError_t ErrorCheck(cudaError_t error_code, const char* filename, int lineNumber)
{if (error_code != cudaSuccess){printf("CUDA error:\r\ncode=%d, name=%s, description=%s\r\nfile=%s, line%d\r\n",error_code, cudaGetErrorName(error_code), cudaGetErrorString(error_code), filename, lineNumber);return error_code;}return error_code;
}

调用错误检测函数：

cudaError_t error = ErrorCheck(cudaSetDevice(iDev), FILE, LINE);

6.2、检查核函数

错误检查函数无法捕捉调用核函数时发生的相关错误，前面也讲到过，核函数的返回值类型时void，即核函数不返回任何值。可以通过在调用核函数之后调用**cudaGetLastError()**函数捕捉核函数错误。

获取cuda程序的最后一个错误—cudaGetLastError
在这里插入图片描述

在调用核函数后，追加如下代码：

ErrorCheck(cudaGetLastError(), __FILE__, __LINE__);
ErrorCheck(cudaDeviceSynchronize(), __FILE__, __LINE__);

7、CUDA记时

通常情况下不仅要关注程序的正确性，还要关注程序的性能（即执行速度）。了解核函数的执行需要多长时间是很有必要的，想要掌握程序的性能，就需要对程序进行精确的记时。

7.1、记时代码

CUDA事件记时代码如下，只需要将需要记时的代码嵌入记时代码之间：

cudaEvent_t start, stop;
ErrorCheck(cudaEventCreate(&start), __FILE__, __LINE__);
ErrorCheck(cudaEventCreate(&stop), __FILE__, __LINE__);
ErrorCheck(cudaEventRecord(start), __FILE__, __LINE__);
cudaEventQuery(start);	//此处不可用错误检测函数/************************************************************
需要记时间的代码
************************************************************/ErrorCheck(cudaEventRecord(stop), __FILE__, __LINE__);
ErrorCheck(cudaEventSynchronize(stop), __FILE__, __LINE__);
float elapsed_time;
ErrorCheck(cudaEventElapsedTime(&elapsed_time, start, stop), __FILE__, __LINE__);
printf("Time = %g ms.\n", elapsed_time);ErrorCheck(cudaEventDestroy(start), __FILE__, __LINE__);
ErrorCheck(cudaEventDestroy(stop), __FILE__, __LINE__);

代码解析：第1行cudaEvent_t start, stop：定义两个CUDA事件类型（cudaEvent_t）的变量；
第2、3行cudaEventCreate函数初始化定义的cudaEvent_t变量；
第4行通过cudaEventRecord函数，在需要记时的代码块之前记录代表时间开始的事件；
第5行cudaEventQuery函数在TCC驱动模式的GPU下可省略，但在处于WDDM驱动模式的GPU必须保留，因此，我们就一直保留这句函数即可。注意：cudaEventQuery函数不可使用错误检测函数；
第8行是需要记时的代码块；
第11行在需要记时的代码块之后记录代表时间结束的事件；
第12行cudaEventSynchronize函数作用是让主机等待事件stop被记录完毕；
第13~15行cudaEventElapsedTime函数的调用作用是计算cudaEvent_t变量start和stop时间差，记录在float变量elapsed_time中，并输出打印到屏幕上；
第17、18行调用cudaEventDestroy函数销毁start和stop这两个类型为cudaEvent_t的CUDA事件。

7.2、核函数记时实例

此代码计算运行核函数10次的平均时间，核函数实际运行11次，由于第一次调用核函数，往往会花费更多的时间，如果将第一次记录进去，可能导致记录的时间不准确，因此忽略第一次调用核函数的时间，取10次平均值。

#include <stdio.h>
#include "../tools/common.cuh"#define NUM_REPEATS 10__device__ float add(const float x, const float y)
{return x + y;
}__global__ void addFromGPU(float *A, float *B, float *C, const int N)
{const int bid = blockIdx.x;const int tid = threadIdx.x;const int id = tid + bid * blockDim.x; if (id >= N) return;C[id] = add(A[id], B[id]);}void initialData(float *addr, int elemCount)
{for (int i = 0; i < elemCount; i++){addr[i] = (float)(rand() & 0xFF) / 10.f;}return;
}int main(void)
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int iElemCount = 4096;                     // 设置元素数量size_t stBytesCount = iElemCount * sizeof(float); // 字节数// （1）分配主机内存，并初始化float *fpHost_A, *fpHost_B, *fpHost_C;fpHost_A = (float *)malloc(stBytesCount);fpHost_B = (float *)malloc(stBytesCount);fpHost_C = (float *)malloc(stBytesCount);if (fpHost_A != NULL && fpHost_B != NULL && fpHost_C != NULL){memset(fpHost_A, 0, stBytesCount);  // 主机内存初始化为0memset(fpHost_B, 0, stBytesCount);memset(fpHost_C, 0, stBytesCount);}else{printf("Fail to allocate host memory!\n");exit(-1);}// （2）分配设备内存，并初始化float *fpDevice_A, *fpDevice_B, *fpDevice_C;ErrorCheck(cudaMalloc((float**)&fpDevice_A, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((float**)&fpDevice_B, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((float**)&fpDevice_C, stBytesCount), __FILE__, __LINE__);if (fpDevice_A != NULL && fpDevice_B != NULL && fpDevice_C != NULL){ErrorCheck(cudaMemset(fpDevice_A, 0, stBytesCount), __FILE__, __LINE__); // 设备内存初始化为0ErrorCheck(cudaMemset(fpDevice_B, 0, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMemset(fpDevice_C, 0, stBytesCount), __FILE__, __LINE__);}else{printf("fail to allocate memory\n");free(fpHost_A);free(fpHost_B);free(fpHost_C);exit(-1);}// 3、初始化主机中数据srand(666); // 设置随机种子initialData(fpHost_A, iElemCount);initialData(fpHost_B, iElemCount);// 4、数据从主机复制到设备ErrorCheck(cudaMemcpy(fpDevice_A, fpHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(fpDevice_B, fpHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(fpDevice_C, fpHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);// 5、调用核函数在设备中进行计算dim3 block(32);dim3 grid((iElemCount + block.x - 1) / 32);float t_sum = 0;for (int repeat = 0; repeat <= NUM_REPEATS; ++repeat){cudaEvent_t start, stop;ErrorCheck(cudaEventCreate(&start), __FILE__, __LINE__);ErrorCheck(cudaEventCreate(&stop), __FILE__, __LINE__);ErrorCheck(cudaEventRecord(start), __FILE__, __LINE__);cudaEventQuery(start);	//此处不可用错误检测函数addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElemCount);    // 调用核函数ErrorCheck(cudaEventRecord(stop), __FILE__, __LINE__);ErrorCheck(cudaEventSynchronize(stop), __FILE__, __LINE__);float elapsed_time;ErrorCheck(cudaEventElapsedTime(&elapsed_time, start, stop), __FILE__, __LINE__);// printf("Time = %g ms.\n", elapsed_time);if (repeat > 0){t_sum += elapsed_time;}ErrorCheck(cudaEventDestroy(start), __FILE__, __LINE__);ErrorCheck(cudaEventDestroy(stop), __FILE__, __LINE__);}const float t_ave = t_sum / NUM_REPEATS;printf("Time = %g ms.\n", t_ave);// 6、将计算得到的数据从设备传给主机ErrorCheck(cudaMemcpy(fpHost_C, fpDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__);// 7、释放主机与设备内存free(fpHost_A);free(fpHost_B);free(fpHost_C);ErrorCheck(cudaFree(fpDevice_A), __FILE__, __LINE__);ErrorCheck(cudaFree(fpDevice_B), __FILE__, __LINE__);ErrorCheck(cudaFree(fpDevice_C), __FILE__, __LINE__);ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);return 0;
}

7.3、nvprof性能刨析

1、nvprof工具说明
CUDA 5.0后有一个工具叫做nvprof的命令行分析工具，nvprof是一个可执行文件。

如下执行命令语句，其中exe_name为可执行文件的名字。

nvprof ./exe_name


#include <stdio.h>
#include "../tools/common.cuh"#define NUM_REPEATS 10__device__ float add(const float x, const float y)
{return x + y;
}__global__ void addFromGPU(float *A, float *B, float *C, const int N)
{const int bid = blockIdx.x;const int tid = threadIdx.x;const int id = tid + bid * blockDim.x; if (id >= N) return;C[id] = add(A[id], B[id]);}void initialData(float *addr, int elemCount)
{for (int i = 0; i < elemCount; i++){addr[i] = (float)(rand() & 0xFF) / 10.f;}return;
}int main(void)
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int iElemCount = 4096;                     // 设置元素数量size_t stBytesCount = iElemCount * sizeof(float); // 字节数// （1）分配主机内存，并初始化float *fpHost_A, *fpHost_B, *fpHost_C;fpHost_A = (float *)malloc(stBytesCount);fpHost_B = (float *)malloc(stBytesCount);fpHost_C = (float *)malloc(stBytesCount);if (fpHost_A != NULL && fpHost_B != NULL && fpHost_C != NULL){memset(fpHost_A, 0, stBytesCount);  // 主机内存初始化为0memset(fpHost_B, 0, stBytesCount);memset(fpHost_C, 0, stBytesCount);}else{printf("Fail to allocate host memory!\n");exit(-1);}// （2）分配设备内存，并初始化float *fpDevice_A, *fpDevice_B, *fpDevice_C;ErrorCheck(cudaMalloc((float**)&fpDevice_A, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((float**)&fpDevice_B, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMalloc((float**)&fpDevice_C, stBytesCount), __FILE__, __LINE__);if (fpDevice_A != NULL && fpDevice_B != NULL && fpDevice_C != NULL){ErrorCheck(cudaMemset(fpDevice_A, 0, stBytesCount), __FILE__, __LINE__); // 设备内存初始化为0ErrorCheck(cudaMemset(fpDevice_B, 0, stBytesCount), __FILE__, __LINE__);ErrorCheck(cudaMemset(fpDevice_C, 0, stBytesCount), __FILE__, __LINE__);}else{printf("fail to allocate memory\n");free(fpHost_A);free(fpHost_B);free(fpHost_C);exit(-1);}// 3、初始化主机中数据srand(666); // 设置随机种子initialData(fpHost_A, iElemCount);initialData(fpHost_B, iElemCount);// 4、数据从主机复制到设备ErrorCheck(cudaMemcpy(fpDevice_A, fpHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(fpDevice_B, fpHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);ErrorCheck(cudaMemcpy(fpDevice_C, fpHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__);// 5、调用核函数在设备中进行计算dim3 block(32);dim3 grid((iElemCount + block.x - 1) / 32);addFromGPU<<<grid, block>>>(fpDevice_A, fpDevice_B, fpDevice_C, iElemCount);    // 调用核函数// 6、将计算得到的数据从设备传给主机ErrorCheck(cudaMemcpy(fpHost_C, fpDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__);// 7、释放主机与设备内存free(fpHost_A);free(fpHost_B);free(fpHost_C);ErrorCheck(cudaFree(fpDevice_A), __FILE__, __LINE__);ErrorCheck(cudaFree(fpDevice_B), __FILE__, __LINE__);ErrorCheck(cudaFree(fpDevice_C), __FILE__, __LINE__);ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__);return 0;
}

nvprof ./nvprofAnalysis主要看GPU activities：[CUDA memcpy HtoD]：主机向设备拷贝数据花费时间占比44.98%；
[CUDA memset]：设备调用cudaMemset函数初始化数据占用时间占比23.76%；
核函数执行占比为18.81%；
[CUDA memcpy DtoH]：设备向主机拷贝数据花费时间占比12.14%；

在这里插入图片描述

7.4、运行时API查询GPU信息

在这里插入图片描述


#include "../tools/common.cuh"
#include <stdio.h>int main(void)
{int device_id = 0;ErrorCheck(cudaSetDevice(device_id), __FILE__, __LINE__);cudaDeviceProp prop;ErrorCheck(cudaGetDeviceProperties(&prop, device_id), __FILE__, __LINE__);printf("Device id:                                 %d\n",device_id);printf("Device name:                               %s\n",prop.name);printf("Compute capability:                        %d.%d\n",prop.major, prop.minor);printf("Amount of global memory:                   %g GB\n",prop.totalGlobalMem / (1024.0 * 1024 * 1024));printf("Amount of constant memory:                 %g KB\n",prop.totalConstMem  / 1024.0);printf("Maximum grid size:                         %d %d %d\n",prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);printf("Maximum block size:                        %d %d %d\n",prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);printf("Number of SMs:                             %d\n",prop.multiProcessorCount);printf("Maximum amount of shared memory per block: %g KB\n",prop.sharedMemPerBlock / 1024.0);printf("Maximum amount of shared memory per SM:    %g KB\n",prop.sharedMemPerMultiprocessor / 1024.0);printf("Maximum number of registers per block:     %d K\n",prop.regsPerBlock / 1024);printf("Maximum number of registers per SM:        %d K\n",prop.regsPerMultiprocessor / 1024);printf("Maximum number of threads per block:       %d\n",prop.maxThreadsPerBlock);printf("Maximum number of threads per SM:          %d\n",prop.maxThreadsPerMultiProcessor);return 0;
}

在这里插入图片描述

说明：Device id: 计算机中GPU的设备代号，我只有一个显卡，所以只能是0；
Device name: 显卡名字，我的显卡是Quadro P620；
Compute capability: GPU计算能力，我的主版本是6，次版本是1；
Amount of global memory: 显卡显存大小，我的是4G的显存；
Amount of constant memory: 常量内存大小；
Maximum grid size: 最大网格大小（三个维度分别的最大值）；
Maximum block size: 最大线程块大小（三个维度分别的最大值）；
Number of SMs: 流多处理器数量；
Maximum amount of shared memory per block: 每个线程块最大共享内存数量；
Maximum amount of shared memory per SM: 每个流多处理器最大共享内存数量；
Maximum number of registers per block: 每个线程块最大寄存器内存数量；
Maximum number of registers per SM: 每个流多处理器最大寄存器内存数量；
Maximum number of threads per block: 每个线程块最大的线程数量；
Maximum number of threads per SM: 每个流多处理器最大的线程数量。

7.5、查询GPU计算核心数量

CUDA运行时API函数是无法查询GPU的核心数量的

#include <stdio.h>
#include "../tools/common.cuh"int getSPcores(cudaDeviceProp devProp)
{  int cores = 0;int mp = devProp.multiProcessorCount;switch (devProp.major){case 2: // Fermiif (devProp.minor == 1) cores = mp * 48;else cores = mp * 32;break;case 3: // Keplercores = mp * 192;break;case 5: // Maxwellcores = mp * 128;break;case 6: // Pascalif ((devProp.minor == 1) || (devProp.minor == 2)) cores = mp * 128;else if (devProp.minor == 0) cores = mp * 64;else printf("Unknown device type\n");break;case 7: // Volta and Turingif ((devProp.minor == 0) || (devProp.minor == 5)) cores = mp * 64;else printf("Unknown device type\n");break;case 8: // Ampereif (devProp.minor == 0) cores = mp * 64;else if (devProp.minor == 6) cores = mp * 128;else if (devProp.minor == 9) cores = mp * 128; // ada lovelaceelse printf("Unknown device type\n");break;case 9: // Hopperif (devProp.minor == 0) cores = mp * 128;else printf("Unknown device type\n");break;default:printf("Unknown device type\n"); break;}return cores;
}int main()
{int device_id = 0;ErrorCheck(cudaSetDevice(device_id), __FILE__, __LINE__);cudaDeviceProp prop;ErrorCheck(cudaGetDeviceProperties(&prop, device_id), __FILE__, __LINE__);printf("Compute cores is %d.\n", getSPcores(prop));return 0;
}

8、组织线程模型

在这里插入图片描述
在这里插入图片描述

8.1、一维网格一维线程块计算二维矩阵加法

#include <stdio.h>
#include "../tools/common.cuh"__global__ void addMatrix(int *A, int *B, int *C, const int nx, const int ny)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;if (ix < nx){for (int iy = 0; iy < ny; iy++){int idx = iy * nx + ix;C[idx] = A[idx] + B[idx];}}
}int main(void)
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int nx = 16;int ny = 8;int nxy = nx * ny;size_t stBytesCount = nxy * sizeof(int);// （1）分配主机内存，并初始化int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount);ipHost_B = (int *)malloc(stBytesCount);ipHost_C = (int *)malloc(stBytesCount);if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL){for (int i = 0; i < nxy; i++){ipHost_A[i] = i;ipHost_B[i] = i + 1;}memset(ipHost_C, 0, stBytesCount); }else{printf("Fail to allocate host memory!\n");exit(-1);}// （2）分配设备内存，并初始化int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int**)&ipDevice_A, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_B, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_C, stBytesCount), __FILE__, __LINE__); if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL){ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); }   else{printf("Fail to allocate memory\n");free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(1);}// calculate on GPUdim3 block(4, 1);dim3 grid((nx + block.x -1) / block.x, 1);printf("Thread config:grid:<%d, %d>, block:<%d, %d>\n", grid.x, grid.y, block.x, block.y);addMatrix<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);  // 调用内核函数ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__); for (int i = 0; i < 10; i++){printf("id=%d, matrix_A=%d, matrix_B=%d, result=%d\n", i + 1,ipHost_A[i], ipHost_B[i], ipHost_C[i]);}free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__); ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__); return 0;
}

8.2、二维网格一维线程块计算二维矩阵加法

#include <stdio.h>
#include "../tools/common.cuh"__global__ void addMatrix(int *A, int *B, int *C, const int nx, const int ny)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = blockIdx.y;unsigned int idx = iy * nx + ix;if (ix < nx && iy < ny){C[idx] = A[idx] + B[idx];}
}int main(void)
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int nx = 16;int ny = 8;int nxy = nx * ny;size_t stBytesCount = nxy * sizeof(int);// （1）分配主机内存，并初始化int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount);ipHost_B = (int *)malloc(stBytesCount);ipHost_C = (int *)malloc(stBytesCount);if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL){for (int i = 0; i < nxy; i++){ipHost_A[i] = i;ipHost_B[i] = i + 1;}memset(ipHost_C, 0, stBytesCount); }else{printf("Fail to allocate host memory!\n");exit(-1);}// （2）分配设备内存，并初始化int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int**)&ipDevice_A, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_B, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_C, stBytesCount), __FILE__, __LINE__); if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL){ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); }   else{printf("Fail to allocate memory\n");free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(1);}// calculate on GPUdim3 block(4, 1);dim3 grid((nx + block.x -1) / block.x, ny);printf("Thread config:grid:<%d, %d>, block:<%d, %d>\n", grid.x, grid.y, block.x, block.y);addMatrix<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);  // 调用内核函数ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__); for (int i = 0; i < 10; i++){printf("id=%d, matrix_A=%d, matrix_B=%d, result=%d\n", i + 1,ipHost_A[i], ipHost_B[i], ipHost_C[i]);}free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__); ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__); return 0;
}

8.3、二维网格二维线程块计算二维矩阵加法

#include <stdio.h>
#include "../tools/common.cuh"__global__ void addMatrix(int *A, int *B, int *C, const int nx, const int ny)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;;unsigned int idx = iy * nx + ix;if (ix < nx && iy < ny){C[idx] = A[idx] + B[idx];}
}int main(void)
{// 1、设置GPU设备setGPU();// 2、分配主机内存和设备内存，并初始化int nx = 16;int ny = 8;int nxy = nx * ny;size_t stBytesCount = nxy * sizeof(int);// （1）分配主机内存，并初始化int *ipHost_A, *ipHost_B, *ipHost_C;ipHost_A = (int *)malloc(stBytesCount);ipHost_B = (int *)malloc(stBytesCount);ipHost_C = (int *)malloc(stBytesCount);if (ipHost_A != NULL && ipHost_B != NULL && ipHost_C != NULL){for (int i = 0; i < nxy; i++){ipHost_A[i] = i;ipHost_B[i] = i + 1;}memset(ipHost_C, 0, stBytesCount); }else{printf("Fail to allocate host memory!\n");exit(-1);}// （2）分配设备内存，并初始化int *ipDevice_A, *ipDevice_B, *ipDevice_C;ErrorCheck(cudaMalloc((int**)&ipDevice_A, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_B, stBytesCount), __FILE__, __LINE__); ErrorCheck(cudaMalloc((int**)&ipDevice_C, stBytesCount), __FILE__, __LINE__); if (ipDevice_A != NULL && ipDevice_B != NULL && ipDevice_C != NULL){ErrorCheck(cudaMemcpy(ipDevice_A, ipHost_A, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_B, ipHost_B, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); ErrorCheck(cudaMemcpy(ipDevice_C, ipHost_C, stBytesCount, cudaMemcpyHostToDevice), __FILE__, __LINE__); }   else{printf("Fail to allocate memory\n");free(ipHost_A);free(ipHost_B);free(ipHost_C);exit(1);}// calculate on GPUdim3 block(4, 4);dim3 grid((nx + block.x -1) / block.x, (ny + block.y - 1) / block.y);printf("Thread config:grid:<%d, %d>, block:<%d, %d>\n", grid.x, grid.y, block.x, block.y);addMatrix<<<grid, block>>>(ipDevice_A, ipDevice_B, ipDevice_C, nx, ny);  // 调用内核函数ErrorCheck(cudaMemcpy(ipHost_C, ipDevice_C, stBytesCount, cudaMemcpyDeviceToHost), __FILE__, __LINE__); for (int i = 0; i < 10; i++){printf("id=%d, matrix_A=%d, matrix_B=%d, result=%d\n", i + 1,ipHost_A[i], ipHost_B[i], ipHost_C[i]);}free(ipHost_A);free(ipHost_B);free(ipHost_C);ErrorCheck(cudaFree(ipDevice_A), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_B), __FILE__, __LINE__); ErrorCheck(cudaFree(ipDevice_C), __FILE__, __LINE__); ErrorCheck(cudaDeviceReset(), __FILE__, __LINE__); return 0;
}

9、内存结构

在这里插入图片描述
在这里插入图片描述
在这里插入图片描述

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/pingmian/41487.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

Google Earth Engine（GEE）——控制台ui.Textbox复制你想要的textbox

Google Earth Engine（GEE）——控制台ui.Textbox复制你想要的textbox

结果 函数： ui.Textbox(placeholder, value, onChange, disabled, style) A textbox that enables the user to input text information. Arguments: placeholder (String, optional): The placeholder text to display when the textbox is empty. Defaults to none. …
阅读更多...
VBA语言専攻T3学员领取资料通知0706

VBA语言専攻T3学员领取资料通知0706

T3学员领取资料通知0706 各位学员∶本周MF系列VBA技术资料增加641-645讲&#xff0c;T3学员看到通知后请免费领取,领取时间7月5日晚上19:00-7月6日晚上19:00。本次增加内容&#xff1a; MF641:前个窗体组合框选项联动下个组合框 MF642:工作表中数据选项联动下个数据验证 MF…
阅读更多...
技术驱动旅游创新！深度解析景区导览小程序的地图渲染与AR导航技术

技术驱动旅游创新！深度解析景区导览小程序的地图渲染与AR导航技术

随着现代生活节奏的加快&#xff0c;人们在外出旅游时更倾向于轻便出行&#xff0c;携带导览地图已成为过去。然而&#xff0c;面对景区广阔的面积和众多景点&#xff0c;游客常常感到迷茫&#xff0c;难以快速定位到自己所需的地点。景区导览小程序让游客只需搜索景区名称&…
阅读更多...
人工智能系列-NumPy（二）

人工智能系列-NumPy（二）

&#x1f308;个人主页&#xff1a;羽晨同学 &#x1f4ab;个人格言:“成为自己未来的主人~” 链接数组 anp.array([[1,2],[3,4]]) print(第一个数组&#xff1a;) print(a) print(\n) bnp.array([[5,6],[7,8]]) print(第二个数组&#xff1a;) print(b) print(\n) print…
阅读更多...
zigbee笔记：六、看门狗定时器（Watch Dog）

zigbee笔记：六、看门狗定时器（Watch Dog）

一、看门狗基础 1、看门狗功能&#xff1a; 由于单片机的工作常常会受到来自外界电磁场的干扰&#xff0c;造成各种寄存器和内存的数据混乱&#xff0c;会导致程序指针错误等&#xff0c;程序运行可能会陷入死循环。程序的正常运行被打断&#xff0c;由单片机控制的系统无法继…
阅读更多...
Motion Guidance: 扩散模型实现图像精确编辑的创新方法

Motion Guidance: 扩散模型实现图像精确编辑的创新方法

在深度学习领域&#xff0c;扩散模型&#xff08;diffusion models&#xff09;因其能够根据文本描述生成高质量图像而备受关注。然而&#xff0c;这些模型在精确编辑图像中对象的布局、位置、姿态和形状方面仍存在挑战。本文提出了一种名为“运动引导”&#xff08;motion gui…
阅读更多...
【数据库了解与学习】

【数据库了解与学习】

1.下载所需版本安装包 1.1将所需文件压缩包以及安装包放在你选择的任意一盘&#xff0c;新建一个没有文字和空格的文件夹 1.2双击打开安装包&#xff0c;选择Custom自定义模式然后点击右下方的Next 1.4三连点击1&#xff0c;再点击箭头出现3&#xff0c;选中3出现4&#xff0c;…
阅读更多...
Infinitar链游新发展新机遇

Infinitar链游新发展新机遇

区块链游戏市场在近年来经历了显著增长&#xff0c;吸引了大量的投资和关注。随着加密货币和NFT&#xff08;非同质化代币&#xff09;概念的普及&#xff0c;越来越多的投资者、游戏开发者和看到了区块链技术在游戏领域的应用潜力&#xff0c;纷纷涌入市场。区块链游戏的用户量…
阅读更多...
Aigtek功率放大器的参数及应用是什么

Aigtek功率放大器的参数及应用是什么

功率放大器是电子电路中的重要组成部分&#xff0c;用于将输入信号的功率增加到更高的水平。它们在各种电子设备和应用中发挥着关键作用。下面Aigtek安泰电子将介绍功率放大器的主要参数以及它们在不同领域的应用。 1.功率放大器的基本参数 增益 功率放大器的增益是指输出信号的…
阅读更多...
android应用的持续构建CI（五）-- 应用管理

android应用的持续构建CI（五）-- 应用管理

一、ER设计 渠道&#xff1a;应用发布的渠道&#xff0c;To B的话&#xff0c;可能是学校、租户等&#xff1b;To C的话&#xff0c;渠道则是不同厂商的应用市场。审核&#xff1a;jenkins构建应用完成后&#xff0c;管理后台生成应用的审核记录。 下面是每个实体详细的字段&a…
阅读更多...
昇思25天学习打卡营第17天（+1）|Diffusion扩散模型

昇思25天学习打卡营第17天（+1）|Diffusion扩散模型

1. 学习内容复盘 本文基于Hugging Face&#xff1a;The Annotated Diffusion Model一文翻译迁移而来&#xff0c;同时参考了由浅入深了解Diffusion Model一文。 本教程在Jupyter Notebook上成功运行。如您下载本文档为Python文件&#xff0c;执行Python文件时&#xff0c;请确…
阅读更多...
Git仓库介绍

Git仓库介绍

1. Github GitHub 本身是一个基于云端的代码托管平台&#xff0c;它提供的是远程服务&#xff0c;而不是一个可以安装在本地局域网的应用程序。因此&#xff0c;GitHub 不可以直接在本地局域网进行安装。 简介&#xff1a;GitHub是最流行的代码托管平台&#xff0c;提供了大量…
阅读更多...
昇思MindSpore学习笔记4-02生成式--DCGAN生成漫画头像

昇思MindSpore学习笔记4-02生成式--DCGAN生成漫画头像

摘要&#xff1a; 记录了昇思MindSpore AI框架使用70171张动漫头像图片训练一个DCGAN神经网络生成式对抗网络&#xff0c;并用来生成漫画头像的过程、步骤。包括环境准备、下载数据集、加载数据和预处理、构造网络、模型训练等。 一、概念 深度卷积对抗生成网络DCGAN Deep C…
阅读更多...
STM32+ESP8266(ESP32)+MQTT+阿里云物联网平台

STM32+ESP8266(ESP32)+MQTT+阿里云物联网平台

1、阿里云物联网平台 - 阿里云物联网平台配置 产品-设备-编辑物模型-设备端开发-查看上报数据 在产品上添加物模型&#xff0c;然后设备是继承自产品的&#xff0c;因此也具有物模型 添加产品、添加设备、产品上添加物模型 - 使用IOT Studio 绘制界面显示温度、湿度、灯开…
阅读更多...
Tomcat（+Servlet）笔记+代码

Tomcat（+Servlet）笔记+代码

Tomcat安装和配置 安装在不含中文的路径&#xff0c;路径不能太长 Apache 官网&#x1f447; Apache Tomcat - Welcome! 配置部分 点击下图红框处&#xff0c;找到Tomcat安装位置 添加项目的文件 配好的话&#xff0c;红框这里有个猫 代码部分 新建jsp文件&#xff0c;里…
阅读更多...
线程（基础概念）

线程（基础概念）

文章目录 一、线程和进程&#xff1f;二、线程初识2.1 线程属性2.2 线程的调度策略2.3 线程的优先级2.3 线程实验 一、线程和进程&#xff1f; 我们经常描述进程&#xff08;process&#xff09;和线程&#xff08;thread&#xff09;&#xff1a; 进程是资源管理的最小单位&a…
阅读更多...
昇思25天学习打卡营第07天 | 函数式自动微分

昇思25天学习打卡营第07天 | 函数式自动微分

昇思25天学习打卡营第07天 | 函数式自动微分 文章目录 昇思25天学习打卡营第07天 | 函数式自动微分函数与计算图微分函数与梯度Stop GradientAuxiliary data 神经网络梯度计算总结打卡 神经网络的训练主要使用反向传播算法&#xff0c;首先计算模型预测值&#xff08;logits&am…
阅读更多...
科普文：微服务之服务网格Service Mesh

科普文：微服务之服务网格Service Mesh

一、ServiceMesh概念 背景 随着业务的发展&#xff0c;传统单体应用的问题越来越严重&#xff1a; 单体应用代码库庞大&#xff0c;不易于理解和修改持续部署困难&#xff0c;由于单体应用各组件间依赖性强&#xff0c;只要其中任何一个组件发生更改&#xff0c;将重新部署整…
阅读更多...
MUNIK解读ISO26262--什么是DFA

MUNIK解读ISO26262--什么是DFA

我们在学习功能安全过程中&#xff0c;经常会听到很多安全分析方法&#xff0c;有我们熟知的FMEA(Failure Modes Effects Analysis)和FTA(Fault Tree Analysis)还有功能安全产品设计中几乎绕不开的FMEDA(Failure Modes Effects and Diagnostic Analysis)&#xff0c;相比于它们…
阅读更多...
【OceanBase】OBProxy 无状态的理解

【OceanBase】OBProxy 无状态的理解

SueWakeup 个人主页&#xff1a;SueWakeup 系列专栏&#xff1a;为祖国的科技进步添砖Java 个性签名&#xff1a;保留赤子之心也许是种幸运吧 本文封面由 凯楠&#x1f4f8;友情提供 目录 前言 OBProxy 无状态的概述 OBProxy 无状态特性带来的优点 1. 高可用 2. 负载均衡…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023