【nvidia CUDA 高级编程】NVSHMEM 直方图——复制式方法

news2024/11/22 6:04:17

博主未授权任何人或组织机构转载博主任何原创文章,感谢各位对原创的支持!
博主链接

本人就职于国际知名终端厂商,负责modem芯片研发。
在5G早期负责终端数据业务层、核心网相关的开发工作,目前牵头6G算力网络技术标准研究。


博客内容主要围绕:
       5G/6G协议讲解
       算力网络讲解(云计算,边缘计算,端计算)
       高级C语言讲解
       Rust语言讲解



NVSHMEM 直方图——复制式方法

在这里插入图片描述

PE:处理单元(process entity)

直方图简介

       我们来了解一个与之前的问题有些相似但却稍显复杂的问题:构造直方图。也就是说,给定一个具有 𝑁 个整数的数组,和 𝑀 个取值范围,计算数组中有多少元素属于 𝑀 个范围中的某一个。不失一般性,我们将指定整数是 [0,𝐾−1] 区间内的正数,范围或桶是均匀的线性间隔(为了简单起见, 𝐾 可以被 𝑀 整除),从而第一个桶可以覆盖 [0,𝐾/𝑀−1] 区间内的数,第二个可以覆盖 [𝐾/𝑀,2𝐾/𝑀−1] 区间内的数,以此类推。

       在我们开始为多个 GPU 重构代码之前,我们将从使用单 GPU 的代码示例入手。解决这个问题最简单的方法还是使用原子操作。我们将对数组进行循环,计算数组中的每个元素应落入哪个桶;给定一个整数 𝑛 ,其所属的直方图数组的索引为 (𝑛𝑀)/𝐾 。然后,以原子操作增加该桶的计数器。检查代码,然后运行代码,看看会得到什么样的输出。您可按需随意调整参数(但要避免数字过大,注意 32 位整数溢出)。

#include <iostream>
#include <cstdlib>

inline void CUDA_CHECK (cudaError_t err) {
    if (err != cudaSuccess) {
        fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
        exit(-1);
    }
}

#define NUM_BUCKETS   16
#define MAX_VALUE     1048576
#define NUM_INPUTS    65536

__global__ void histogram_kernel(const int* input, int* histogram, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    if (idx < N) {
        int value = input[idx];

        int histogram_index = (value * NUM_BUCKETS) / MAX_VALUE;

        atomicAdd(&histogram[histogram_index], 1);
    }
}

int main(int argc, char** argv) {
    const int N = NUM_INPUTS;

    // 在主机上构建直方图输入数据
    int* input = (int*) malloc(N * sizeof(int));

    // 输入数据范围从 0 至 MAX_VALUE - 1 不等
    for (int i = 0; i < N; ++i) {
        input[i] = rand() % MAX_VALUE;
    }

    // 复制到设备
    int* d_input;
    CUDA_CHECK(cudaMalloc((void**) &d_input, N * sizeof(int)));
    CUDA_CHECK(cudaMemcpy(d_input, input, N * sizeof(int), cudaMemcpyHostToDevice));

    // 分配直方图数组
    int* histogram = (int*) malloc(NUM_BUCKETS * sizeof(int));
    memset(histogram, 0, NUM_BUCKETS * sizeof(int));

    int* d_histogram;
    CUDA_CHECK(cudaMalloc((void**) &d_histogram, NUM_BUCKETS * sizeof(int)));
    CUDA_CHECK(cudaMemset(d_histogram, 0, NUM_BUCKETS * sizeof(int)));

    // 执行直方图
    int threads_per_block = 256;
    int blocks = (NUM_INPUTS + threads_per_block - 1) / threads_per_block;

    histogram_kernel<<<blocks, threads_per_block>>>(d_input, d_histogram, N);
    CUDA_CHECK(cudaDeviceSynchronize());

    // 将数据复制回主机,并检查一些值
    CUDA_CHECK(cudaMemcpy(histogram, d_histogram, NUM_BUCKETS * sizeof(int), cudaMemcpyDeviceToHost));

    std::cout << "Histogram counters:" << std::endl << std::endl;
    int num_buckets_to_print = 4;
    for (int i = 0; i < NUM_BUCKETS; i += NUM_BUCKETS / num_buckets_to_print) {
        std::cout << "Bucket [" << i * (MAX_VALUE / NUM_BUCKETS) << ", " << (i + 1) * (MAX_VALUE / NUM_BUCKETS) - 1 << "]: " << histogram[i];
        std::cout << std::endl;
        if (i < NUM_BUCKETS - NUM_BUCKETS / num_buckets_to_print - 1) {
            std::cout << "..." << std::endl;
        }
    }

    free(input);
    free(histogram);
    CUDA_CHECK(cudaFree(d_input));
    CUDA_CHECK(cudaFree(d_histogram));

    return 0;
}

