首页 > 其他分享 >高性能计算-CUDA单流多流控制

高性能计算-CUDA单流多流控制

时间：2025-01-07 11:57:35浏览次数：1  
标签：end hostA 单流 CudaSafecCall 传输 CUDA 流控制 deviceA

1. 介绍：

(1) 用CUDA计算 pow(sin(id),2)+ pow(cos(id),2)的结果
(2) 对比单流(同步传输、异步传输)、多流深度优先调度、多流广度优先调度的效率(包含数据传输和计算)

核心代码

1. 用CUDA计算 pow(sin(id),2)+ pow(cos(id),2)的结果
2. 对比单流(同步传输、异步传输)、多流深度优先调度、多流广度优先调度的效率(包含数据传输和计算)
3. 使用接口错误检查宏
*/

#include <stdio.h>

#define CUDA_ERROR_CHECK    //API检查控制宏

#define BLOCKSIZE 256
int N = 1<<28;              //数据个数
int NBytes = N*sizeof(float); //数据字节数


//宏定义检查API调用是否出错
#define CudaSafecCall(err) __cudaSafeCall(err,__FILE__,__LINE__)
inline void __cudaSafeCall(cudaError_t err,const char* file,const int line)
{
    #ifdef CUDA_ERROR_CHECK
    if(err!=cudaSuccess)
    {
        fprintf(stderr,"cudaSafeCall failed at %s:%d :(%d) %s\n",file,line,err,cudaGetErrorString(err));
        exit(-1);
    }
    #endif
}

//宏定义检查获取流中的执行错误，主要是对核函数
#define CudaCheckError() _cudaCheckError(__FILE__,__LINE__)
inline void _cudaCheckError(const char * file,const int line)
{
    #ifdef CUDA_ERROR_CHECK
    cudaError_t err = cudaGetLastError();
    if(err != cudaSuccess)
    {
        fprintf(stderr,"cudaCheckError failed at %s:%d :(%d) %s\n",file,line,err,cudaGetErrorString(err));
        exit(-1);
    }
    #endif
}

__global__ void kernel_func(float * arr,int offset,const int n)
{
    int id = offset + threadIdx.x + blockIdx.x * blockDim.x;
    if(id<n)
        arr[id] = pow(sinf(id),2) + pow(cosf(id),2);
}

//单流主机非锁页内存,同步传输
float gpu_base()
{
    //开辟主机非锁页内存空间
    float* hostA,*deviceA;
    hostA = (float*)calloc(N,sizeof(float));
    CudaSafecCall(cudaMalloc((void**)&deviceA,NBytes));

    
    float gpuTime = 0.0;
    cudaEvent_t start,end;
    CudaSafecCall(cudaEventCreate(&start));
    CudaSafecCall(cudaEventCreate(&end));
    CudaSafecCall(cudaEventRecord(start));
    
    CudaSafecCall(cudaMemcpy(deviceA,hostA,NBytes,cudaMemcpyHostToDevice));
    kernel_func<<<(N-1)/BLOCKSIZE + 1,BLOCKSIZE>>>(deviceA,0,N);
    CudaCheckError();

    CudaSafecCall(cudaEventRecord(end));
    CudaSafecCall(cudaEventSynchronize(end));
    CudaSafecCall(cudaEventElapsedTime(&gpuTime,start,end));
    CudaSafecCall(cudaEventDestroy(start));
    CudaSafecCall(cudaEventDestroy(end));

    CudaSafecCall(cudaMemcpy(hostA,deviceA,NBytes,cudaMemcpyDeviceToHost));

    printf("gpu_base 单流非锁页内存,数据传输和计算耗时 %f ms\n",gpuTime);
    CudaSafecCall(cudaFree(deviceA));
    free(hostA);
    return gpuTime;
}

