CUDA编程基础

news2024/11/26 4:25:47

文章目录

  • 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.coloradmin.cn/o/1897865.html

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

相关文章

ret2syscall简单总结

ret2syscall简单总结

主要是自己的简单的学习总结。 知识点 关于系统调用如何传递参数问题&#xff0c;即系统调用约定&#xff08;syscall&#xff0c;int 80h&#xff0c;svc&#xff09;_int 80h intel汇编用法-CSDN博客 ret2syscall的做题思路&#xff08;以32位程序为例&#xff09; - ZikH…
阅读更多...
如何从腾讯云迁移到AWS

如何从腾讯云迁移到AWS

随着跨境出海潮不断扩大&#xff0c;企业越来越意识到将工作负载迁移到海外节点的必要性&#xff0c;以获取更多功能、灵活性和性能。然而&#xff0c;顺利迁移业务主机并确保业务稳定访问是一项具有挑战性的任务。在此挑战中&#xff0c;借助AWS迁移工具和迁移流程的强大支持&…
阅读更多...
一文了解常见DNS结构

一文了解常见DNS结构

很多企业忽略DNS这个关键的组件&#xff0c;而当DNS出现问题是&#xff0c;就会导致网站无法访问、电子邮件无法发送和接收&#xff0c;从而影响到企业的正常运行。而网络团队成员如果想要处理DNS问题就必须对所在网络的DNS架构有一定的理解。 主DNS服务器&#xff1a; 主DNS服…
阅读更多...
微信开放能力 微信小程序获取微信头像、昵称、转发功能、分享到朋友圈、手机号验证组件、客服功能等

微信开放能力 微信小程序获取微信头像、昵称、转发功能、分享到朋友圈、手机号验证组件、客服功能等

一、获取微信头像 当小程序需要让用户完善个人资料时&#xff0c;我们可以通过提供的头像、昵称填写能力快速完善示例 在 wxml 文件中 <view><button class"btn" open-type"chooseAvatar" bindchooseavatar"chooseAvatar"><im…
阅读更多...
深度报告 | 百度安全携手极越安全发布《整车安全渗透测试白皮书》

深度报告 | 百度安全携手极越安全发布《整车安全渗透测试白皮书》

注重点&#xff0c;如何确保车辆全生命周期的安全已成为整个行业亟待解决的问题。对于车企而言&#xff0c;通过渗透测试尽量多地发现安全威胁&#xff0c;是确保车辆信息系统的稳定运行、保障用户安全驾驶至关重要的措施。然而&#xff0c;传统的渗透测试方法已无法满足智能网…
阅读更多...
谈大语言模型动态思维流程编排

谈大语言模型动态思维流程编排

尽管大语言模型已经呈现出了强大的威力&#xff0c;但是如何让它完美地完成一个大的问题&#xff0c;仍然是一个巨大的挑战。 需要精心地给予大模型许多的提示&#xff08;Prompt&#xff09;。对于一个复杂的应用场景&#xff0c;编写一套完整的&#xff0c;准确无误的提示&am…
阅读更多...
Ubuntu 24.04 上安装 Kubernetes，超级详细的教程！

Ubuntu 24.04 上安装 Kubernetes，超级详细的教程！

Kubernetes 是一个免费的开源容器编排工具&#xff0c;它允许基于容器的应用程序的自动化部署、扩展和管理。 我们将介绍如何使用 Kubeadm 逐步在 Ubuntu 24.04 上安装 Kubernetes 此次演示中&#xff0c;我们将使用以下三个 Ubuntu 24.04 实例 Instance 1 : Master Node (k…
阅读更多...
【Mathematica14.0】快速从下载安装到使用

【Mathematica14.0】快速从下载安装到使用

目录 1.简介 2.下载安装 下载 安装 3.一小时掌握mathematica使用 单元模式 内置函数 符号表达式 迭代器 赋值 通配符及查找替换 函数定义 匿名函数&#xff08;拉姆达表达式&#xff09; 函数映射 函数式与运算符 函数自定义选项 图形可视化 交互式界面 数值…
阅读更多...
window系统openssl开发环境搭建(VS2017)

window系统openssl开发环境搭建(VS2017)

window系统openssl开发环境搭建 VS2017 一、下载openssl二、安装openssl三、openssl项目配置3.1 配置include文件3.2 配置openssl动态库四、编写openssl测试代码五、问题总结5.1 问题 一5.2 问题二一、下载openssl https://slproweb.com/products/Win32OpenSSL.html 根据自己…
阅读更多...
本迪戈和阿德莱德银行与MongoDB合作， 利用生成式AI对银行核心技术进行现代化改造

