【CUDA】五、基础概念：Coalescing合并用于内存优化

最新推荐文章于 2024-05-02 21:20:50 发布

prinTao

最新推荐文章于 2024-05-02 21:20:50 发布

阅读量813

点赞数 12

分类专栏： CUDA编程文章标签：算法

本文链接：https://blog.csdn.net/prinTao/article/details/135620784

版权

CUDA编程专栏收录该内容

11 篇文章 2 订阅

订阅专栏

先来看之前的例子：

矩阵乘法中的 Coalescing writes

刚才的代码中，可以观察到两个for循环，这里可以进行优化。

“coalescing writes”（合并写操作）

“coalescing writes”（合并写操作）是一种优化内存访问模式的技术，它能显著提高内存带宽的利用效率。这种技术尤其对于全局内存访问非常重要，因为全局内存访问速度相比于核心计算速度要慢得多。

底层原理

内存事务：当GPU的线程尝试访问全局内存时，这些访问被分组为内存事务。每个事务可以一次性读取或写入多个连续的字节。使用适当大小的数据类型以匹配内存事务的大小。
内存对齐：为了有效地合并写操作，线程访问的内存地址应该是对齐的，并且连续线程访问的地址也应该是连续的。确保数据结构和数组在内存中对齐。
线程访问模式：如果一个线程块中的所有线程都按照一定的模式（例如，线程i访问地址i）访问连续的内存地址，则这些访问可以被合并成一个或几个内存事务。设计线程块和线程索引以便线程以线性和连续的顺序访问内存。减少线程内的条件分支，以保持连续的内存访问模式。

代码

__global__ void MatrixMultiplyCoalesced(float *A, float *B, float *C, int N) {
    // 计算行和列索引
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;

    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; k++) {
            // 累加计算矩阵C中(row, col)位置的值
            sum += A[row * N + k] * B[k * N + col];
        }
        // 写入计算结果到矩阵C中，利用合并写操作优化
        // 每个线程按照顺序写入连续的内存地址
        C[row * N + col] = sum;
    }
}

优化点：

合并写操作：在写入结果到矩阵C时，每个线程写入的是连续的内存位置（C[row * N + col]）。这样，当多个线程同时写入时，由于它们访问的是连续的内存地址，这些写操作可以被合并成较少的内存事务。这种访问模式对于全局内存来说是高效的。
线程索引的布局：通过合理的线程索引布局（即row和col的计算方式），我们确保了线程以线性和有序的方式访问全局内存，这对于实现高效的合并写操作至关重要。

深入分析

先来一点基础概念

Row-Major Order（行主序存储）

Row-Major Order（行主序）是一种在计算机内存中存储多维数组数据的方法。在行主序排列中，多维数组的行元素是连续存储的。这意味着二维数组的第一行的所有元素在内存中紧接着排列，其次是第二行的所有元素，依此类推。这种存储方式非常普遍，尤其是在C和C++等编程语言中。

连续性：在行主序排列中，数组的每一行元素在内存中是连续存储的。这意味着如果你有一个二维数组，数组中的第一行的所有元素在内存中是连续的。
内存地址计算：在行主序排列中，二维数组元素**array[i][j]**的内存地址可以通过下面的公式计算：
```
Address = BaseAddress + (i * NumberOfColumns + j) * ElementSize
```
其中，**BaseAddress是数组的起始地址，NumberOfColumns是数组的列数，ElementSize**是数组每个元素的大小。
优势：行主序排列的一个主要优势是它可以提高访问数组行时的内存访问效率。当按行顺序遍历数组时，由于内存预取和缓存的效果，可以减少缓存未命中的情况。

由于CUDA的内存访问模式，合理地组织数据可以显著提高性能。例如，在处理二维数据结构时，应该尽量确保线程以行主序访问内存，这样可以利用合并内存访问的优势，减少内存延迟。

Aligned Accesses（对齐访问）

Aligned Accesses（对齐访问）：内存访问操作（如读取或写入）是针对特定边界上的地址进行的，这些边界通常是数据类型大小的整数倍。

对齐访问的基本原理

对齐的定义：如果一个地址是N字节数据类型大小（例如，4字节整型）的整数倍，则称该地址为“对齐”的。例如，对于4字节整型，地址0, 4, 8, 12等是对齐的。
访问效率：对齐的内存访问比非对齐的内存访问更高效。非对齐访问可能需要额外的内存周期来处理，因为它可能跨越两个或更多的内存字。

在CUDA中的应用

在CUDA编程中，对齐访问对于优化性能尤其重要：

合并内存访问：在CUDA中，对齐的全局内存访问可以在编译器中被合并为更少的内存事务，从而提高内存带宽的利用率。
数据结构设计：在设计数据结构时，应当考虑内存对齐，以确保高效的内存访问。例如，使用**__align__(n)**指令来确保CUDA结构体或数组是按照特定边界对齐的。

对齐方式

CUDA和C++中都有强制对齐的手段。

CUDA中的数据对齐方式

方法/关键字	描述	例子
align(n)	显式指定变量或结构体的对齐方式为n字节	align(16) float2 a;
cudaMallocPitch	分配二维数组，确保每行都是对齐的	cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);

在CUDA中，__align__(n) 关键字是特定于NVIDIA的编译器扩展，用于指定变量或结构体在内存中的对齐方式。cudaMallocPitch 函数用于分配二维数组，同时确保每行数据在内存中是对齐的，这对于优化二维数据的内存访问效率非常关键。

C++中的数据对齐方式

方法/关键字	描述	例子
alignas	C++11标准中引入，用于指定变量或类型的对齐要求	alignas(16) float a;
#pragma pack(n)	指定结构体、联合体和类成员的字节对齐方式	#pragma pack(push, 1) struct MyStruct { char a; int b; }; #pragma pack(pop)
attribute((aligned(n)))	GCC特有的属性，用于指定变量或类型的对齐要求	int a attribute((aligned(16)));
std::aligned_storage	C++标准库模板类型，用于创建具有特定对齐要求的存储空间	using AlignedType = std::aligned_storage<sizeof(MyType), alignof(MyType)>::type;
std::aligned_union	C++标准库模板类型，创建可容纳其成员中任何一个的对齐存储空间	using AlignedUnion = std::aligned_union<0, Type1, Type2>::type;
alignof	C++11引入的操作符，返回其参数类型的对齐要求	size_t alignment = alignof(double);

在C++中，alignas 是C++11标准中引入的关键字，用于指定变量或类型的对齐要求。#pragma pack 和 __attribute__((aligned(n))) 是编译器特定的指令，分别用于Visual Studio和GCC。std::aligned_storage 和 std::aligned_union 是C++标准库中的模板类型，用于创建具有特定对齐要求的存储空间。alignof 是用于查询类型对齐要求的操作符。