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.xdnf.cn/news/1472899.html

如若内容造成侵权/违法违规/事实不符，请联系一条长河网进行投诉反馈，一经查实，立即删除！

相关文章

知名品牌因商标痛失市场：114家直营店山寨店7000多家！

知名品牌因商标痛失市场：114家直营店山寨店7000多家！

奶茶知名品牌“鹿角巷”当年红遍大江南北&#xff0c;是最早的新茶饮品牌&#xff0c;但是当年商标注册存在问题&#xff0c;被同行奶茶品牌抢占了先机&#xff0c;发声明“对大陆商标注册细则不详&#xff0c;在商标注册过程中让假店钻了法律空档”&#xff0c;最夸张的时候全…
阅读更多...
python实践-实现TTS项目的应用

python实践-实现TTS项目的应用

#coqui-ai TTS 项目地址# 版本&#xff1a; V0.22.0 博主使用的python版本&#xff1a;3.10.6 一、项目下载安装 下载&#xff1a; git或者直接下载都可 git clone https://github.com/coqui-ai/TTS.git 建议选择压缩包下载&#xff0c;选好对应版本。方便github会卡的小…
阅读更多...
51单片机STC89C52RC——14.1 直流电机调速

51单片机STC89C52RC——14.1 直流电机调速

目录 目的/效果 1&#xff1a;电机转速同步LED呼吸灯 2 通过独立按键 控制直流电机转速。 一&#xff0c;STC单片机模块 二&#xff0c;直流电机 2.1 简介 2.2 驱动电路 2.2.1 大功率器件直接驱动 2.2.2 H桥驱动 正转 反转 2.2.3 ULN2003D 引脚、电路 2.3 PWM&…
阅读更多...
【刷题汇总--游游的you、腐烂的苹果、孩子们的游戏(圆圈中最后剩下的数)】

【刷题汇总--游游的you、腐烂的苹果、孩子们的游戏(圆圈中最后剩下的数)】

C日常刷题积累 今日刷题汇总 - day0051、游游的you1.1、题目1.2、思路1.3、程序实现 - 蛮力法1.4、程序实现 - 贪心(优化) 2、腐烂的苹果2.1、题目2.2、思路2.3、程序实现 - bfs 3、孩子们的游戏(圆圈中最后剩下的数)3.1、题目3.2、思路3.3、程序实现 -- 环形链表3.4、程序实现…
阅读更多...
鸿蒙开发：Universal Keystore Kit（密钥管理服务）【生成密钥(C/C++)】

鸿蒙开发：Universal Keystore Kit（密钥管理服务）【生成密钥(C/C++)】

生成密钥(C/C) 以生成ECC密钥为例&#xff0c;生成随机密钥。具体的场景介绍及支持的算法规格。 注意&#xff1a; 密钥别名中禁止包含个人数据等敏感信息。 开发前请熟悉鸿蒙开发指导文档&#xff1a;gitee.com/li-shizhen-skin/harmony-os/blob/master/README.md点击或者复…
阅读更多...
（2024）KAN: Kolmogorov–Arnold Networks：评论

（2024）KAN: Kolmogorov–Arnold Networks：评论

KAN: Kolmogorov–Arnold Networks: A review 公和众与号&#xff1a;EDPJ&#xff08;进 Q 交流群&#xff1a;922230617 或加 VX&#xff1a;CV_EDPJ 进 V 交流群&#xff09; 目录 0. 摘要 1. MLP 也有可学习的激活函数 2. 标题的意义 3. KAN 是具有样条基激活函数的 M…
阅读更多...
基于Python爬虫的城市二手房数据分析可视化

基于Python爬虫的城市二手房数据分析可视化

基于Python爬虫的城市二手房数据分析可视化 一、前言二、数据采集（爬虫，附完整代码）三、数据可视化（附完整代码）3.1 房源面积-总价散点图3.2 各行政区均价3.3 均价最高的10个小区3.4 均价最高的10个地段3.5 户型分布3.6 词云图四、如何更换城市一、前言 二手房具有价格普…
阅读更多...
《安全行业大模型技术应用态势发展报告（2024）》

《安全行业大模型技术应用态势发展报告（2024）》

人工智能技术快速迭代发展&#xff0c;大模型应用场景不断拓展&#xff0c;随着安全行业对人工智能技术的应用程度日益加深&#xff0c;大模型在网络安全领域的应用潜力和挑战逐渐显现。安全行业大模型技术的应用实践不断涌现&#xff0c;其在威胁检测、风险评估和安全运营等方…
阅读更多...
idm下载慢怎么回事 idm批量导入下载使用方法

idm下载慢怎么回事 idm批量导入下载使用方法

IDM (Internet Download Manager)是一款兼容性大&#xff0c;支持多种语言的下载管理软件&#xff0c;它可以自动检测并下载网页上的内容&#xff0c;这正是这一优点&#xff0c;使得它受到了广大用户的喜爱。在日常使用互联网的过程中&#xff0c;快速下载文件对我们来说非常重…
阅读更多...
Patch SCN使用说明---惜分飞

Patch SCN使用说明---惜分飞

