04、cuda reduce了解内存使用

最新推荐文章于 2024-06-03 10:18:50 发布
最新推荐文章于 2024-06-03 10:18:50 发布
阅读量402 收藏 9
点赞数 6
分类专栏： Cuda 文章标签： cuda gpu
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_45063703/article/details/135679284
版权
本文介绍了从CPU计算到GPU计算的性能提升，详细讨论了全局内存、共享内存、动态共享内存以及原子函数在数组规约中的应用，旨在展示如何通过这些方法加速计算并减少内存使用。最后提到未来可能的研究方向——在线程束层面的规约计算。
摘要由CSDN通过智能技术生成

文章目录

前言

从数组规约中了解内存使用情况

1、CPU操作

real reduce(const real *x,const int N){
    real sum = 0.0;
    for(int n=0;n<N;++n){
        sum += x[n];
    }
    return sum;
}

2、全局内存

void __global__ reduce_global(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    real *x = d_x + blockDim.x * blockIdx.x;

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {
        if (tid < offset)
        {
            x[tid] += x[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[blockIdx.x] = x[0];
    }
}

3、共享内存

共享内存可以被程序员直接操控

减少对全局内存的访问，实现高效的线程块内部通信

也可以提高全局内存访问的合并度。

void __global__ reduce_shared(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    __shared__ real s_y[128];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {

        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[bid] = s_y[0];
    }
}

4、动态共享内存

void __global__ reduce_dynamic(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {

        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[bid] = s_y[0];
    }
}

5、完整代码

#include "error.cuh"
#include <stdio.h>

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

const int NUM_REPEATS = 100;
const int N = 100000000;
const int M = sizeof(real) * N;
const int BLOCK_SIZE = 128;

void timing(real *h_x, real *d_x, const int method);

int main(void)
{
    real *h_x = (real *) malloc(M);
    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1.23;
    }
    real *d_x;
    CHECK(cudaMalloc(&d_x, M));

    printf("\nUsing global memory only:\n");
    timing(h_x, d_x, 0);
    printf("\nUsing static shared memory:\n");
    timing(h_x, d_x, 1);
    printf("\nUsing dynamic shared memory:\n");
    timing(h_x, d_x, 2);

    free(h_x);
    CHECK(cudaFree(d_x));
    return 0;
}

void __global__ reduce_global(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    real *x = d_x + blockDim.x * blockIdx.x;

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {
        if (tid < offset)
        {
            x[tid] += x[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[blockIdx.x] = x[0];
    }
}

void __global__ reduce_shared(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    __shared__ real s_y[128];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {

        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[bid] = s_y[0];
    }
}

void __global__ reduce_dynamic(real *d_x, real *d_y)
{
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n < N) ? d_x[n] : 0.0;
    __syncthreads();

    for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1)
    {

        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncthreads();
    }

    if (tid == 0)
    {
        d_y[bid] = s_y[0];
    }
}

real reduce(real *d_x, const int method)
{
    int grid_size = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
    const int ymem = sizeof(real) * grid_size;
    const int smem = sizeof(real) * BLOCK_SIZE;
    real *d_y;
    CHECK(cudaMalloc(&d_y, ymem));
    real *h_y = (real *) malloc(ymem);

    switch (method)
    {
        case 0:
            reduce_global<<<grid_size, BLOCK_SIZE>>>(d_x, d_y);
            break;
        case 1:
            reduce_shared<<<grid_size, BLOCK_SIZE>>>(d_x, d_y);
            break;
        case 2:
            reduce_dynamic<<<grid_size, BLOCK_SIZE, smem>>>(d_x, d_y);
            break;
        default:
            printf("Error: wrong method\n");
            exit(1);
            break;
    }

    CHECK(cudaMemcpy(h_y, d_y, ymem, cudaMemcpyDeviceToHost));

    real result = 0.0;
    for (int n = 0; n < grid_size; ++n)
    {
        result += h_y[n];
    }

    free(h_y);
    CHECK(cudaFree(d_y));
    return result;
}

void timing(real *h_x, real *d_x, const int method)
{
    real sum = 0;

    for (int repeat = 0; repeat < NUM_REPEATS; ++repeat)
    {
        CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));

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

        sum = reduce(d_x, method);

        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));
    }

    printf("sum = %f.\n", sum);
}

