CUDA C/C++ 从入门到入土 CUDA中一些小tricks

已于 2022-07-09 20:58:10 修改
阅读量697 收藏 2
点赞数 2
文章标签： c++ c语言 算法
于 2022-07-08 20:23:46 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/supersyzz/article/details/125684721
版权

本文介绍一些CUDA C/C++编程过程中的一些小工具~，无痛给自己cuda代码配置上

目录

1. 查看cuda代码运行时间，定位优化问题。

1.1 NVPROF

2. 错误检测

3.事件，用于标记，可以测算核函数运行时间

4.unified memory 简化流程，不用显示传输数据

1. 查看cuda代码运行时间，定位优化问题。

1.1 NVPROF

nvprof a.exe
nvprof --print-gpu-trace a.exe
nvprof --print-api-trace a.exe

To do :
nvprof 结合nvvp或者nsight进行可视化分析
在这里插入图片描述

2. 错误检测

在这里插入图片描述
选择第二种，封装到error.cuh中：

#pragma once
#include <stdio.h>

#define CHECK(call)                                   \
do                                                    \
{                                                     \
    const cudaError_t error_code = call;              \
    if (error_code != cudaSuccess)                    \
    {                                                 \
        printf("CUDA Error:\n");                      \
        printf("    File:       %s\n", __FILE__);     \
        printf("    Line:       %d\n", __LINE__);     \
        printf("    Error code: %d\n", error_code);   \
        printf("    Error text: %s\n",                \
            cudaGetErrorString(error_code));          \
        exit(1);                                      \
    }                                                 \
} while (0)

然后需要做三件事：

  1. 在.cu文件中 #include “error.cuh”
  2. 在.cu文件中，将所有涉及CUDA的操作用CHECK()函数封装，实例如下：
  3. 编译时加上error.cuh
    注释中用！！！check标明了调用check的地方
int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    //!!!check
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    //!!!check
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));

    // copy matrix A and B from host to device memory
    //!!!check
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
   
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);  
	CHECK(cudaGetLastError());
	//!!!check
    CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    //cudaThreadSynchronize();

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    //!!!check
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_b));
    CHECK(cudaFree(d_c));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_b));
    CHECK(cudaFreeHost(h_c));
    return 0;
}

3. 事件

作为一个标记，比较常用的用法是可以用来测算核函数运行的时间。步骤如下：
在这里插入图片描述
实例如下，注释中用！！！开头标明了流程：
note：显然事件的函数都是操作CUDA的，所以都可以用CHECK函数包装（cudaEventQuery除外）

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc;
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
    
    // !!! step 1 声明
    cudaEvent_t start, stop;
    
    // !!! step2 创建一个start 一个stop事件
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));


    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    int *d_a, *d_b, *d_c;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));


    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    
	//!!! 添加事件到执行流，可以理解start事件为标注一下此处
    CHECK(cudaEventRecord(start));
    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);    
    //!!! 用stop标记此处
    CHECK(cudaEventRecord(stop));
    //!!! 同步
    CHECK(cudaEventSynchronize(stop));
    float elapsed_time;
    // !!!计算start事件和stop事件之间的时间间隔，即cuda核函数的运行时间。
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    printf("Time = %g ms.\n", elapsed_time);
	//!!! 销毁事件
     CHECK(cudaEventDestroy(start));
     CHECK(cudaEventDestroy(stop));
    CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    //cudaThreadSynchronize();
    

    cpu_matrix_mult(h_a, h_b, h_cc, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cc[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_b));
    CHECK(cudaFree(d_c));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_b));
    CHECK(cudaFreeHost(h_c));
    return 0;
}

4.unified memory

数据要在cpu和gpu之间反复传输累不累呀~
用unified memory可以做到，一次定义，cpu和gpu都可调用！

具体步骤：

  1. 将想要同时cpu和gpu使用的变量用__managed__进行定义
  2. 删除对cuda中变量内存分配空间函数cudaMallocHost
  3. 删除原本cpu中内存分配空间函数cudaMalloc
  4. 删除cpu向gpu传输数据函数cudaMemcpy。请注意！在jetson系列中CPU和GPU集成在一个芯片，所以这种情况下，确实删除了拷贝过程。但是，在其他更多的硬件产品，CPU和GPU不在同一芯片上时，删除了传输函数不代表不传输，只不过unified memory会内部隐式的调用数据传输，我们只是把显式的数据传输给删除掉了。
  5. 删除gpu向cpu传输数据函数 cudaMemcpy。请和上面一样注意
  6. 删除cudaFree和cudaFreeHost函数。
    实例如下：