编译运行指令如下:

nvcc -x cu -arch=sm_70 -o histogram histogram.cpp
./histogram

运行结果如下:

Histogram counters:

Bucket [0, 65535]: 4083
...
Bucket [262144, 327679]: 4107
...
Bucket [524288, 589823]: 4015
...
Bucket [786432, 851967]: 4045

复制式方法的 NVSHMEM 实现

       在多个 GPU 上分配工作量的一种方法与我们在 𝜋 估算器上使用的方法相同:给定 𝑁 个整数,我们即可把它们均匀地分配到所有 GPU 上,然后可以对所有 PE 进行归约。我们将此称为**“复制式”方法**,因为在所有 GPU 上都存在完整的直方图副本。我们将第一步,即增加每个直方图桶内的计数值,命名为“列表”步骤;将合并所有 PE 上的结果的第二步命名为“结合”步骤,并分别计算时间(以便与下一个方法进行比较)。

请添加图片描述

       我们将使用归约 API nvshmem_int_sum_reduce() 来归约直方图的所有桶:

nvshmem_int_sum_reduce(team, destination, source, nelems);

如果 destination == source,那么这就变成了就地归约,是 NVSHMEM 中的合理做法;这样做的好处在于,与创建临时目标缓冲区相比,其代码更加干净,所以我们建议在此练习中这样做。

练习

代码如下(file name:histogram_step1.cpp)

#include <iostream>
#include <cstdlib>
#include <chrono>

#include <nvshmem.h>
#include <nvshmemx.h>

inline void CUDA_CHECK (cudaError_t err) {
    if (err != cudaSuccess) {
        fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
        exit(-1);
    }
}

#define NUM_BUCKETS   16
#define MAX_VALUE     1048576
#define NUM_INPUTS    65536

__global__ void histogram_kernel(const int* input, int* histogram, int N)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    if (idx < N) {
        int value = input[idx];

        int histogram_index = ((size_t) value * NUM_BUCKETS) / MAX_VALUE;

	    atomicAdd(&histogram[histogram_index], 1);
    }
}

