【cuda学习日记】2.2 使用2维网络(grid)和2维块(block)对矩阵进行求和

news2025/1/22 9:27:22

在2.0中进行了用一维网格和块对一维向量进行了求和。
在2.1中例化了二维的网格和块。
接下来进行2维网络(grid)和2维块(block)对矩阵进行求和。

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <windows.h>

typedef unsigned long DWORD;

#define CHECK(call) \
 {\
    const cudaError_t error = call; \
    if (error != cudaSuccess)\
    {\
        printf("Error: %s: %d\n", __FILE__, __LINE__);\
        printf("code :%d reason :%s\n", error , cudaGetErrorString(error));\
        exit(1);\
    }\
}

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;
    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i])> epsilon)
        {
            match = 0;
            printf("Array do not match\n");
            printf("host %5.2f gpu % 5.2f at current %d\n", hostRef[i], gpuRef[i], i);
            break;

        }
    }
    if (match) printf("array matches\n");
}

void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int) time(&t));

    for (int i = 0; i < size; i++) {
        ip[i] = (float) (rand() & 0xff) / 10.0f;
    }
}

void sumMatrixOnHost(float *A, float *B, float *C, const int nx, const int ny){
    float *ia = A;
    float *ib = B;
    float *ic = C;

    for (int iy = 0; iy < ny; iy++)
    {
        for (int ix =0; ix < nx; ix++){
            ic[ix] = ia[ix] + ib[ix];
        }
        ia += nx; 
        ib += nx;
        ic += nx;
    }
}

__global__ void sumMatrixOnGPU2D(float *MatA, float *MatB, float *MatC, int nx, int ny){
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy*nx + ix;

    if (ix < nx && iy < ny){
        MatC[idx] = MatA[idx] + MatB[idx];
    }
}


int main(int argc , char **argv)
{
    printf("%s starting\n", argv[0]);

    int dev = 0;
    cudaDeviceProp deviceprop;
    CHECK(cudaGetDeviceProperties(&deviceprop,dev));
    printf("Using Device %d : %s\n", dev, deviceprop.name);
    CHECK(cudaSetDevice(dev));

    //set up data
    int nx  = 1<<14;
    int ny  = 1<<14;
    int nxy = nx * ny;
    size_t nBytes = nxy  * sizeof(float);
    printf("matrix size %d %d\n", nx, ny);

    float *h_A, *h_B, *hostRef, *gpuRef;

    h_A = (float *) malloc (nBytes);
    h_B = (float *) malloc (nBytes);
    hostRef = (float *) malloc (nBytes);
    gpuRef = (float *) malloc (nBytes);

    initialData(h_A, nxy);
    initialData(h_B, nxy);

    memset(hostRef,0, nBytes);
    memset(gpuRef,0, nBytes);

    // malloc device global memory
    float *d_MatA, *d_MatB, *d_MatC;
    cudaMalloc((float**)&d_MatA, nBytes);
    cudaMalloc((float**)&d_MatB, nBytes);
    cudaMalloc((float**)&d_MatC, nBytes);

    //transfer data from host to device
    cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_MatB, h_B, nBytes, cudaMemcpyHostToDevice);


    int dimx = 32;
    int dimy = 32;
    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1)/block.x, (ny + block.y - 1)/block.y);
    
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    
    cudaEventRecord(start);

    sumMatrixOnGPU2D<<<grid,block>>>(d_MatA, d_MatB, d_MatC, nx, ny);
    cudaDeviceSynchronize();
    
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);
    
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("execution config <<<(%d,%d), (%d,%d)>>>\n", grid.x,grid.y, block.x, block.y);
    printf("Kernel execution time: %f ms\n", milliseconds);

    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    
    //copy kernel result back to host
    cudaMemcpy(gpuRef, d_MatC, nBytes, cudaMemcpyDeviceToHost);

    sumMatrixOnHost(h_A, h_B, hostRef, nx,ny);

    checkResult(hostRef, gpuRef, nxy);

    cudaFree(d_MatA);
    cudaFree(d_MatB);
    cudaFree(d_MatC);

    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return 0;
}

基本流程和1维向量求和类似
输出结果:
Using Device 0 : NVIDIA GeForce RTX 4090
matrix size 16384 16384
execution config <<<(512,512), (32,32)>>>
Kernel execution time: 5.351136 ms
array matches

