【CUDA】CUDA中缓存机制对计时的影响

news2024/11/13 23:13:29

笔者在阅读知乎上一个关于CUDA编程的专栏时,发现作者写的很多文章中都会附带计时的模块用于评估程序的运行效率,然而笔者发现,在运行这篇文章中的代码时时,得到的结果和作者的结果有较大差异,主要体现在:使用第三种thrust库的方法得到的结果比第二种使用共享内存的结果要差。
观察每次的运行时间,发现第三种方法首次运行的时间大约是后续运行时间的7倍,而第二种方法的每次运行时间基本一致,因此造成了结果的不一致,笔者将文中的代码进行修改,展示两个版本,其中version1是前20次的结果(即原文结果),version2是第1次到第21次的运行结果(即后20次的结果,不包括第一次),可以看到,此时第三种方法的结果比第二种方法的结果快,且三种方法的结果差距不大,那么可以说在计算层面,对该程序的优化提升不大,而在访存方面,thrust使用的方法效率要比共享内存低。

#include <stdio.h>
#include "error.cuh"
#include <thrust/scan.h>
#include <thrust/execution_policy.h>

#ifdef USE_DP
    typedef double real;
#else
    typedef float real;
#endif


const int N = 1000000;
const int M = sizeof(int) * N;
const int NUM_REAPEATS = 20;
const int BLOCK_SIZE = 1024;
const int GRID_SIZE = (N - 1) / BLOCK_SIZE + 1;


__global__ void globalMemScan(real *d_x, real *d_y) {
    real *x = d_x + blockDim.x * blockIdx.x;
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    real y = 0.0;
    if (n < N) {
        for (int offset = 1; offset < blockDim.x; offset <<= 1) {
            if (threadIdx.x >= offset) y = x[threadIdx.x] + x[threadIdx.x - offset];
            __syncthreads();
            if (threadIdx.x >= offset) x[threadIdx.x] = y;
        }
        if (threadIdx.x == blockDim.x - 1) d_y[blockIdx.x] = x[threadIdx.x];
    } 
}

__global__ void addBaseValue(real *d_x, real *d_y) {
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    real y = blockIdx.x > 0 ? d_y[blockIdx.x - 1] : 0.0;
    if (n < N) {
        d_x[n] += y;
    } 
}

__global__ void sharedMemScan(real *d_x, real *d_y) {
    extern __shared__ real s_x[];
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    s_x[threadIdx.x] = n < N ? d_x[n] : 0.0;
    __syncthreads();
    real y = 0.0;
    if (n < N) {
        for (int offset = 1; offset < blockDim.x; offset <<= 1) {
            if (threadIdx.x >= offset) y = s_x[threadIdx.x] + s_x[threadIdx.x - offset];
            __syncthreads();
            if (threadIdx.x >= offset) s_x[threadIdx.x] = y;
        }
        d_x[n] = s_x[threadIdx.x];
        if (threadIdx.x == blockDim.x - 1) d_y[blockIdx.x] = s_x[threadIdx.x];
    } 
}


void scan(real *d_x, real *d_y, real *d_z, const int method) {
    switch (method)
    {
    case 0:
        globalMemScan<<<GRID_SIZE, BLOCK_SIZE>>>(d_x, d_y);
        globalMemScan<<<1, GRID_SIZE>>>(d_y, d_z);
        addBaseValue<<<GRID_SIZE, BLOCK_SIZE>>>(d_x, d_y);
        break;
    case 1:
        sharedMemScan<<<GRID_SIZE, BLOCK_SIZE, sizeof(real) * BLOCK_SIZE>>>(d_x, d_y);
        sharedMemScan<<<1, GRID_SIZE, sizeof(real) * GRID_SIZE>>>(d_y, d_z);
        addBaseValue<<<GRID_SIZE, BLOCK_SIZE>>>(d_x, d_y);
        break;
    case 2:
        thrust::inclusive_scan(thrust::device, d_x, d_x + N, d_x);
        break;
    
    default:
        break;
    }
}