6、原子函数

void __global__ reduce(const real *d_x,real *d_y,const  int  N ){
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n<N) ? d_x[n]:0.0;
    __syncthreads();

    for(int offset = blockDim.x >>1;offset >0;offset >>=1){
        if (tid < offset){
            s_y[tid] += s_y[tid + offset]; 
        }
        __syncthreads();
    }

    if(tid==0){
        atomicAdd(d_y,s_y[0]);
    }
}

完整代码

#include "error.cuh"
#include<stdio.h>

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

const int NUM_REPEAT = 100;
const int N = 10000000;
const int M = sizeof(real) * N ;
const int BLOCK_SIZE = 128;

void timing(const real *d_x);

int main(void){
    real *h_x = (real *) malloc(M);
    for(int n=0;n<N; ++n){
        h_x[n] = 1.23;
        }
    real *d_x;
    CHECK(cudaMalloc(&d_x,M));
    CHECK(cudaMemcpy(d_x,h_x,M,cudaMemcpyHostToDevice));

    printf("\nusing atomicAdd:\n");
    timing(d_x);

    free(h_x);
    CHECK(cudaFree(d_x));
    return 0;

}

void __global__ reduce(const real *d_x,real *d_y,const  int  N ){
    const int tid = threadIdx.x;
    const int bid = blockIdx.x;
    const int n = bid * blockDim.x + tid;
    extern __shared__ real s_y[];
    s_y[tid] = (n<N) ? d_x[n]:0.0;
    __syncthreads();

    for(int offset = blockDim.x >>1;offset >0;offset >>=1){
        if (tid < offset){
            s_y[tid] += s_y[tid + offset]; 
        }
        __syncthreads();
    }

    if(tid==0){
        atomicAdd(d_y,s_y[0]);
    }
}

real reduce(const real *d_x){
    const int gird_size = (N+BLOCK_SIZE -1) / BLOCK_SIZE;
    const int smem = sizeof(real) * BLOCK_SIZE;

    real h_y[1] = {0};
    real *d_y;
    CHECK(cudaMalloc(&d_y,sizeof(real)));
    CHECK(cudaMemcpy(d_y,h_y,sizeof(real),cudaMemcpyHostToDevice));

    reduce<<<gird_size,BLOCK_SIZE,smem>>>(d_x,d_y,N);

    CHECK(cudaMemcpy(h_y,d_y,sizeof(real),cudaMemcpyDeviceToHost));
    CHECK(cudaFree(d_y));
    return h_y[0];
}

void timing(const real *d_x){
    real sum = 0;
    for(int repeat=0;repeat < NUM_REPEAT;++repeat){
        cudaEvent_t start,stop;
        CHECK(cudaEventCreate(&start));
        CHECK(cudaEventCreate(&stop));
        CHECK(cudaEventRecord(start));
        cudaEventQuery(start);

        sum = reduce(d_x);

        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));

    }
    printf("sum = %f.\n",sum);

}

7、总结

1、从cpu计算到GPU计算，加快计算速度

2、从全局内存、共享内存、动态共享内存、原子函数，不断加速计算和减少内存使用

3、数组规约的加速示例

4、进一步研究，还可以从线程束方面进行规约计算。

