cuda核函数 调用函数_CUDA编程

CUDA编程优秀博客收集

希葛格的韩少君：CMake之构建cuda应用程序​zhuanlan.zhihu.com

CUDA Toolkit Documentation​docs.nvidia.com

ZihaoZhao：CUDA编程入门（一）CUDA编程模型​zhuanlan.zhihu.com

小小将：CUDA编程入门极简教程​zhuanlan.zhihu.com

1. 线程管理

线程网格和线程块从逻辑上代表了一个核函数的线程层次结构，这种组织方式可以帮助我们有效地利用资源，优化性能

合适的线程网格和线程块大小来正确地组织线程，内核的性能可以得到大大地提高

GPU中线程的组织结构

2.内存管理

CUDA编程另一个显著的特点就是解释了内存层次结构，每一个GPU设备都会有用于不同用途的存储类型。

只需先关注这三种内存类型：寄存器（Registers）、共享内存（Shared Memory）和全局内存（Global Memory）

2.1、寄存器是GPU上运行速度最快的内存空间，通常其带宽为8TB/s左右，延迟为1个时钟周期

2.2、共享内存是GPU上可受用户控制的一级缓存，对于共享内存的使用，主要考虑数据的重用性。当存在着数据的重复利用时，使用共享内存是比较合适的。如果数据不被重用，则直接将数据从全局内存或常量内存读入寄存器即可。

2.3 全局内存是GPU中最大、延迟最高并且最常使用的内存。全局内存类似于CPU的系统内存。在编程中对全局内存访问的优化以最大化程度提高全局内存的数据吞吐量是十分重要的。

GPU内存层次结构

作者：ZihaoZhao
链接：https://zhuanlan.zhihu.com/p/97192227
来源：知乎
著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。

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


//CPU对照组，用于对比加速比
void sumMatrix2DonCPU(float * MatA,float * MatB,float * MatC,int nx,int ny)
{
    float* a = MatA;
    float* b = MatB;
    float* c = MatC;
    for(int j=0; j<ny; j++)
    {
        for(int i=0; i<nx; i++)
        {
          c[i] = a[i]+b[i];
        }
        c += nx;
        b += nx;
        a += nx;
    }
}

//核函数，每一个线程计算矩阵中的一个元素。
__global__ void sumMatrix(float * MatA,float * MatB,float * MatC,int nx,int ny)
{
    int ix = threadIdx.x+blockDim.x*blockIdx.x;
    int iy = threadIdx.y+blockDim.y*blockIdx.y;
    int idx = ix+iy*ny;
    if (ix<nx && iy<ny)
    {
        MatC[idx] = MatA[idx]+MatB[idx];
    }
}

//主函数
int main(int argc,char** argv)
{
    //设备初始化
    printf("strating...n");
    initDevice(0);

    //输入二维矩阵，4096*4096，单精度浮点型。
    int nx = 1<<12;
    int ny = 1<<12;
    int nBytes = nx*ny*sizeof(float);

    //Malloc，开辟主机内存
    float* A_host = (float*)malloc(nBytes);
    float* B_host = (float*)malloc(nBytes);
    float* C_host = (float*)malloc(nBytes);
    float* C_from_gpu = (float*)malloc(nBytes);
    initialData(A_host, nx*ny);
    initialData(B_host, nx*ny);

    //cudaMalloc，开辟设备内存
    float* A_dev = NULL;
    float* B_dev = NULL;
    float* C_dev = NULL;
    CHECK(cudaMalloc((void**)&A_dev, nBytes));
    CHECK(cudaMalloc((void**)&B_dev, nBytes));
    CHECK(cudaMalloc((void**)&C_dev, nBytes));

    //输入数据从主机内存拷贝到设备内存
    CHECK(cudaMemcpy(A_dev, A_host, nBytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(B_dev, B_host, nBytes, cudaMemcpyHostToDevice));

    //二维线程块，32×32
    dim3 block(32, 32);
    //二维线程网格，128×128
    dim3 grid((nx-1)/block.x+1, (ny-1)/block.y+1);

    //测试GPU执行时间
    double gpuStart = cpuSecond();
    //将核函数放在线程网格中执行
    sumMatrix<<<grid,block>>>(A_dev, B_dev, C_dev, nx, ny);
    CHECK(cudaDeviceSynchronize());
    double gpuTime = cpuSecond() - gpuStart;
    printf("GPU Execution Time: %f secn", gpuTime);

    //在CPU上完成相同的任务
    cudaMemcpy(C_from_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost);
    double cpuStart=cpuSecond();
    sumMatrix2DonCPU(A_host, B_host, C_host, nx, ny);
    double cpuTime = cpuSecond() - cpuStart;
    printf("CPU Execution Time: %f secn", cpuTime);

    //检查GPU与CPU计算结果是否相同
    CHECK(cudaMemcpy(C_from_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost));
    checkResult(C_host, C_from_gpu, nx*ny);

    //释放内存
    cudaFree(A_dev);
    cudaFree(B_dev);
    cudaFree(C_dev);
    free(A_host);
    free(B_host);
    free(C_host);
    free(C_from_gpu);
    cudaDeviceReset();
    return 0;
}