本迪戈和阿德莱德银行与MongoDB合作， 利用生成式AI对银行核心技术进行现代化改造

MongoDB公司&#xff08;纳斯达克股票代码&#xff1a;MDB&#xff09;近日宣布与本迪戈和阿德莱德银行 (Bendigo and Adelaide Bank&#xff0c;澳大利亚证券交易所股票代码&#xff1a;BEN&#xff09;建立合作伙伴关系。 本迪戈和阿德莱德银行将使用MongoDB Atlas对其银行核…
阅读更多...
自闭症孩子的语言之旅：最晚几岁会说话的探索与思考

自闭症孩子的语言之旅：最晚几岁会说话的探索与思考

作为在自闭症学校工作的教育者&#xff0c;我深知自闭症这一神经发展性障碍给孩子们带来的挑战&#xff0c;尤其是他们在语言发展方面的困难。自闭症孩子的语言发展轨迹各不相同&#xff0c;有的孩子可能早早地展现出语言天赋&#xff0c;而有的孩子则可能迟迟不开口。那么&…
阅读更多...
Windows搭建本地对象存储服务MinIO并且使用springboot整合

Windows搭建本地对象存储服务MinIO并且使用springboot整合

开发文档&#xff1a; MinIO Windows中文文档 MinIO Object Storage for Windows &#xff08;英文文档&#xff09; 1、准备工作 准备一个空的文件夹&#xff0c;用来存放minio相关的内容&#xff1b; 这里是在D盘创建一个minio的文件夹&#xff1b; 后续所有跟MinIO相关…
阅读更多...
【LLM】一、利用ollama本地部署大模型

【LLM】一、利用ollama本地部署大模型

目录 前言 一、Ollama 简介 1、什么是Ollama 2、特点&#xff1a; 二、Windows部署 1.下载 2.安装 3.测试安装 4.模型部署&#xff1a; 5.注意 三、 Docker部署 1.docker安装 2.ollama镜像拉取 3.ollama运行容器 4.模型部署&#xff1a; 5.注意&#xff1a; 总结 前言…
阅读更多...
Ubuntu系统复制文件到共享文件夹出错

Ubuntu系统复制文件到共享文件夹出错

1、问题描述 Ubuntu系统复制文件到共享文件夹时&#xff0c;出现拼接文件时出错&#xff1a;输入/输出错误。 使用cp命令&#xff1a; cp -Rf XXX YYY 也是出错&#xff1a; cp: 写入 xxx 出错: 输入/输出错误 2、查看磁盘空间 查看磁盘空间&#xff0c;显示空间还有剩余…
阅读更多...
【C语言】qsort()函数详解：能给万物排序的神奇函数

【C语言】qsort()函数详解：能给万物排序的神奇函数

&#x1f984;个人主页:修修修也 &#x1f38f;所属专栏:C语言 ⚙️操作环境:Visual Studio 2022 目录 一.qsort()函数的基本信息及功能 二.常见的排序算法及冒泡排序 三.逐一解读qsort()函数的参数及其原理 1.void* base 2.size_t num 3.size_t size 4.int (*compar)(c…
阅读更多...
ANN文献综述

ANN文献综述

人工神经网络文献综述 摘要 人工神经网络&#xff08;Artificial Neural Networks, ANNs&#xff09;是由多个简单的、相互连接的处理单元组成的自适应系统&#xff0c;通过调整这些单元之间的连接强度&#xff0c;ANNs能够实现对复杂数据的建模和预测。本文综述了ANNs的基本…
阅读更多...
进程、程序、应用程序之间的关系

进程、程序、应用程序之间的关系

文章目录 进程和程序进程和应用程序总结参考资料 进程和程序 程序&#xff1a;程序是存放在硬盘中的可执行文件&#xff0c;主要包括代码指令和数据。程序本身是一个静态的文件&#xff0c;只有在被操作系统加载到内存中并执行时才会变成进程。 进程&#xff1a;进程是程序在…
阅读更多...
换根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;得到两棵子树…
阅读更多...
Django自动生成Swagger接口文档 —— Python

Django自动生成Swagger接口文档 —— Python

1. 前言 当接口开发完成&#xff0c;紧接着需要编写接口文档。传统的接口文档通常都是使用Word或者一些接口文档管理平台进行编写&#xff0c;但此类接口文档维护更新比较麻烦&#xff0c;每次接口有变更&#xff0c;需要手动修改接口文档。在实际的工作中&#xff0c;经常会遇…
阅读更多...
不错的用户需求访谈方法

不错的用户需求访谈方法

不错的用户需求访谈方法&#xff0c;可以用如下的矩阵&#xff0c;用来引导用户访谈&#xff1a;
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023