//单流主机锁页内存,异步传输
float gpu_base_pinMem()
{
    //开辟主机非锁页内存空间
    float* hostA,*deviceA;
    CudaSafecCall(cudaMallocHost((void**)&hostA,NBytes));
    CudaSafecCall(cudaMalloc((void**)&deviceA,NBytes));
    
    float gpuTime = 0.0;
    cudaEvent_t start,end;
    CudaSafecCall(cudaEventCreate(&start));
    CudaSafecCall(cudaEventCreate(&end));
    CudaSafecCall(cudaEventRecord(start));
    
    CudaSafecCall(cudaMemcpy(deviceA,hostA,NBytes,cudaMemcpyHostToDevice));
    kernel_func<<<(N-1)/BLOCKSIZE + 1,BLOCKSIZE>>>(deviceA,0,N);
    CudaCheckError();

    CudaSafecCall(cudaEventRecord(end));
    CudaSafecCall(cudaEventSynchronize(end));
    CudaSafecCall(cudaEventElapsedTime(&gpuTime,start,end));
    CudaSafecCall(cudaEventDestroy(start));
    CudaSafecCall(cudaEventDestroy(end));

    CudaSafecCall(cudaMemcpy(hostA,deviceA,NBytes,cudaMemcpyDeviceToHost));

    printf("gpu_base_pinMem 单流锁页内存,数据传输和计算耗时 %f ms\n",gpuTime);

    CudaSafecCall(cudaFreeHost(hostA));
    CudaSafecCall(cudaFree(deviceA));
    return gpuTime;
}

//多流深度优先调度
float gpu_MStream_deep(int nStreams)
{
    //开辟主机非锁页内存空间
    float* hostA,*deviceA;
    //异步传输必须用锁页主机内存
    CudaSafecCall(cudaMallocHost((void**)&hostA,NBytes));
    CudaSafecCall(cudaMalloc((void**)&deviceA,NBytes));
    
    float gpuTime = 0.0;
    cudaEvent_t start,end;
    cudaStream_t* streams = (cudaStream_t*)calloc(nStreams,sizeof(cudaStream_t));
    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamCreate(streams+i));
    CudaSafecCall(cudaEventCreate(&start));
    CudaSafecCall(cudaEventCreate(&end));
    CudaSafecCall(cudaEventRecord(start));
    
    //传输、计算,流间最多只有一个传输和一个计算同时进行
    // 每个流处理的数据量
    int nByStream = N/nStreams;
    for(int i=0;i<nStreams;i++)
    {
        int offset = i * nByStream;
        CudaSafecCall(cudaMemcpyAsync(deviceA+offset,hostA+offset,nByStream*sizeof(float),cudaMemcpyHostToDevice,streams[i]));
        kernel_func<<<(nByStream-1)/BLOCKSIZE + 1,BLOCKSIZE,0,streams[i]>>>(deviceA,offset,(i+1)*nByStream);
        CudaCheckError();
        CudaSafecCall(cudaMemcpyAsync(hostA+offset,deviceA+offset,nByStream*sizeof(float),cudaMemcpyDeviceToHost,streams[i]));
    }

    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamSynchronize(streams[i]));

    CudaSafecCall(cudaEventRecord(end));
    CudaSafecCall(cudaEventSynchronize(end));
    CudaSafecCall(cudaEventElapsedTime(&gpuTime,start,end));
    CudaSafecCall(cudaEventDestroy(start));
    CudaSafecCall(cudaEventDestroy(end));

    printf("gpu_MStream_deep %d个流深度优先调度,数据传输和计算耗时 %f ms\n",nStreams,gpuTime);

    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamDestroy(streams[i]));

    CudaSafecCall(cudaFreeHost(hostA));
    CudaSafecCall(cudaFree(deviceA));
    free(streams);
    return gpuTime;
}