block的尺寸为32x32。//block(dimx,dimy)定义的。
改变block尺寸为32x16:
execution config <<<(512,1024), (32,16)>>>
Kernel execution time: 3.778752 ms

进一步改变block尺寸为16x16:
execution config <<<(1024,1024), (16,16)>>>
Kernel execution time: 3.712736 ms

在之前尝试使用nvprof测试kernl性能时,report
======= Warning: nvprof is not supported on devices with compute capability 8.0 and higher.

参考 https://blog.csdn.net/TH_NUM/article/details/109952643 使用nsys
将C:\Program Files\NVIDIA Corporation\Nsight Systems 2024.5.1\target-windows-x64加入环境变量即可

nsys profile --stats=true .\sum_matrix_on_gpu_timer.exe

输出:

Collecting data...
Generating 'C:\Users\ADMINI~1\AppData\Local\Temp\nsys-report-ffa3.qdstrm'
[1/8] [========================100%] report2.nsys-rep
[2/8] [========================100%] report2.sqlite
[3/8] Executing 'nvtx_sum' stats report
SKIPPED: C:\Users\Administrator\Desktop\edward_temp\chapter2\report2.sqlite does not contain NV Tools Extension (NVTX) data.
[4/8] Executing 'osrt_sum' stats report
SKIPPED: No data available.
[5/8] Executing 'cuda_api_sum' stats report

 Time (%)  Total Time (ns)  Num Calls   Avg (ns)     Med (ns)   Min (ns)  Max (ns)   StdDev (ns)           Name
 --------  ---------------  ---------  -----------  ----------  --------  ---------  -----------  ----------------------
     93.3        321764988          3  107254996.0  91069908.0  83897570  146797510   34432084.1  cudaMemcpy
      4.0         13772507          3    4590835.7   4393180.0   3984976    5394351     725179.5  cudaFree
      1.5          5118078          3    1706026.0   1249576.0    819401    3049101    1182856.9  cudaMalloc
      1.0          3496955          1    3496955.0   3496955.0   3496955    3496955          0.0  cudaDeviceSynchronize
      0.1           459711          1     459711.0    459711.0    459711     459711          0.0  cudaLaunchKernel
      0.0            49593          2      24796.5     24796.5       707      48886      34067.7  cudaEventCreate
      0.0            22341          1      22341.0     22341.0     22341      22341          0.0  cuLibraryUnload
      0.0            18196          2       9098.0      9098.0      7920      10276       1665.9  cudaEventRecord
      0.0            15060          1      15060.0     15060.0     15060      15060          0.0  cudaEventSynchronize
      0.0             1961          1       1961.0      1961.0      1961       1961          0.0  cuCtxSynchronize
      0.0             1434          1       1434.0      1434.0      1434       1434          0.0  cuModuleGetLoadingMode
      0.0             1012          2        506.0       506.0       205        807        425.7  cudaEventDestroy      
      0.0              181          1        181.0       181.0       181        181          0.0  cuDeviceGetLuid

[6/8] Executing 'cuda_gpu_kern_sum' stats report

 Time (%)  Total Time (ns)  Instances  Avg (ns)   Med (ns)   Min (ns)  Max (ns)  StdDev (ns)                          Name
 --------  ---------------  ---------  ---------  ---------  --------  --------  -----------  -----------------------------------------------------
    100.0          3453326          1  3453326.0  3453326.0   3453326   3453326          0.0  sumMatrixOnGPU2D(float *, float *, float *, int, int)

[7/8] Executing 'cuda_gpu_mem_time_sum' stats report

 Time (%)  Total Time (ns)  Count   Avg (ns)    Med (ns)   Min (ns)  Max (ns)  StdDev (ns)           Operation
 --------  ---------------  -----  ----------  ----------  --------  --------  -----------  ----------------------------
     68.3        180949528      2  90474764.0  90474764.0  89939258  91010270     757319.8  [CUDA memcpy Host-to-Device]
     31.7         83834368      1  83834368.0  83834368.0  83834368  83834368          0.0  [CUDA memcpy Device-to-Host]

[8/8] Executing 'cuda_gpu_mem_size_sum' stats report

 Total (MB)  Count  Avg (MB)  Med (MB)  Min (MB)  Max (MB)  StdDev (MB)           Operation
 ----------  -----  --------  --------  --------  --------  -----------  ----------------------------
   2147.484      2  1073.742  1073.742  1073.742  1073.742        0.000  [CUDA memcpy Host-to-Device]
   1073.742      1  1073.742  1073.742  1073.742  1073.742        0.000  [CUDA memcpy Device-to-Host]