软件说明 该软件是惜分飞&#xff08;https://www.xifenfei.com&#xff09;开发&#xff0c;仅用来查看和修改Oracle数据库SCN(System Change Number),主要使用在数据库因为某种原因导致无法正常启动的情况下使用该工具进行解决.特别是Oracle新版本中使用隐含参数,event,orad…
阅读更多...
RH850系列芯片深度剖析 1.8-内存管理之MPU

RH850系列芯片深度剖析 1.8-内存管理之MPU

RH850系列芯片深度剖析 1.8-内存管理之MPU 文章目录 RH850系列芯片深度剖析 1.8-内存管理之MPU一、MPU简介1.1 功能特性1.2 系统保护标识符（SPID）二、保护区域设置2.1 保护区域属性设置2.2 保护区域设置注意事项2.2.1 跨越保护区域边界2.2.2 无效的保护区域设置2.2.3 保护违规…
阅读更多...
2-3 图像分类数据集

2-3 图像分类数据集

MNIST数据集是图像分类任务中广泛使用的数据集之一&#xff0c;但作为基准数据集过于简单&#xff0c;我们将使用类似但更复杂的Fashion-MNIST数据集。 %matplotlib inline import torch import torchvision # pytorch模型关于计算机视觉模型实现的一个库 from torch.utils i…
阅读更多...
【Unity 人性动画的复用性】

【Unity 人性动画的复用性】

Unity的动画系统&#xff0c;通常称为Mecanim&#xff0c;提供了强大的动画复用功能&#xff0c;特别是针对人型动画的重定向技术。这种技术允许开发者将一组动画应用到不同的角色模型上&#xff0c;而不需要为每个模型单独制作动画。这通过在模型的骨骼结构之间建立对应关系来…
阅读更多...
【Odoo开源ERP】别把ERP与进销存软件混为一谈

【Odoo开源ERP】别把ERP与进销存软件混为一谈

导读&#xff1a;企业使用ERP软件能够实现管理升级&#xff0c;多方信息集成&#xff0c;按照既定策略逻辑运算&#xff0c;生成计划建议&#xff0c;减少人力成本&#xff0c;提高准确率的同时提高经营能力。 ERP&#xff0c;是MRP II的下一代软件&#xff0c;除了MRP II已有的…
阅读更多...
JAVA 获取客户端信息工具类

JAVA 获取客户端信息工具类

获取客户端信息工具类 import com.wangyao.common.constant.ConstantNumeral; import jakarta.annotation.Nullable; import jakarta.servlet.http.HttpServletRequest; import jakarta.validation.constraints.NotNull; import lombok.extern.slf4j.Slf4j;import java.io.IOEx…
阅读更多...
[单master节点k8s部署]18.监控系统构建（三）Grafana安装

[单master节点k8s部署]18.监控系统构建（三）Grafana安装

Grafana是一个跨平台的开源的度量分析和可视化工具。支持多种数据源&#xff0c;比如OpenTSDB&#xff0c;Prometheus&#xff0c;ElasticResearch&#xff0c;Cloudwatch等。 Grafana安装 通过yaml配置grafana的pod和service&#xff0c;grafana工作在kube-system的命名空间…
阅读更多...
【C++】 解决 C++ 语言报错：Memory Leak

【C++】 解决 C++ 语言报错：Memory Leak

文章目录 引言 内存泄漏&#xff08;Memory Leak&#xff09;是 C 编程中常见且严重的内存管理问题之一。当程序分配了内存而没有正确释放&#xff0c;导致内存无法被重新利用时&#xff0c;就会发生内存泄漏。这种错误会导致程序占用越来越多的内存&#xff0c;最终可能导致系…
阅读更多...
NFT音乐版权系统的主要功能

NFT音乐版权系统的主要功能

NFT音乐版权系统是指利用区块链技术和NFT技术来管理和交易音乐版权的系统。该系统的主要功能包括以下几个方面。北京木奇移动技术有限公司&#xff0c;专业的软件外包开发公司&#xff0c;欢迎交流合作。 1. 音乐版权确权 NFT音乐版权系统可以为音乐作品的版权提供独特的标识和…
阅读更多...
FineBI在线学习资源-数据处理

FineBI在线学习资源-数据处理

FineBI在线学习资源汇总&#xff1a; 学习资源 视频课程 帮助文档 问答 数据处理学习文档&#xff1a; 相关资料&#xff1a; 故事背景概述-https://help.fanruan.com/finebi6.0/doc-view-1789.html 基础表处理-https://help.fanruan.com/finebi6.0/doc-view-1791.html …
阅读更多...
什么是数字化产科管理平台？

什么是数字化产科管理平台？

什么是数字化产科管理平台&#xff1f; 数字化产科管理平台是为医院产科量身定制的信息管理系统&#xff0c;旨在提高医院产科的服务质量和管理效率。该平台全面覆盖了孕妇从怀孕开始到生产结束以及产后42天以内的一系列医院保健服务信息。以下是数字产科管理平台的详细介绍&a…
阅读更多...
最新文章

Copyright @ 2022~2023