CUDA学习笔记(6) 共享内存与全局内存

文章浏览阅读1.4w次,点赞16次,收藏45次。共享内存(Sharedmemory)是位于每个流处理器组(SM)中的高速内存空间,主要作用是存放一个线程块(Block)中所有线程都会频繁访问的数据。流处理器(SP)访问它的速度仅比寄存器(R

  共享内存(Shared memory)是位于每个流处理器组(SM)中的高速内存空间,主要作用是存放一个线程块(Block)中所有线程都会频繁访问的数据。流处理器(SP)访问它的速度仅比寄存器(Register)慢,它的速度远比全局显存快。但是他也是相当宝贵的资源,一般只有几十KByte,

  这里以我电脑上的“Quadro K620”为例:

硬件资源 参数
流处理器(SP) 128 * 3 = 384
流处理器组(SM/SMM) 3
全局显存(Global memory) 2048MByte
每个SM中的共享内存(Shared memory) 64KByte
每个SM中的寄存器(Register) 65536个32bit寄存器

  可以看到全局显存(Global memory)的大小比共享显存(Shared memory)大了几个数量级。当然,共享显存的访问速度也一定比全局显存快。再从硬件上看:

这里写图片描述

  每个流处理器组(SM)都拥有属于自己的共享显存(Shared memory)且大小只有64KByte,而这些流处理器组(SM)都位于同一块芯片上(on the same chip),这块芯片又通过PCB电路连接内存芯片(DRAM)。

  • 流处理器(SP)共享显存(Shared memory)上数据的访问属于片上访问,可以立刻取得数据。
  • 流处理器(SP)内存芯片(DRAM)的访问要通过请求内存控制器等一系列操作,然后才能得到数据。

  为了能够明显的看到共享显存(Shared memory)的速度优势,我们就需要把那些被频繁访问的数据直接放到共享显存(Shared memory)中,这样我们的程序在执行时就可以避免频繁访问全局显存(Global memory),导致时间浪费在数据的传输和访问上。

  举个例子,我们需要统计一张超高清图片或者20Mbyte大小的一维数组中每个数值的个数,也就是做直方图统计,结合前面的原子操作(Atomic functions)可以将计算时间至少减半,当然不同架构的GPU会有不同程度的优化。

*注4:直方图统计用于记录一组数据中的数值分布情况,即每个数值出现了多少次,统计数值的分布。

  方案一:

  1. 统计数值分布的数组直接放在全局显存(Global memory)中,每个线程依次读取全局显存中的待统计数据。
  2. 然后使用原子操作对统计数组(Global)中与源数据中数值所对应的元素执行“++”。

  方案二:

  1. 统计数值分布的数组放在全局显存(Global memory)中,每个线程块拥有一个自己的统计数组(Shared)放在共享显存(Shared memory)中。
  2. 线程块中的每个线程依次读取全局显存中的待统计数据,然后使用原子操作对统计数组(Shared)中与源数据中数值所对应的元素执行“++”。
  3. 对线程块中的线程同步,以保证该线程块中的所有线程都已经完成原子操作。
  4. 将当前线程块中的统计数组(Shared)加到统计数组(Global)中。

  这样,方案二就通过使用共享显存(Shared memory)避免了对全局显存(Global memory)的频繁访问,以达到提高程序运行速度的效果。

  方案一核函数的实现:

// This is kernel function file !
#define hist_MAX_VALUE 256

extern "C" __global__ void kernel_func(unsigned int * thread_index_array, unsigned char *source_data_array, \
    unsigned int * histogram_array, unsigned int *clock_counter_array)
{
    // 计算线程号
    unsigned int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    unsigned int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
    // 总线程数
    unsigned int thread_num = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z;

    // 记录线程号
    thread_index_array[thread_index] = thread_index;

    unsigned int counter_1, counter_2;

    counter_1 = clock();

    unsigned int value = source_data_array[thread_index];
    atomicAdd(&histogram_array[value], 1);

    counter_2 = clock();

    clock_counter_array[thread_index] = counter_1;
    clock_counter_array[thread_index + thread_num] = counter_2;
}

这里写图片描述

结果是:方案一统计10MByte大小数据的直方图使用了11.35毫秒

  方案二核函数的实现:

// This is kernel function file !
#define hist_MAX_VALUE 256

extern "C" __global__ void kernel_func(unsigned int * thread_index_array, unsigned char *source_data_array, \
    unsigned int * histogram_array, unsigned int *clock_counter_array)
{
    // 计算线程号
    unsigned int block_index = blockIdx.x + blockIdx.y * gridDim.x

版权声明:本文内容由互联网用户自发贡献,该文观点仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 举报,一经查实,本站将立刻删除。
如需转载请保留出处:https://bianchenghao.cn/88368.html

(0)
编程小号编程小号

相关推荐

发表回复

您的电子邮箱地址不会被公开。 必填项已用 * 标注