Generated:
    C:\Users\Administrator\Desktop\edward_temp\chapter2\report2.nsys-rep
    C:\Users\Administrator\Desktop\edward_temp\chapter2\report2.sqlite

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.coloradmin.cn/o/2276129.html

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

相关文章

【Linux系列】`find / -name cacert.pem` 文件搜索

&#x1f49d;&#x1f49d;&#x1f49d;欢迎来到我的博客&#xff0c;很高兴能够在这里和您见面&#xff01;希望您在这里可以感受到一份轻松愉快的氛围&#xff0c;不仅可以获得有趣的内容和知识&#xff0c;也可以畅所欲言、分享您的想法和见解。 推荐:kwan 的首页,持续学…

【论文笔记】Sign Language Video Retrieval with Free-Form Textual Queries

&#x1f34e;个人主页&#xff1a;小嗷犬的个人主页 &#x1f34a;个人网站&#xff1a;小嗷犬的技术小站 &#x1f96d;个人信条&#xff1a;为天地立心&#xff0c;为生民立命&#xff0c;为往圣继绝学&#xff0c;为万世开太平。 基本信息 标题: Sign Language Video Retr…

Observability:将 OpenTelemetry 添加到你的 Flask 应用程序

作者&#xff1a;来自 Elastic jessgarson 待办事项列表可以帮助管理与假期计划相关的所有购物和任务。使用 Flask&#xff0c;你可以轻松创建待办事项列表应用程序&#xff0c;并使用 Elastic 作为遥测后端&#xff0c;通过 OpenTelemetry 对其进行监控。 Flask 是一个轻量级…

网站目录权限加固

说明 在一个入侵链路中&#xff0c;往往是利用某个安全漏洞&#xff0c;向服务器写入或上传一个webshell&#xff08;后门&#xff09;&#xff0c;再通过webshell提权或进行后续渗透入侵行为。 这个过程中&#xff0c;获取webshell是最关键最重要的一个步骤&#xff0c;如能在…

qt QPainter setViewport setWindow viewport window

使用qt版本5.15.2 引入viewport和window目的是用于实现QPainter画出来的内容随着窗体伸缩与不伸缩两种情况&#xff0c;以及让QPainter在widget上指定的区域(viewport)进行绘制/渲染&#xff08;分别对应下方demo1&#xff0c;demo2&#xff0c;demo3&#xff09;。 setViewpo…

一些计算机零碎知识随写(25年1月)-1

我原以为世界上有技术的那批人不会那么闲&#xff0c;我错了&#xff0c;被脚本真实了。 今天正隔着画画呢&#xff0c;手机突然弹出几条安全告警通知。 急忙打开服务器&#xff0c;发现问题不简单&#xff0c;直接关服务器重装系统..... 首先&#xff0c;不要认为小网站&…

分布式锁Redisson详解,Redisson如何解决不可重入,不可重试,超时释放,主从一致问题的分析解决(包括源码简单分析)

目录 1. Redisson解决不可重入锁导致的死锁问题 2. 不可重试问题 Pub/Sub 的优势 锁释放的发布逻辑 3. 超时释放的问题 1. 锁的超时释放机制背景 2. 源码分析 2.1 锁的获取 2.2 看门狗机制 2.3 看门狗续期实现 2.4 手动设置锁的过期时间 总结 4. 主从一致性 问题…

【微服务】面试 4、限流

微服务限流技术总结 一、微服务业务面试题引入 在微服务业务面试中&#xff0c;限流是重要考点&#xff0c;常与分布式事务、分布式服务接口幂等解决方案、分布式任务调度等一同被考查。面试官一般会询问项目中是否实施限流及具体做法&#xff0c;回答需涵盖限流原因、采用的方…

爬虫基础之爬取歌曲宝歌曲批量下载

声明&#xff1a;本案列仅供学习交流使用 任何用于非法用途均与本作者无关 需求分析: 网站:邓紫棋-mp3在线免费下载-歌曲宝-找歌就用歌曲宝-MP3音乐高品质在线免费下载 (gequbao.com) 爬取 歌曲名 歌曲 实现歌手名称下载所有歌曲 本案列所使用的模块 requests (发送…

树莓派-5-GPIO的应用实验之GPIO的编码方式和SDK介绍