int main(int argc, char const *argv[])
{
    int m=30;
    int n=96;
    int k=160;

   // int *h_a, *h_b, *h_c, *h_cc;
    //CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
   // CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    //CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
   // CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));

    cudaEvent_t start, stop;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    

    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = rand() % 1024;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = rand() % 1024;
        }
    }

    //int *d_a, *d_b, *d_c;
    //CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
   // CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
   // CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));

    // copy matrix A and B from host to device memory
    //CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    //CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    
    CHECK(cudaEventRecord(start));
    //CHECK(cudaEventQuery(start));
    
    
    //gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m, n, k);    
    gpu_shared_mult_try2<<<dimGrid, dimBlock>>>(h_a, h_b, h_gpu, m, n, k);   
    CHECK(cudaGetLastError());
    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));
    
    float elapsed_time;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    
    printf("Time = %g ms.\n", elapsed_time);
    
    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));
    
    //CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    //cudaThreadSynchronize();

    cpu_matrix_mult(h_a, h_b, h_cpu, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    {
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cpu[i*k + j] - h_gpu[i*k + j])>(1.0e-10))
            {
                
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }

    // free memory
   // CHECK(cudaFree(d_a));
   // CHECK(cudaFree(d_b));
   // CHECK(cudaFree(d_c));
    //CHECK(cudaFreeHost(h_a));
    //CHECK(cudaFreeHost(h_b));
    //CHECK(cudaFreeHost(h_cpu));
    //CHECK(cudaFreeHost(h_gpu));
    return 0;
}
glider312
关注
  • 2
    点赞
  • 踩
  • 2
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
cuda-optimization-tips-tricks-and-techniques.pdf
11-18
《CUDA优化技巧与技术》...在进行CUDA程序优化时，开发者应遵循一种结构化的方法，从明显的性能问题开始，逐步深入到更复杂的优化技术，同时保持对性能指标的持续监控，以确保优化措施真正带来了性能提升，而不是相反。
c_transliteration_and_tricks：C89中从俄语到拉丁语的音译
03-04
"c_transliteration_and_tricks"项目专注于C89标准下的一个特殊任务——从俄语到拉丁语的音译。这个任务涉及到字符编码、字符串处理以及算法设计等多个知识点。 首先，我们需要理解C89，这是C语言的一个早期标准，...
CUDA C/C++ 优化数据传输效率
平凡
04-24 2104
设备内存和GPU之间的峰值带宽（例如，在NVIDIA Tesla C2050上为144 GB / s）比主机内存和设备内存之间的峰值带宽（在PCIe x16 Gen2上为8 GB / s）高得多。这种差异意味着在主机和GPU设备之间进行数据传输可能会影响或破坏整体的应用程序性能。 首先，让我们了解主机设-备数之间据传输的一些基本准则。 （1）尽可能减少在主机和设备之间传输的数据量，即使这意味着与在主机CPU上运行相比，在GPU上运行内核的速度几乎没有提高。 （2）使用页面锁定（或“固定”）内存时，主机
CUDA C++扩展的详细描述
lovely_yoshino的博客
02-17 7714
函数执行空间说明符表示函数是在主机上执行还是在设备上执行，以及它是可从主机调用还是从设备调用。 执行空间说明符将函数声明为内核。 它的功能是：对 函数的任何调用都必须指定其执行配置，如执行配置中所述。对 函数的调用是异步的，这意味着它在设备完成执行之前返回。 执行空间说明符声明了一个函数： 执行空间说明符声明了一个函数： 和 执行空间说明符不能一起使用。但是， 和 执行空间说明符可以一起使用，在这种情况下，该函数是为主机和设备编译的。 Application Compatibility 中引入的
C++开发基础之初探CUDA计算环境搭建
最新发布
小乖兽技术的博客
06-06 992
项目中有使用到CUDA计算的相关内容。但是在早期CUDA计算环境搭建的过程中，并不是非常顺利，编写此篇文章记录下。对于刚刚开始研究的你可能会有一定的帮助。在这个项目中，我们主要体会框架的用法，以及CUDA计算环境搭建的。通过编码实践，构建项目成功实验了CUDA计算环境搭建，为接下来的工作准备好环境。
(二)CUDA C++ 编程指南-编程模型
FlyToClouds的博客
04-03 759
学习CUDA C++编程
CUDA - CUDA C&C++简介
guaaaaaaa的博客
06-05 378
CUDA C&C++编程简介
c++和cuda混合编程总结
m0_51661400的博客
08-10 760
c++和cuda混合编程，重点介绍生成cuda动态库静态库命令并给出具体案例
tricks-and-hacks:C ++中的漂亮程序
03-09
本文将深入探讨C++中的一些关键知识点，帮助你编写更高效、更优雅的程序。 1. **模板（Templates）**: C++的模板功能允许我们创建泛型代码，适用于不同数据类型。这不仅提高了代码的重用性，还能减少冗余。例如，你...
西门子PLC例程源码S7-200中文实例Tips&tricks
03-24
西门子PLC例程源码S7_200中文实例Tips&tricks本资源系百度网盘分享地址
Tricks-html-css:实现一些小组件页面以及特效〜
03-26
以下是一些常见的CSS特效： 1. **响应式设计**：利用`@media`查询，可以根据设备屏幕尺寸调整布局，确保在不同设备上的良好显示。 2. **过渡与动画**：`transition`属性定义元素在状态改变时的平滑过渡，而`...
在c++项目中集成cuda程序
u013569930的博客
02-23 4144
1.新建或者打开一个c++项目，这里采用VS2017的一个helloworld程序作为实例， 如下： 2.在项目目录【头文件】和【源文件】下添加自己的cuda程序“.cuh”和“.cu”文件； 3.在项目上右键->生成依赖项->生成自定义，如下： 在生成自定义中勾选“CUDA 11.6(对应自己的版本)”，如下： 4.在“.cu”文件上右键，属性， 在链接器->输入->附加依赖项中，添加“cudart.lib”，如下： 5.打开菜单..
【CUDA】两幅图像叠加
hujingshuang
11-10 5674
【CUDA】两幅图像叠加
Linux 下qt creator 对c++ 和cuda混合编译生成.so 文件
weixin_39484607的博客
08-04 741
问题一 ：creator 上如何设置可以编译输出.so 文件，要做哪些工作呢？解决如下 1、TEMPLATE = lib DEFINES += TESTPCL_LIBRARY //(我的项目名 TARGET = testPCL) 2、在我要开放出去的接口类那里 做一个定义，就可以在编译的时候自动支持 -shared -fPIC 了 #include <QtCore/qglobal.h> //必要 #if defined(TESTPCL_LIBRARY)...
FFmpeg 09 overlay画中画
qq_45675910的博客
07-10 836
Video filter: overlay 描述：在指定的位置上覆盖第一个输入 语法：overlay[=x:[y[:rgb={0,1}]]] 参数描述 x：第二个视频左上角的水平坐标，默认值为0 y：第二个视频左上角的竖直坐标，默认值为0 rgb: 0 输入的颜色控件不改变，默认值；1 输入的颜色空间设置为RGB 变量-可用在x和y的表达式中 main_w/W 主要的输入宽度 main_h/H 主要的输入高度 overlay_w/w overlay输入的宽度 overlay_h/h overlay输入的高度
CUDA动态库封装以及调用
DreamerZhang的专栏
05-12 8984
参考：http://blog.sina.com.cn/s/blog_618941f701016d26.html 通过将CUDA相关计算操作放在库中，方便在项目中调用，省去了每次编译cu文件的麻烦，也便于集成到其他平台上。 本文配置：VS2015   CUDA8.0 一、封装CUDA动态库 主要步骤：修改自定义方式、设置cu文件项类型为CDUA CC++ ，添加依赖库cudart.lib.
c++ 和cuda混合编程
lily1234567的博客
12-03 7132
  装好CUDA 5.5 sdk后，默认会自动添加好系统环境变量。     因此不需要额外配置，不过为了保险起见，可以选择性地添加以下环境变量：   CUDA_BIN_PATH　　%CUDA_PATH%\bin CUDA_LIB_PATH　　%CUDA_PATH%\lib\Win32 CUDA_SDK_BIN　　%CUDA_SDK_PATH%\bin\Win32 CUDA_SDK...
transformer-xl小tricks
08-03
回答: Transformer-XL是一种改进的Transformer模型，它在处理语言模型任务时引入了一些小技巧。这些小技巧旨在提高模型的效率和性能。其中一项技巧是使用多个损失函数来改善模型的训练速度和性能。另外，Transformer-XL还增加了更多的层来提高模型的表现，最多可以达到64层。相比之下，传统的Vanilla Transformer模型在计算量和记忆性方面存在一些缺点。[2][3] 通过这些小技巧，Transformer-XL能够在处理语言模型任务时更加高效和有效地利用上下文信息。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值