int main(int argc, char** argv) {
    // 初始化 NVSHMEM
    nvshmem_init();

    // 获取 NVSHMEM 处理元素 ID 和 PE 数量
    int my_pe = nvshmem_my_pe();
    int n_pes = nvshmem_n_pes();

    // 每个 PE(任意)选择与其 ID 对应的 GPU
    int device = my_pe;
    CUDA_CHECK(cudaSetDevice(device));

    // 每台设备处理 1 / n_pes 的部分工作。
    const int N = NUM_INPUTS / n_pes;

    // 在主机上构建直方图输入数据
    int* input = (int*) malloc(N * sizeof(int));

    // 为每个 PE 初始化一个不同的随机数种子。
    srand(my_pe);

    // 输入数据范围从 0 至 MAX_VALUE - 1 不等
    for (int i = 0; i < N; ++i) {
        input[i] = rand() % MAX_VALUE;
    }

    // 复制到设备
    int* d_input;
    d_input = (int*) nvshmem_malloc(N * sizeof(int));
    CUDA_CHECK(cudaMemcpy(d_input, input, N * sizeof(int), cudaMemcpyHostToDevice));

    // 分配直方图数组
    int* histogram = (int*) malloc(NUM_BUCKETS * sizeof(int));
    memset(histogram, 0, NUM_BUCKETS * sizeof(int));

    int* d_histogram;
    d_histogram = (int*) nvshmem_malloc(NUM_BUCKETS * sizeof(int));
    CUDA_CHECK(cudaMemset(d_histogram, 0, NUM_BUCKETS * sizeof(int)));

    // 为合理准确的计时执行一次同步
    nvshmem_barrier_all();

    using namespace std::chrono;

    high_resolution_clock::time_point tabulation_start = high_resolution_clock::now();

    // 执行直方图
    int threads_per_block = 256;
    int blocks = (NUM_INPUTS / n_pes + threads_per_block - 1) / threads_per_block;

    histogram_kernel<<<blocks, threads_per_block>>>(d_input, d_histogram, N);
    CUDA_CHECK(cudaDeviceSynchronize());

    nvshmem_barrier_all();

    high_resolution_clock::time_point tabulation_end = high_resolution_clock::now();

    high_resolution_clock::time_point combination_start = high_resolution_clock::now();

    // 在所有 PE 上执行归约
    nvshmem_int_sum_reduce(NVSHMEM_TEAM_WORLD, d_histogram, d_histogram, NUM_BUCKETS);

    high_resolution_clock::time_point combination_end = high_resolution_clock::now();

    // 打印 PE 0 上的结果
    if (my_pe == 0) {
        duration<double> tabulation_time = duration_cast<duration<double>>(tabulation_end - tabulation_start);
        std::cout << "Tabulation time = " << tabulation_time.count() * 1000 << " ms" << std::endl << std::endl;

        duration<double> combination_time = duration_cast<duration<double>>(combination_end - combination_start);
        std::cout << "Combination time = " << combination_time.count() * 1000 << " ms" << std::endl << std::endl;

        // 将数据复制回主机
        CUDA_CHECK(cudaMemcpy(histogram, d_histogram, NUM_BUCKETS * sizeof(int), cudaMemcpyDeviceToHost));

        std::cout << "Histogram counters:" << std::endl << std::endl;
        int num_buckets_to_print = 4;
        for (int i = 0; i < NUM_BUCKETS; i += NUM_BUCKETS / num_buckets_to_print) {
            std::cout << "Bucket [" << i * (MAX_VALUE / NUM_BUCKETS) << ", " << (i + 1) * (MAX_VALUE / NUM_BUCKETS) - 1 << "]: " << histogram[i];
            std::cout << std::endl;
            if (i < NUM_BUCKETS - NUM_BUCKETS / num_buckets_to_print - 1) {
                std::cout << "..." << std::endl;
            }
        }
    }

    free(input);
    free(histogram);
    nvshmem_free(d_input);
    nvshmem_free(d_histogram);

    // 最终确定 nvshmem
    nvshmem_finalize();

    return 0;
}

编译和运行命令:

nvcc -x cu -arch=sm_70 -rdc=true -I $NVSHMEM_HOME/include -L $NVSHMEM_HOME/lib -lnvshmem -lcuda -o histogram_step1 exercises/histogram_step1.cpp
nvshmrun -np $NUM_DEVICES ./histogram_step1

运行结果如下:

Tabulation time = 0.033777 ms

Combination time = 0.042937 ms

Histogram counters:

Bucket [0, 65535]: 4135
...
Bucket [262144, 327679]: 4028
...
Bucket [524288, 589823]: 4088
...
Bucket [786432, 851967]: 4100


在这里插入图片描述

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

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

相关文章

pandas时间序列,案列

一&#xff1a;pandas时间序列 1.1为什么要学习pandas中的时间序列 不管在什么行业&#xff0c;时间序列都是一种非常重要的数据形式&#xff0c;很多统计数据以及数据的规律也都和时间序列有着非常重要的联系&#xff0c;而且在pandas中处理时间序列是非常简单的 1.2生成一段时…

【Python技巧】:cmd查看Python版本号居然与自己电脑装的版本不一致,特此提出解决方案