void timing(const real *h_x, real *d_x, real *d_y, real *d_z, real *h_ret, const int method) {

    float tSum = 0.0;
    float t2Sum = 0.0;
    float tSumVersion2 = 0.0;
    float t2SumVersion2 = 0.0;
    for (int i=0; i<=NUM_REAPEATS; ++i) {
        CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
        cudaEvent_t start, stop;
        CHECK(cudaEventCreate(&start));
        CHECK(cudaEventCreate(&stop));
        CHECK(cudaEventRecord(start));
        cudaEventQuery(start);

        scan(d_x, d_y, d_z, method);

        CHECK(cudaEventRecord(stop));
        CHECK(cudaEventSynchronize(stop));
        float elapsedTime;
        CHECK(cudaEventElapsedTime(&elapsedTime, start, stop));
        if(i == 0) {
            // do nothing
        } else {
            tSumVersion2 += elapsedTime;
            t2SumVersion2 += elapsedTime * elapsedTime;
        }

        if(i == NUM_REAPEATS) {
            // do nothing
        } else {
            tSum += elapsedTime;
            t2Sum += elapsedTime * elapsedTime;
        }
        
        printf("%g\t", elapsedTime);
        CHECK(cudaEventDestroy(start));
        CHECK(cudaEventDestroy(stop));
    }
    printf("\n======version1======\n");
    printf("\n%g\t", tSum);
    float tAVG = tSum / NUM_REAPEATS;
    float tERR = sqrt(t2Sum / NUM_REAPEATS - tAVG * tAVG);
    printf("Time = %g +- %g ms.\n", tAVG, tERR);
    printf("\n======version2======\n");
    printf("\n%g\t", tSumVersion2);
    float tAVGVersion2 = tSumVersion2 / NUM_REAPEATS;
    float tERRVersion2 = sqrt(t2SumVersion2 / NUM_REAPEATS - tAVGVersion2 * tAVGVersion2);
    printf("Time = %g +- %g ms.\n", tAVGVersion2, tERRVersion2);
    CHECK(cudaMemcpy(h_ret, d_x, M, cudaMemcpyDeviceToHost));
}

int main() {
    real *h_x = new real[N];
    real *h_y = new real[N];    
    real *h_ret = new real[N];
    for (int i=0; i<N; i++) h_x[i] = 1.23;
    real *d_x, *d_y, *d_z;
    CHECK(cudaMalloc((void **)&d_x, M));
    CHECK(cudaMalloc((void **)&d_y, sizeof(real) * GRID_SIZE));
    CHECK(cudaMalloc((void **)&d_z, sizeof(real)));
    
    printf("using global mem:\n");
    timing(h_x, d_x, d_y, d_z, h_ret, 0);
    for (int i = N - 10; i < N; i++) printf("%f  ", h_ret[i]);
    printf("\n");
    printf("using shared mem:\n");
    timing(h_x, d_x, d_y, d_z, h_ret, 1);
    for (int i = N - 10; i < N; i++) printf("%f  ", h_ret[i]);
    printf("\n");
    printf("using thrust lib:\n");
    timing(h_x, d_x, d_y, d_z, h_ret, 2);
    for (int i = N - 10; i < N; i++) printf("%f  ", h_ret[i]);
    printf("\n");

    CHECK(cudaFree(d_x));
    CHECK(cudaFree(d_y));
    CHECK(cudaFree(d_z));
    delete[] h_x;
    delete[] h_y;
    delete[] h_ret;
}

贴一张运行结果:
在这里插入图片描述
局部放大图:
在这里插入图片描述
可以看到,三种方法在计算层面基本上都稳定到0.03-0.04ms这个数量级。

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

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

相关文章

《战甲神兵》开发者报告:游戏崩溃问题80%发生在Intel可超频酷睿i9处理器上——酷睿i7 K系列CPU也表现出高崩溃率

在Intel持续面临第13代和第14代CPU崩溃问题的背景下&#xff0c;近日&#xff0c;《战甲神兵》(Warframe)的开发者们于7月9日披露了游戏崩溃的统计数据&#xff0c;并描述了诊断该问题的过程。根据开发团队的说法&#xff0c;一名未进行超频且使用全新PC的员工&#xff0c;即便…

