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);
    }

    // 设置执行GPU
    int 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);  // 主机内存初始化为0
        memset(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__); // 设备内存初始化为0
        ErrorCheck(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);  // 主机内存初始化为0
        memset(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__); // 设备内存初始化为0
        ErrorCheck(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: // Fermi
      if (devProp.minor == 1) cores = mp * 48;
      else cores = mp * 32;
      break;
     case 3: // Kepler
      cores = mp * 192;
      break;
     case 5: // Maxwell
      cores = mp * 128;
      break;
     case 6: // Pascal
      if ((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 Turing
      if ((devProp.minor == 0) || (devProp.minor == 5)) cores = mp * 64;
      else printf("Unknown device type\n");
      break;
     case 8: // Ampere
      if (devProp.minor == 0) cores = mp * 64;
      else if (devProp.minor == 6) cores = mp * 128;
      else if (devProp.minor == 9) cores = mp * 128; // ada lovelace
      else printf("Unknown device type\n");
      break;
     case 9: // Hopper
      if (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 GPU
    dim3 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 GPU
    dim3 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 GPU
    dim3 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.mfbz.cn/a/774875.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

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

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

基于Python爬虫的城市二手房数据分析可视化 一、前言二、数据采集（爬虫，附完整代码）三、数据可视化（附完整代码）3.1 房源面积-总价散点图3.2 各行政区均价3.3 均价最高的10个小区3.4 均价最高的10个地段3.5 户型分布3.6 词云图四、如何更换城市一、前言 二手房具有价格普…
阅读更多...
博途通讯笔记1：1200与1200之间S7通讯

博途通讯笔记1：1200与1200之间S7通讯

目录 一、添加子网连接二、创建PUT GET三、各个参数的意义 一、添加子网连接 二、创建PUT GET 三、各个参数的意义
阅读更多...
换根dp，CF 633F - The Chocolate Spree

换根dp，CF 633F - The Chocolate Spree

一、题目 1、题目描述 2、输入输出 2.1输入 2.2输出 3、原题链接 633F - The Chocolate Spree 二、解题报告 1、思路分析 2600的题&#xff0c;但是不算很困难。 先考虑暴力做法&#xff0c;如何得到两条不相交的路径&#xff1f; 枚举删除的边&#xff0c;得到两棵子树…
阅读更多...
idea集成tomcat报错kotlin/Result

idea集成tomcat报错kotlin/Result

记一次报错 解决办法
阅读更多...
鼠标自动点击器怎么用？鼠标连点器入门教程！

鼠标自动点击器怎么用？鼠标连点器入门教程！

鼠标自动点击器是适用于Windows电脑的自动执行鼠标点击操作的工具&#xff0c;主要用于模拟鼠标点击操作&#xff0c;实现鼠标高速点击的操作。通过模拟鼠标点击&#xff0c;可以在用户设定的位置、频率和次数下自动执行点击动作。 鼠标自动点击器主要的应用场景&#xff1a; …
阅读更多...
数据操作10-15题（30 天 Pandas 挑战）

数据操作10-15题（30 天 Pandas 挑战）

数据操作 1. 相关知识点1.12 分组与连表1.13 排名 2. 题目2.10 第N高的薪水2.11 第二高的薪水2.12 部门工资最高的员工2.13 分数排名2.14 删除重复的电子邮箱2.15 每个产品在不同商店的价格 1. 相关知识点 1.12 分组与连表 分组max_salaryemployee.groupby(departmentId)[sal…
阅读更多...
超简易SpringBoot工程构建与部署 ( 图解 - 零基础专属 )

超简易SpringBoot工程构建与部署 ( 图解 - 零基础专属 )

目录 简单了解MVC架构 模型&#xff08;Model&#xff09; 视图&#xff08;View&#xff09; 控制器&#xff08;Controller&#xff09; 基本环境准备 MYSQL建库建表 创库创表 智能生成数据 创建SpringBoot项目 配置pox.xml 代码提供 补充(IDEA的Maven要配置正确…
阅读更多...
用kimi和claude自动生成时间轴图表

用kimi和claude自动生成时间轴图表

做时间轴图表并不难&#xff0c;但是很麻烦&#xff0c;先要大量收集相关事件&#xff0c;然后在一些图表软件中反复调整操作。现在借助AI工具&#xff0c;可以自动生成了。 首先&#xff0c;在kimi中输入提示词来获取某个企业的大事记&#xff1a; 联网检索&#xff0c;元语…
阅读更多...
前后端数据交互流程

前后端数据交互流程

一、前言 用户在浏览器访问一个网站时&#xff0c;会有前后端数据交互的过程&#xff0c;前后端数据交互也有几种的情况&#xff0c;一下就简单的来说明一下 二、原理 介绍前后端交互前先来了解一下浏览器的功能&#xff0c;浏览器通过渲染引擎和 JavaScript 引擎协同工作&am…
阅读更多...
uboot ethernet初始化

uboot ethernet初始化

在Board_r.c 使用initr_net,先初始化phy,然后初始化gmac,driver/net/gmacv300/gmac.c实现管脚复用和gmac设备注册 Board_r.c #ifdef CONFIG_CMD_NETstatic int initr_net(void){puts("Net: ");eth_initialize();#if defined(CONFIG_RESET_PHY_R)debug("Reset…
阅读更多...
计算机视觉 图像融合技术概览

计算机视觉 图像融合技术概览

在许多计算机视觉应用中（例如机器人运动和医学成像），需要将来自多幅图像的相关信息集成到一幅图像中。这种图像融合将提供更高的可靠性、准确性和数据质量。 多视图融合可以提高图像的分辨率，同时恢复场景的 3D 表示。多模态融合结合了来自不同传感器的图像，称为多传感器融…
阅读更多...
短视频文案提取神器怎么提取抖音视频文案！

短视频文案提取神器怎么提取抖音视频文案！

很多编导以及视频内容创作者为了提高自己的工作效率还会使用视频转文字提取神器&#xff0c;我们都清楚短视频领域每个平台人群熟悉都有所不同&#xff0c;在分发内容的时候也会调整内容已符合平台属性。 短视频文案提取神器怎么提取抖音视频文案 短视频常见的平台有抖音、西瓜…
阅读更多...
SpringBoot实战：轻松实现XSS攻击防御（注解和过滤器）

SpringBoot实战：轻松实现XSS攻击防御（注解和过滤器）

文章目录 引言一、XSS攻击概述1.1 XSS攻击的定义1.2 XSS攻击的类型1.3 XSS攻击的攻击原理及示例 二、Spring Boot中的XSS防御手段2.1 使用注解进行XSS防御2.1.1 引入相关依赖2.1.2 使用XSS注解进行参数校验2.1.3 实现自定义注解处理器2.1.4 使用注解 2.2 使用过滤器进行XSS防御…
阅读更多...
现在这个行情怎么理解股票期权？一个守住底线的工具！

现在这个行情怎么理解股票期权？一个守住底线的工具！

今天带你了解现在这个行情怎么理解股票期权&#xff1f;一个守住底线的工具&#xff01;股票期权是一种金融衍生品&#xff0c;给予持有者在未来特定时间以特定价格购买或出售股票的权利。 行情看down可以买入看跌期权&#xff0c;看跌期权的买方是因为预计行情的价格会在近期…
阅读更多...
imx6ull/linux应用编程学习(11）CAN应用编程基础

imx6ull/linux应用编程学习(11）CAN应用编程基础

关于裸机的can通信&#xff0c;会在其他文章发&#xff0c;这里主要讲讲linux上的can通信。 与I2C,SPI等同步通讯方式不同&#xff0c;CAN通讯是异步通讯&#xff0c;也就是没有时钟信号线来保持信号接收同步&#xff0c;也就是所说的半双工&#xff0c;无法同时发送与接收&…
阅读更多...
【python】PyQt5控件尺寸大小位置，内容边距等API调用方法实战解析

【python】PyQt5控件尺寸大小位置，内容边距等API调用方法实战解析

✨✨ 欢迎大家来到景天科技苑✨✨ &#x1f388;&#x1f388; 养成好习惯&#xff0c;先赞后看哦~&#x1f388;&#x1f388; &#x1f3c6; 作者简介&#xff1a;景天科技苑 &#x1f3c6;《头衔》&#xff1a;大厂架构师&#xff0c;华为云开发者社区专家博主&#xff0c;…
阅读更多...
小（微）间距P1.538COB渠道现货销售将加速全面升级替换SMD产品。

小（微）间距P1.538COB渠道现货销售将加速全面升级替换SMD产品。

COB&#xff08;Chip on Board&#xff09;技术&#xff0c;如一颗璀璨的星辰&#xff0c;在上世纪60年代的科技夜空中悄然升起。它巧妙地将LED芯片镶嵌在PCB电路板的怀抱中&#xff0c;再用特种树脂为其披上一层坚韧的外衣&#xff0c;宛如一位精心雕琢的艺术家在创作一幅完美…
阅读更多...
【Python机器学习】处理文本数据——用tf-idf缩放数据

【Python机器学习】处理文本数据——用tf-idf缩放数据

为了按照我们预计的特征信息量大小来缩放特征&#xff0c;而不是舍弃那些认为不重要的特征&#xff0c;最常见的一种做法就是使用词频-逆向文档频率&#xff08;tf-idf&#xff09;。这一方法对某个特定文档中经常出现的术语给与很高的权重&#xff0c;但是堆在语料库的许多文档…
阅读更多...
【UE5.1】Chaos物理系统基础——04 事件驱动粒子效果

【UE5.1】Chaos物理系统基础——04 事件驱动粒子效果

目录 效果 步骤 一、炸开时的烟雾效果 二、炸开时的碎片效果 效果 步骤 一、炸开时的烟雾效果 1. 选中场景中的几何体集&#xff0c;勾选“通知中断”&#xff0c;这样当几何体集解体时就能产生一个事件&#xff1b;勾选“通知碰撞”&#xff0c;当几何体集检测到碰撞也能…
阅读更多...
Java+前后端分离架构+ MySQL8.0.36产科信息管理系统 产科电子病历系统源码

Java+前后端分离架构+ MySQL8.0.36产科信息管理系统 产科电子病历系统源码

Java前后端分离架构 MySQL8.0.36产科信息管理系统 产科电子病历系统源码 产科信息管理系统—住院管理 数字化产科住院管理是现代医院管理中的重要组成部分&#xff0c;它利用数字化技术优化住院流程&#xff0c;提升医疗服务质量和效率。以下是对数字化产科住院管理的详细阐述…
阅读更多...
最新文章

Copyright @ 2022~2023

友情链接： 我的编程经验分享 一条长河网 尧图网站开发 尧图建网站