//多流广度优先调度
float gpu_MStream_wide(int nStreams)
{
    //开辟主机非锁页内存空间
    float* hostA,*deviceA;
    //异步传输必须用锁页主机内存
    CudaSafecCall(cudaMallocHost((void**)&hostA,NBytes));
    CudaSafecCall(cudaMalloc((void**)&deviceA,NBytes));
    
    float gpuTime = 0.0;
    cudaEvent_t start,end;
    cudaStream_t* streams = (cudaStream_t*)calloc(nStreams,sizeof(cudaStream_t));
    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamCreate(streams+i));
    CudaSafecCall(cudaEventCreate(&start));
    CudaSafecCall(cudaEventCreate(&end));
    CudaSafecCall(cudaEventRecord(start));
    
    //传输、计算,流间并行
    // 每个流处理的数据量
    int nByStream = N/nStreams;
    for(int i=0;i<nStreams;i++)
    {
        int offset = i * nByStream;
        CudaSafecCall(cudaMemcpyAsync(deviceA+offset,hostA+offset,nByStream*sizeof(float),cudaMemcpyHostToDevice,streams[i]));
    }
    for(int i=0;i<nStreams;i++)
    {
        int offset = i * nByStream;
        kernel_func<<<(nByStream-1)/BLOCKSIZE + 1,BLOCKSIZE,0,streams[i]>>>(deviceA,offset,(i+1)*nByStream);
        CudaCheckError();
    }
    for(int i=0;i<nStreams;i++)
    {
        int offset = i * nByStream;
        CudaSafecCall(cudaMemcpyAsync(hostA+offset,deviceA+offset,nByStream*sizeof(float),cudaMemcpyDeviceToHost,streams[i]));
    }

    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamSynchronize(streams[i]));

    CudaSafecCall(cudaEventRecord(end));
    CudaSafecCall(cudaEventSynchronize(end));
    CudaSafecCall(cudaEventElapsedTime(&gpuTime,start,end));
    CudaSafecCall(cudaEventDestroy(start));
    CudaSafecCall(cudaEventDestroy(end));

    printf("gpu_MStream_wide %d个流广度优先调度,数据传输和计算耗时 %f ms\n",nStreams,gpuTime);

    for(int i=0;i<nStreams;i++)
        CudaSafecCall(cudaStreamDestroy(streams[i]));

    CudaSafecCall(cudaFreeHost(hostA));
    CudaSafecCall(cudaFree(deviceA));
    free(streams);
    return gpuTime;
}

int main(int argc,char* argv[])
{
    int nStreams = argc==2? atoi(argv[1]):4;

    //gpu默认单流,主机非锁页内存,同步传输
    float gpuTime1 = gpu_base();

    //gpu默认单流,主机锁页内存,异步传输
    float gpuTime2 = gpu_base_pinMem();

    //gpu多流深度优先调度,异步传输
    float gpuTime3 = gpu_MStream_deep(nStreams);

    //gpu多流广度优先调度,异步传输
    float gpuTime4 = gpu_MStream_wide(nStreams);

    printf("相比默认单流同步传输与计算,单流异步传输及运算加速比为 %f\n",nStreams,gpuTime1/gpuTime2);
    printf("相比默认单流同步传输与计算,%d 个流深度优先调度异步传输及运算加速比为 %f\n",nStreams,gpuTime1/gpuTime3);
    printf("相比默认单流同步传输与计算,%d 个流广度优先调度异步传输及运算加速比为 %f\n",nStreams,gpuTime1/gpuTime4);
    return 0;
}

3. 测试结果

各项测试耗时及与单流同步传输和计算加速比

项目\流数量 4 8 16 32 64
单流同步传输(耗时ms) 312.43 298.41 305.48 306.54 310.75
单流异步传输(耗时ms/加速比) 197.15/1.58 195.89/1.52 201.88/1.51 202.87/1.51 201.53/1.54
多流深度优先调度(耗时ms/加速比) 151.04/2.06 129.95/2.29 131.49/2.32 123.08/2.49 126.48/2.45
多流广度优先调度(耗时ms/加速比) 149.29/2.09 129.6/2.3 134.55/2.27 122.82/2.49 126.42/2.45

4. 结果分析

(1) 单流异步传输比同步传输明显效率更高，这是因为同步传输PCIE 通过 DMA 只能访问锁页内存，同步传输使用的主机内存地址是虚拟非锁页内存地址，相比异步传输同步传输额外增加了非锁页向锁页内存转换的开销；

(2) 多流的连个该测试中随着流数量的增加，加速比会提升；

(3) 多流广度优先相比深度优先调度在 228数据规模下效率几乎一致，可能因为数据规模较大，硬件资源紧张无法真正实现多流并发的优势。经多次测试，使用数据规模220，流2-8个时 ,广度优先的加速比能提升 2%左右，随着流数的增加广度优先效率反而不如深度优先。

标签：end,hostA,单流,CudaSafecCall,传输,CUDA,流控制,deviceA
From： https://www.cnblogs.com/anluo8/p/18656498