网络安全设备——EDR

网络安全中的EDR&#xff08;Endpoint Detection and Response&#xff0c;端点检测与响应&#xff09;是一种主动式的端点安全解决方案&#xff0c;它专注于监控、检测和响应计算机和终端设备上的安全威胁。以下是EDR的详细解释&#xff1a; 一、定义与功能 EDR是一种网络安…

【C++】入门基础(引用、inline、nullptr)

目录 一.引用 1.引用的定义 2.引用的特性 3.引用的使用场景 4.const引用 5.引用和指针的区别 二.inline 三.nullptr 一.引用 1.引用的定义 引用不是新定义一个变量&#xff0c;而是给已经存在的变量取一个别名&#xff0c;编译器不会给引用变量开辟内存空间&#xff0c…

(实测可用)(3)Git的使用——RT Thread Stdio添加的软件包,github与gitee冲突造成无法上传文件到gitee

★硬件资源&#xff1a;本文章以STM32L431RCT6做主控芯片做验证&#xff1b; ★IDE开发环境&#xff1a;RT Thread stdio&#xff1b; ★RT Thread 版本&#xff1a;V4.0.3 一、RT Thread Stdio加载软件包 1、如下图所示&#xff0c;通过RT Thread Stdio加载的软件包&#…

[FPGA]-时序传输模型分析

时序传输模型分析 FPGA内部时钟树 clk到达每个寄存器的时间不一致。 内部时钟树内部示意图如下所示&#xff1a; 在实际FPGA芯片内部&#xff0c;时钟到达每个寄存器的时钟偏差很小&#xff0c;但仍然存在&#xff1b;比如clk到达REG1花费时间0.2ns&#xff0c;到达REG6花费…

【Linux】基于环形队列RingQueue的生产消费者模型

提示&#xff1a;文章写完后&#xff0c;目录可以自动生成&#xff0c;如何生成可参考右边的帮助文档 目录 前言 环形队列的概念及定义 POSIX信号量 RingQueue的实现方式 RingQueue.hpp的构建 Thread.hpp Main.cc主函数的编写 Task.hpp function包装器的使用 总结 前言…

Torch-Pruning 库入门级使用介绍

项目地址&#xff1a;https://github.com/VainF/Torch-Pruning Torch-Pruning 是一个专用于torch的模型剪枝库&#xff0c;其基于DepGraph 技术分析出模型layer中的依赖关系。DepGraph 与现有的修剪方法&#xff08;如 Magnitude Pruning 或 Taylor Pruning&#xff09;相结合…

Python:setattr()函数和__setattr__()魔术方法

相关阅读 Pythonhttps://blog.csdn.net/weixin_45791458/category_12403403.html?spm1001.2014.3001.5482 setattr()是一个python内置的函数&#xff0c;用于设置一个对象的属性值&#xff0c;一般情况下&#xff0c;可以通过点运算符(.)完成相同的功能&#xff0c;但是getat…

[激光原理与应用-112]:南京科耐激光-激光焊接-焊中检测-智能制程监测系统IPM介绍 - 16 - 常见的产品指标

目录 一、光学传感器技术指标&#xff1a;实时信号采集与信号处理 &#xff08;1&#xff09;适用激光器的功率范围宽&#xff1a; &#xff08;2&#xff09;感光范围&#xff1a;350nm~1750nm&#xff1a;从可见光到红外光 &#xff08;3&#xff09;信号类型&#xff1a…

NSSCTF_RE(二)暑期

[CISCN 2021初赛]babybc LLVM是那个控制流平坦化混淆&#xff0c;bc是IR指令文件 得到64位elf文件 然后就慢慢分析&#xff0c;感觉太妙了我靠 一个数独游戏&#xff0c;用二个二维数组添加约束&#xff0c;一个二维数组作地图&#xff0c;慢慢看 最后用 z3 来解数独&#xf…

k8s快速部署一个网站