项目场景&#xff1a; 大家好&#xff01;欢迎大家看我的博客&#xff0c;最近学习Python的GUI&#xff08;PyQt5&#xff09;的时候发现了自己电脑的一个python问题&#xff0c;我本想装一下PyQt5&#xff0c;顺手查了一下自己电脑的Python版本&#xff0c;没想到居然是Python…

哈希表(二)—— 开散列 / 拉链法 / 哈希桶的模拟实现

哈希表的基本思路是通过某种方式将某个值映射到对应的位置&#xff0c;这里的采取的方式是除留余数法&#xff0c;即将原本的值取模以后再存入到数组的对应下标&#xff0c;即便存入的值是一个字符串&#xff0c;也可以根据字符串哈希算法将字符串转换成对应的ASCII码值&#x…

Week 6 hw3-1 全连接网络反向传播推导

Week 6 hw3-1 全连接网络反向传播推导 折腾了半天&#xff0c;记录一下。 作业中网络由若干全连接层ReLU组成&#xff0c;输出层的函数为softmax&#xff0c;损失函数为交叉熵。 一、记号 设网络有nnn层。如图&#xff0c;当i<ni<ni<n时&#xff0c;我们有如下几条…

机器学习/人工智能 实验一:典型监督学习方法分类实践与比较分析

一、实验目的与要求 (1)利用所学习的监督学习方法完成目标识别实验方案的设计。 (2)编程并利用相关软件完成实验测试&#xff0c;得到实验结果。 (3)通过对实验数据的分析﹑整理&#xff0c;方法的对比&#xff0c;得出实验结论&#xff0c;培养学生创新思维和编写实验报告的能…

【PyTorch深度学习实践】09_卷积神经网络基础

文章目录1.卷积操作1.1 卷积操作1.2 padding-填充1.3 stride-步长1.4 pooling-池化1.5 基础版CNN代码示例1.6 完整CNN代码示例1.卷积操作 卷积神经网络概览 1.1 卷积操作 输入通道数卷积核通道数&#xff0c;卷积核个数输出通道数 1.2 padding-填充 padding是为了让源图像最…

FPGA图像处理HLS实现三种图像缩放算法,线性插值、双线性插值、双三次插值,提供HLS工程和vivado工程源码

目录一、三种图像缩放算法介绍线性插值双线性插值双三次插值二、HLS实现线性插值图像缩放三、HLS实现双线性插值图像缩放四、HLS实现双三次插值图像缩放五、HLS在线仿真并导出IP六、其他FPGA型号HLS在线仿真并导出IP七、zynq7100开发板vivado工程八、上板调试验证九、福利&…

纪念QT可直接安装的离线版最后版本5.14.2

为什么说纪念呢&#xff1f;因为&#xff0c;这个版本之后再也没有可下载下来安装的版本了&#xff0c;因为我们以后再也没有这么方便了。为是很么说纪念呢&#xff1f;因为我们从QT还很柔弱的时候开始就是使用的离线版。 以前用c#来做组态&#xff0c;自定义控件开发起来也还…

基础知识一览2

这里写目录标题1.XML2.1 XML中的转义字符2.2 CDATA区2.3 如何去约束XMl:DTD2.3.1 xml文件内部引用DTD约束2.3.2 xml文件引用外部DTD约束2.3.3 xml文件引用公共DTD约束1.XML xml的文件后缀名是.xmlxml有且只有一个根标签xml的标签是尖括号包裹关键字成对出现的&#xff0c;有开…

如何做好banner设计(banner设计要点包括哪些)

网页设计的Banner作为表达网站价值或者传达广告信息的视觉主体&#xff0c;一直在根据网络环境的变化而变化着&#xff0c;从表现形式到尺寸大小&#xff0c;再到创意的多元化&#xff0c;因此更需要我们网页设计师们对其设计创意进行丰富和完善&#xff0c;才能真正达到宣传的…

Elasticsearch入门——Elasticsearch7.8.0版本和Kibana7.8.0版本的下载、安装(win10环境)