nsq_ai
关注
专栏目录
CUDA 基础算法之reduce、scan、histogram
lei_h11的博客
12-30 2946
前言 之前对于CUDA的学习基本上就是不会就查，拿来就用的状况，对一些基础算法的了解不是特别深，之前在面试的时候还有被问到scan扫描算法计算数组的前缀和，表示还没有详细的了解以致只能尴尬地说不清楚，是真的贼尴尬啊，后来去学了些视频课，才逐渐有了一些些基础。（说起来之前还不知道有step complexity 和 work complexity这两种复杂度呢） 1、step and work co...
CUDA 11功能清单
吴建明wujianming_110117
04-10 1229
CUDA 11功能清单 基于NVIDIA Ampere GPU架构的新型NVIDIA A100 GPU在加速计算方面实现了最大的飞跃。A100 GPU具有革命性的硬件功能，CUDA 11与A100一起发布。 CUDA 11能够利用新的硬件功能来加速HPC，基因组学，5G，渲染，深度学习，数据分析，数据科学，机器人技术以及更多不同的工作负载。 CUDA 11包含了所有功能-从平台系统软件到入门和开发GPU加速的应用程序所需的一切。本文概述了此版本中的主要软件功能： • 支持NVIDIA Ampere GPU
Reduce内存不足的解决方案
劫无忧
01-27 1万+
MapReduce作业运行时，任务可能会失败，报out of memory错误。这个时候可以采用以下几个过程调优简单粗暴： 加大内存哪个阶段报错就增加那个阶段的内存。以reduce阶段为例，map阶段的类似mapreduce.reduce.memory.mb=5120 //设置reduce container的内存大小 mapreduce.reduce.java.opts=-Xms2000m -
【hadoop运维】running beyond physical memory limits：正确配置yarn中的mapreduce内存
hiliang521的博客
09-03 1695
【hadoop运维】running beyond physical memory limits：正确配置yarn中的mapreduce内存
cuda reduce学习
09-12 1285
这个程序实现的是加法的并行运算：#include <stdio.h> #include <time.h> #include <stdlib.h> #include <cuda_runtime.h> __global__ void parallel_reduce_kernel(float* d_out, float* d_in){ int myID = threadIdx.x + blockI
CUDA reduce 并行规约求和
MAX的专栏
12-28 1176
例如, 1024 * 1024 大小的数据, 每一个block处理1024个数据, 一共1024 个block, 处理完后剩下1024个数据, 接着在处理, 就有加速的效果:每一个block 块作为。
CUDA编程练习（九） 归约
SimpleBright_Man的博客
06-03 470
拖延症晚期，但是还有得救
CUDA高性能编程之驱动程序编写之MPI
12-26
例如，使用MPI进行负载均衡，或者通过Allreduce操作进行全局结果的聚合。 5. 实战案例：cudaMPI中的示例代码可以帮助初学者理解CUDA和MPI如何协作。这些示例可能涵盖简单的点对点通信，如传递矩阵的一部分到另一个...
cuda by example 例程+教材+《GPU高性能计算之CUDA》实例
05-13
1. **CUDA架构**：了解CUDA的多线程执行模型，包括线程块（Thread Block）、网格（Grid）、流式多处理器（SM, Streaming Multiprocessor）以及寄存器、全局内存、共享内存和常量内存等内存层次。 2. **CUDA编程模型...
CUDA(六). 从并行排序方法理解并行化思维——冒泡、归并、双调排序的GPU实现
热门推荐
Rachel Zhang的专栏
09-14 4万+
在第五讲中我们学习了GPU三个重要的基础并行算法： Reduce, Scan 和 Histogram，分析了 其作用与串并行实现方法。 在第六讲中，本文以冒泡排序 Bubble Sort、归并排序 Merge Sort 和排序网络中的双调排序 Bitonic Sort 为例， 讲解如何从数据结构课上学的串行并行排序方法转换到并行排序，并附GPU实现代码。 在并行方法中，我们将考虑到并行方法需要注意
mapreduce内存分配
AI蜗牛之家
07-05 1万+
稍微有点mapreduce使用经验的同学肯定对OOM不陌生，对的，我目前在mapReduce里面遇到的最多的报错也是内存分配出错，所以看到好多hadoop执行脚本里面有好多关于内存的参数，虽然是知道和内存分配有关系，但是我依然不太清楚具体的原理，从网上查阅相关资料，看到博主整理的，甚是欣慰，稍作整理如下： 关于mapreduce程序运行在yarn上时内存的分配一直是一个让我蒙圈的事情，单独查任何...
YARN和MapReduce内存分配计算公式
TomAndersen的博客
06-08 2708
前言 hadoop：2.7.7 本文内容均来自：HDP Command Line Installation 2.6.5 中第1.10节 HDP（Hortonworks Data Platform）是最常见的Hadoop的第三方发行版之一，类似的Hadoop发行版还有CDH、MapR等 计算YARN和MapReduce的内存需求 在Hadoop集群中，YARN管理着集群中的每个节点上的可用资源，并为运行在集群中的应用程序（如MapReduce）提供需要的资源，其中Container是YARN中最小的资源
MapReduce的reduce内存详解
王傲旗的大数据之路
11-19 5286
前提:map阶段已经合并成大文件存入到磁盘中 1.Reduce发送拷贝请求到Map端 2.Reduce加载环形缓冲区中 3.达到一定的阈值(默认0.8)写出一个小文件 4.小文件达到一定的个数进行合并 5.将最终的结果传送给reduce输出到HDFS ...
CUDA系列学习（五）GPU基础算法: Reduce, Scan, Histogram
Rachel Zhang的专栏
06-25 4万+
喵~不知不觉到了CUDA系列学习第五讲，前几讲中我们主要介绍了基础GPU中的软硬件结构，内存管理，task类型等；这一讲中我们将介绍3个基础的GPU算法：reduce，scan，histogram，它们在并行算法中非常常用，我们在本文中分别就其功能用处，串行与并行实现进行阐述。 1. Task complexitytask complexity包括step complexity（可以并行成几个操作
CUDA编程——Mars：MapReduce on GPU
Zhang Junior 的博客
01-19 5061
CUDA编程——Mars：MapReduce on GPUGPU加速机器学习　　GPU是一种SIMT（单指令多线程）体系结构，即多个线程执行同一个指令，而每个线程操作的数据不同。这种结构令GPU天生具有大规模计算能力。GPU出色的浮点计算性能特别提高了深度学习两大关键活动：分类和卷积的性能，同时又达到所需的精准度。深度学习需要很高的内在并行度、大量的浮点计算能力以及矩阵预算，而GPU可以提供这些能力
MapReduce 在Shuffle阶段 内存溢出原因分析及处理方法
houzhizhen的专栏
12-03 8199
现象 在Reduce运行中，有时出现内存溢出错误，抛出的异常信息如下： ··· Error: org.apache.hadoop.mapreduce.task.reduce.Shuffle$ShuffleError: error in shuffle in fetcher#1 at org.apache.hadoop.mapreduce.task.reduce.Shuffle.run(Shuffl...
[YARN] Yarn下Mapreduce的内存参数理解
一只刚刚上路的猿
11-08 1万+
这篇文章算是给自己重新缕清MR下内存参数的含义 Container是什么？ Container就是一个yarn的java进程，在Mapreduce中的AM，MapTask，ReduceTask都作为Container在Yarn的框架上执行，你可以在RM的网页上看到Container的状态 基础 Yarn的ResourceManger（简称RM）通过逻辑上的队列分配内存，CP
关于组织参加“第八届‘泰迪杯’数据挖掘挑战赛”的通知-4页
最新发布
11-03
关于组织参加“第八届‘泰迪杯’数据挖掘挑战赛”的通知-4页
PyMySQL-1.1.0rc1.tar.gz
11-03
PyMySQL-1.1.0rc1.tar.gz
cuda锁页内存使用方法
10-25
在 CUDA 中，为了有效地管理内存，特别是对于那些需要在主机和设备之间频繁传输的数据，锁页内存（Host-Registered Memory）是一个有用的功能。 **使用方法**: 1. **分配**: 使用 `cudaMallocManaged` 函数动态地...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

nsq_ai

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值