文章目录 1 GPIO编码方式1.1 管脚信息1.2 使用场合1.3 I2C总线1.4 SPI总线2 RPI.GPIO2.1 PWM脉冲宽度调制2.2 静态函数2.2.1 函数setmode()2.2.2 函数setup()2.2.3 函数output()2.2.4 函数input()2.2.5 捕捉引脚的电平改变2.2.5.1 函数wait_for_edge()2.2.5.2 函数event_detect…

Scala分布式语言二(基础功能搭建、面向对象基础、面向对象高级、异常、集合)

章节3基础功能搭建 46.函数作为值三 package cn . itbaizhan . chapter03 // 函数作为值&#xff0c;函数也是个对象 object FunctionToTypeValue { def main ( args : Array [ String ]): Unit { //Student stu new Student() /*val a ()>{"GTJin"…

CVE-2025-22777 (CVSS 9.8):WordPress | GiveWP 插件的严重漏洞

漏洞描述 GiveWP 插件中发现了一个严重漏洞&#xff0c;该插件是 WordPress 最广泛使用的在线捐赠和筹款工具之一。该漏洞的编号为 CVE-2025-22777&#xff0c;CVSS 评分为 9.8&#xff0c;表明其严重性。 GiveWP 插件拥有超过 100,000 个活跃安装&#xff0c;为全球无数捐赠平…

支付宝租赁小程序提升租赁行业效率与用户体验

内容概要 在当今数字化的世界里&#xff0c;支付宝租赁小程序的出现构建了一种新的租赁模式&#xff0c;使得用户在使用过程中体验更加流畅。想象一下&#xff0c;你在寻找租赁服务时&#xff0c;不再需要繁琐的流程和冗长的等待&#xff0c;只需通过手机轻松点击几下&#xf…

关于使用FastGPT 摸索的QA

近期在通过fastGPT&#xff0c;创建一些基于特定业务场景的、相对复杂的Agent智能体应用。 工作流在AI模型的基础上&#xff0c;可以定义业务逻辑&#xff0c;满足输出对话之外的需求。 在最近3个月来的摸索和实践中&#xff0c;一些基于经验的小问题点&#xff08;自己也常常…

服务器/电脑与代码仓gitlab/github免密连接

git config --global user.name "xxxx" git config --global user.email "xxxxxx163.com" #使用注册GitHub的邮箱 生成对应邮箱的密码对 ssh-keygen -t rsa -b 4096 -C "xxxxxx163.com" 把公钥id_rsa.pub拷贝到github中 Setting----->…

【C语言系列】函数递归

函数递归 一、递归是什么&#xff1f;1.1尾递归 二、递归的限制条件三、递归举例3.1举例一&#xff1a;求n的阶乘3.2举例二&#xff1a;顺序打印一个整数的每一位 四、递归与迭代4.1举例三&#xff1a;求第n个斐波那契数 五、拓展学习青蛙跳台问题 一、递归是什么&#xff1f; …

springboot 默认的 mysql 驱动版本

本案例以 springboot 3.1.12 版本为例 <parent><groupId>org.springframework.boot</groupId><artifactId>spring-boot-starter-parent</artifactId><version>3.1.12</version><relativePath/> </parent> 点击 spring-…

[QCustomPlot] 交互示例 Interaction Example

本文是官方例子的分析: Interaction Example 推荐笔记: qcustomplot使用教程–基本绘图 推荐笔记: 4.QCustomPlot使用-坐标轴常用属性 官方例子需要用到很多槽函数, 这里先一次性列举, 自行加入到qt的.h中.下面开始从简单的开始一个个分析. void qcustomplot_main_init(void); …

openMetaData docker方式安装部署记录

OpenMetadata一站式元数据管理平台&#xff0c;是一款功能强大的开源元数据管理平台&#xff0c;旨在帮助企业更好地发现、理解和管理其数据资产。它提供了一套全面的工具和功能&#xff0c;涵盖了数据发现、数据血缘、数据质量、数据探查、数据治理和团队协作等多个方面。 那…

57. Three.js案例-创建一个带有聚光灯和旋转立方体的3D场景

57. Three.js案例-创建一个带有聚光灯和旋转立方体的3D场景 实现效果 该案例实现了使用Three.js创建一个带有聚光灯和旋转立方体的3D场景。 知识点 WebGLRenderer&#xff08;WebGL渲染器&#xff09; THREE.WebGLRenderer 是 Three.js 中用于将场景渲染为 WebGL 内容的核…