目录一、Elasticsearch7.8.0版本下载、安装1.1、官网下载地址1.2、下载步骤1.3、安装步骤(需要jdk11及以上版本支持)1.4、启动后&#xff0c;控制台中文乱码问题解决二、Node下载、安装&#xff08;安装Kibana之前需要先安装Node&#xff09;2.1、Node官网下载地址2.2、Node下载…

Linux文字处理和文件编辑(三)

1、Linux里的配置文件&#xff1a; /etc/bashrc文件&#xff1a;该配置文件在root用户下&#xff0c;权限很高。~/.bashrc文件&#xff1a;只有当前用户登录时才会执行该配置文件。每次打开终端&#xff0c;都会自动执行配置文件里的代码。比如&#xff0c;alias md‘mkdir’就…

《2022年终总结》

2022年终总结 笔者成为社畜的一年&#xff0c;整整打了一年工&#xff01; 之前都说每年都有点变化&#xff0c;今年的变化可能就是更加懒散了&#xff0c;玩了更多的手机 就是运动的坚持更加多了&#xff0c;收入也增加了&#xff0c;哈哈&#xff01; 其实今年的变化不大&am…

41. 【农产品溯源项目前后端Demo】后端目录结构

本节介绍下后端代码的目录结构。 1. 实现用户管理、菜单管理、角色管理、代码自动生成等服务,归结为系统管理,是若依框架提供的能力。 2. ruoyi-traces实现农产品溯源应用的代码,如果要引入其他Java包,修改本模块的pom.xml文件。 1)config包加载配置文件数据,配置文件路…

FPGA:IIC验证镁光EEPROM仿真模型(纯Verilog)

目录日常唠嗑一、程序设计二、镁光模型仿真验证三、testbench文件四、完整工程下载日常唠嗑 IIC协议这里就不赘述了&#xff0c;网上很多&#xff0c;这里推荐两个&#xff0c;可以看看【接口时序】6、IIC总线的原理与Verilog实现 &#xff0c;还有IIC协议原理以及主机、从机Ve…

基于SpringBoot的车牌识别系统(附项目地址)

yx-image-recognition: 基于spring boot maven opencv 实现的图像深度学习Demo项目&#xff0c;包含车牌识别、人脸识别、证件识别等功能&#xff0c;贯穿样本处理、模型训练、图像处理、对象检测、对象识别等技术点 介绍 spring boot maven 实现的车牌识别及训练系统 基于…

3-1存储系统-存储器概述主存储器

文章目录一.存储器概述&#xff08;一&#xff09;存储器分类1.按在计算机中的作用&#xff08;层次&#xff09;分类2.按存储介质分类3.按存取方式分类4.按信息的可保存性分类&#xff08;二&#xff09;存储器的性能指标二.主存储器&#xff08;一&#xff09;基本组成1.译码…

6 个必知必会高效 Python 编程技巧

编写更好的Python 代码需要遵循Python 社区制定的最佳实践和指南。遵守这些标准可以使您的代码更具可读性、可维护性和效率。 本文将展示一些技巧&#xff0c;帮助您编写更好的 Python 代码 文章目录遵循 PEP 8 风格指南1.遵守 PEP 8 命名约定2. 使用描述性的和有意义的变量名…

读书笔记--- ggplot2:数据分析与图形艺术

最近看了这本书《ggplot2&#xff1a;数据分析与图形艺术》&#xff08;第2版&#xff09;&#xff0c;实际上网页在线版本已经更新到第3版了&#xff08;https://ggplot2-book.org/&#xff09;。 这本书页数不多&#xff0c;但是整体还是值得阅读&#xff0c;不愧是Hadley W…

【Proteus仿真】【STM32单片机】酒精浓度检测系统设计

文章目录一、功能简介二、软件设计三、实验现象联系作者一、功能简介 本项目使用Proteus8仿真STM32单片机控制器&#xff0c;使用LCD1602显示模块、按键模块、LED和蜂鸣器、MQ-3酒精传感器模块等。 主要功能&#xff1a; 系统运行后&#xff0c;LCD1602显示酒精浓度值和阈值&…