相关文章

  • cuda kernel启动的反汇编
    原始代码//Typeyourcodehere,orloadanexample.extern"C"__global__voidsquare(int*array,intn){inttid=blockDim.x*blockIdx.x+threadIdx.x;if(tid<n)array[tid]=array[tid]*array[tid];}voidsquare_do(int......
  • Linux服务器无Root权限安装Cuda方法及问题解决
    CUDA简介什么是CUDA？CUDA（ComputeUnifiedDeviceArchitecture）是由NVIDIA提供的一种并行计算平台和编程模型，用于加速计算密集型任务。CUDA允许开发者使用GPU的计算能力，通过并行处理来快速执行复杂的计算任务。CUDA包括以下主要组成部分：CUDAToolkit：为开发人员提供工......
  • 深度学习CUDA环境安装教程---动手学深度学习
    首先说明我安装的是《动手学深度学习》中的环境本人是小白，一次安装，可能有不对的地方，望包含。安装CUDA因为我们是深度学习，很多时候要用到gpu进行训练，所以我们需要一种方式加快训练速度。通俗地说，CUDA是一种协助“CPU任务分发+GPU并行处理”的编程模型/平台，用于加速GPU和CPU之......
  • 高性能计算-CUDA矩阵加法及优化测试
    1.目标：对16384*16384规模的矩阵进行加法运算，对比CPU和GPU计算的效率，还有不同线程块大小规模下对效率的影响；并做可能的优化测试。2.核心代码/*用GPU对二维矩阵做加法，分析不同线程块规模下的性能变化*/#include<stdio.h>#include<stdlib.h>#include<sys/time.h>#......
  • CUDA编程【5】获取GPU设备信息
    文章目录通过cudaAPI获取1.获取设备数量2.获取当前设备ID3.设置当前设备4.获取设备属性5.获取设备限制6.获取设备共享内存配置7.获取设备缓存配置8.获取设备是否支持统一内存9.获取设备是否支持并发内核执行10.获取设备的最大线程块数11.获取设备的时钟频率......
  • 编译CUDA时的ARCH参数
    https://blog.csdn.net/Vingnir/article/details/135255072在编译CUDA程序时，ARCH是指定给nvcc（NVIDIACUDACompiler）的一个重要参数。ARCH代表着目标GPU的计算能力（ComputeCapability），这是一个特定于NVIDIAGPU架构的指标，用于表明GPU支持的特性和指令集。关于CUDA计算能力（Com......
  • 折腾笔记[4]-cuda的hello-world
    摘要在window11上搭建cuda开发环境并编译helloworld程序;关键信息编译器:cudanvcc12.4.131平台:windows11原理简介cuda简介CUDA(ComputeUnifiedDeviceArchitecture，统一计算架构)是由英伟达所推出的一种集成技术，向用户提供了可以很优雅地调用GPU进行并行计算的编程......
  • 利用CUDA编程实现在GPU中对图像的极坐标变换加速
    问题来源：1.需要对输入图像中的一个环形区域，进行极坐标逆变换，将该环形区域转换为一张新的矩形图像2.opencv没有直接对环形区域图像进行变换的函数，需要通过循环遍历的方式，利用polarToCart进行转换3.循环遍历不可避免的带来速度上的问题，尤其是图片较大时解决思路1：使用open......
  • 【CUDA】cuDNN：加速深度学习的核心库
    【CUDA】cuDNN：加速深度学习的核心库1.什么是cuDNN？cuDNN（CUDADeepNeuralNetworklibrary）是NVIDIA提供的一个高性能GPU加速库，专为深度学习任务中常见的操作进行了高度优化。它不仅提供了单操作的高效实现，还支持多操作融合（fusion），旨在最大化地利用NVIDIAGPU的计算......
  • 十亿行数据挑战：CUDA申请出战（从17分钟到17秒）
    文章结尾有最新热度的文章，感兴趣的可以去看看。本文是经过严格查阅相关权威文献和资料，形成的专业的可靠的内容。全文数据都有据可依，可回溯。特别申明：数据和资料已获得授权。本文内容，不涉及任何偏颇观点，用中立态度客观事实描述事情本身导读在我学习CUDA的过程中，我决定用......

赞助商

阅读排行

网站主页关于我们联系我们网站地图

本网站内容转载自其他媒体，侵权联系[admin##ips99.com]。

Copyright © 2020-2023 IPS99 版权所有 IPS99