🍪
cookielau
  • Introduction
  • Machine Learning
    • Distributed
      • Bookmarks
    • NLP
      • Transformers
    • MLC
      • Tensor Program Abstraction
      • End-to-End Module Execution
  • Framework
    • PyTorch
      • Bookmarks
      • Model
      • Shared
      • Miscellaneous
    • Tensorflow
      • Bookmarks
      • Model
      • Shared
      • Miscellaneous
    • CUDA
      • Bookmarks
    • DeepSpeed
    • Bagua
      • Model
      • Optimizer
    • Others
      • Bookmarks
  • About Me
    • 2022-04-28
  • Random Thoughts
  • Archives
    • CPP
      • Bookmarks
      • Container
      • Algorithm
      • FILE CONTROL
      • Virtual Table
      • Assembly
      • Key Words
      • Problems
      • Others
    • JAVA
      • String Container
      • Maps
    • PYTHON
      • Bookmarks
      • Python Tools
        • Batch Rename
        • Combine Excel
        • Excel Oprations
        • Read Write Excel
        • Rotate PDF
      • Library
        • Pandas Notes
        • Numpy Notes
        • Json Notes
      • Spider
        • Selenium Install
        • Selenium Locating
        • Selenium Errors
        • Selenium Basics
      • Django
        • Start Up
      • Others
    • LINUX
      • Installation
      • Cli Tools
      • WSL
      • Bugs
    • JUNIOR2
      • Economics
        • Chapter 0x01 经济管理概述
        • Chapter 0x02 微观市场机制分析
        • Chapter 0x03 生产决策与市场结构
        • Chapter 0x04 宏观经济市场分析
        • Chapter 0x05 管理的职能
        • Chapter 0x06 生产系统结构与战略
        • Chapter 0x0b 投资项目经济评价
        • Chapter 0x0f 投资项目经济评价
      • Computer Network
        • 概述
        • 分层模型
        • 物理层
        • 数据链路层
        • 网络层
        • 传输层
        • 应用层
        • HTTP(s)实验
        • [Practice]
      • Software Engineering
        • Introduction
        • Demand Analysis
        • Task Estimation
        • Presentation
      • Network Security
        • Chapter 0x01 概述
        • Chapter 0x02 密码学
        • Chapter 0x03 公钥体制
        • Chapter 0x04 消息认证
        • Chapter 0x05 密钥管理
        • Chapter 0x06 访问控制
        • Assignments
      • x86 Programming
        • Basic Knowledge
        • Program Design
        • System Interruption
        • Frequently used functions
    • MD&LaTex
      • Markdown
      • LaTex
    • NPM
      • NPM LINK
    • MyBlogs
      • 2020BUAA软工——“停下来,回头看”
      • 2020BUAA软工——“初窥构建之法”
      • 2020BUAA软工——“上手软件工程,PSP初体验!”
      • 2020BUAA软工——“深度评测官”
      • 2020BUAA软工——“并肩作战,平面交点Pro”
    • SC
      • PAC 2022
        • Lectures
      • OpenMP & MPI
        • MPI Overview
        • Message Passing Programming
        • OpenMP Overview
        • Work Sharing Directives
        • Annual Challenge
        • Future Topics in OpenMP
        • Tasks
        • OpenMP & MPI
    • Hardware
      • Nvidia GPU
        • Frequent Error
        • Memory Classification
        • CUDA_7_Streams_Simplify_Concurrency
        • Optimize_Data_Transfers_in_CUDA
        • Overlap_Data_Transfers_in_CUDA
        • Write_Flexible_Kernels_with_Grid-Stride_Loops
        • How_to_Access_Global_Memory_Efficiently
        • Using_Shared_Memory
      • Intel CPU
        • Construction
        • Optimization
        • Compilation
        • OpenMP
    • English
      • Vocab
      • Composition
    • Interview
      • Computer Network
Powered by GitBook
On this page
  • testCode
  • Misaligned Data Access
  • Stride Memory Access

Was this helpful?

  1. Archives
  2. Hardware
  3. Nvidia GPU

How_to_Access_Global_Memory_Efficiently

PreviousWrite_Flexible_Kernels_with_Grid-Stride_LoopsNextUsing_Shared_Memory

Last updated 5 years ago

Was this helpful?

How to Access Global Memory Efficiently in CUDA C/C++ Kernels

2020 Jan 28th CookieLau

Source:

[toc]

testCode

分别用一下代码测试 offset 访问和 stride 访问的间隔对带宽的影响:

__global__ void offset(T* a, int s)
{
  int i = blockDim.x * blockIdx.x + threadIdx.x + s;
  a[i] = a[i] + 1;
}

__global__ void stride(T* a, int s)
{
  int i = (blockDim.x * blockIdx.x + threadIdx.x) * s;
  a[i] = a[i] + 1;
}

Misaligned Data Access

C870: Compute Capability 1.0 C1060: Compute Capability 1.3 C2050: Compute Capability 2.0

在 device 中分配的数组都被 cuda Driver 按照 256字节 对齐,当访存 global Memory 的时候可以通过 32字节、64字节 或 128字节等分块进行数据交换。

对于C870这种 compute capability 在 1.0 及以下的 GPU,其 warp size 只有16,而且当发生 misaligned access 的时候,会对每个 misaligned 的 data 单独进行存取,所以会从原来的 16 thread 变成 16次的 32-bytes 访存。对于存取 float 数据而言,每次取 32-bytes 的数据中只有 4-bytes 是有效的,所以带宽减少到原来(offset=0) 的 1/8.

对于C1060这种 compute capability 较好的,对带宽的影响没有 C870 那么严重,只要 misaligned 落在访存的 segment,如 32,64,128-bytes 只会降低部分的性能。

对于C2050这种 compute capability 大于等于 2.0 的设备而言,其每个 multiprocessor 都配有一个 128-bytes 的 L1 cache,所以offset的改变几乎不会对带宽产生影响。

Stride Memory Access

对于C870完全不能处理misaligned,只能处理linear+aligned的架构来说,除了stride=1的情况,都出现了 7/8 的落差,变为正常情况下的 1/8。

对于CC(Compute Capability>1.0) 的来说,都能处理部分 misaligned,所以带宽曲线是 smoothly 下降,但对于相隔很远的访问也无能为力。

但是我们又时常需要进行 stride 访问,比如 grid-stride,则可以通过 shared memory 来解决这一问题。 Shared memory 是 on-chip 的、被一个 thread block 中的所有 threads 共享的一部分内存。

举一个例子:将 2D 的数组的数据加载到 shared memory 里面进行访存降低对带宽的伤害。在共享内存中是没有 stride access penalty 的。

https://devblogs.nvidia.com/how-access-global-memory-efficiently-cuda-c-kernels/