1&#xff09;使用Deployment控制器部署镜像&#xff1a; kubectl create deployment web-demo --imagelizhenliang/web-demo:v1 kubectl get deployment,pods[rootk8s-matser ~]# kubectl get pods NAME READY STATUS RESTARTS A…

25届平安产险校招测评IQ新16PF攻略:全面解析与应试策略

尊敬的读者&#xff0c;您好。随着平安产险校招季的到来&#xff0c;许多应届毕业生正积极准备着各项测评。本文旨在提供一份详尽的测评攻略&#xff0c;帮助您更好地理解平安产险的校招测评流程&#xff0c;以及如何有效应对。 25届平安产险平安IQ&#xff08;新&#xff09;测…

Java 设计模式系列:外观模式

简介 外观模式&#xff08;Facade Pattern&#xff09;是一种设计模式&#xff0c;又名门面模式&#xff0c;是一种通过为多个复杂的子系统提供一个一致的接口&#xff0c;而使这些子系统更加容易被访问的模式。该模式对外有一个统一接口&#xff0c;外部应用程序不用关心内部…

2024-07-14 Unity插件 Odin Inspector2 —— Essential Attributes

文章目录 1 说明2 重要特性2.1 AssetsOnly / SceneObjectsOnly2.2 CustomValueDrawer2.3 OnValueChanged2.4 DetailedInfoBox2.5 EnableGUI2.6 GUIColor2.7 HideLabel2.8 PropertyOrder2.9 PropertySpace2.10 ReadOnly2.11 Required2.12 RequiredIn&#xff08;*&#xff09;2.…

基于Python thinker GUI界面的股票评论数据及投资者情绪分析设计与实现

1.绪论 1.1背景介绍 Python 的 Tkinter 库提供了创建用户界面的工具&#xff0c;可以用来构建股票评论数据及投资者情绪分析的图形用户界面&#xff08;GUI&#xff09;。通过该界面&#xff0c;用户可以输入股票评论数据&#xff0c;然后通过情感分析等技术对评论进行情绪分析…

昇思25天学习打卡营第14天 | ShuffleNet图像分类

昇思25天学习打卡营第14天 | ShuffleNet图像分类 文章目录 昇思25天学习打卡营第14天 | ShuffleNet图像分类ShuffleNetPointwise Group ConvolutionChannel ShuffleShuffleNet模块网络构建 模型训练与评估数据集训练模型评估模型预测 总结打卡 ShuffleNet ShuffleNetV1是旷世科…

大模型系列3--pytorch dataloader的原理

pytorch dataloader运行原理 1. 背景2. 环境搭建2.1. 安装WSL & vscode2.2. 安装conda & pytorch_gpu环境 & pytorch 2.112.3 命令行验证python环境2.4. vscode启用pytorch_cpu虚拟环境 3. 调试工具3.1. vscode 断点调试3.2. py-spy代码栈探测3.3. gdb attach3.4. …

基于锚框的物体检测过程

说明&#xff1a;基于锚框的物体检测过程&#xff1a;分为单阶段和两阶段 整体步骤&#xff1a; 提供目标候选区域&#xff1a; 锚框提供了一组预定义的候选区域&#xff0c;这些区域可以覆盖各种尺度和长宽比的目标。通过这些锚框&#xff0c;可以在不同的位置和不同的尺度上…

02-Charles的安装与配置

一、Charles的安装 Charles的下载地址&#xff1a;https://www.charlesproxy.com/。 下载之后&#xff0c;傻瓜式安装即可。 二、Charles组件介绍 主导航栏介绍&#xff1a; 请求导航栏介绍&#xff1a; 请求数据栏介绍&#xff1a; 三、Charles代理设置 四、客户端-windows代理…

【Linux】多线程_6

文章目录 九、多线程7. 生产者消费者模型生产者消费者模型的简单代码结果演示 未完待续 九、多线程 7. 生产者消费者模型 生产者消费者模型的简单代码 Makefile&#xff1a; cp:Main.ccg -o $ $^ -stdc11 -lpthread .PHONY:clean clean:rm -f cpThread.hpp&#xff1a; #i…