40、通用GPU编程：高效内存访问、平铺技术与OpenCL入门

GPU编程：内存优化与OpenCL入门

root9

于 2025-11-02 15:06:37 发布

阅读量20

点赞数

CC 4.0 BY-SA版权

分类专栏：并行编程：从多核到集群文章标签： GPU编程 CUDA OpenCL

本文链接：https://blog.youkuaiyun.com/root9/article/details/154638676

并行编程：从多核到集群专栏收录该内容

48 篇文章 ¥499.90

订阅专栏¥69.90

会员秒杀 ¥9.9 重磅福利

超级会员免费看

通用GPU编程：高效内存访问、平铺技术与OpenCL入门

1. 高效内存访问与平铺技术

在GPU编程中，内核函数的执行通常需要从全局内存中访问大量数据，这是由线程的数据并行执行模型所导致的。然而，全局内存访问的开销较大，因此数据应被复制到访问速度更快的共享内存或寄存器中。CUDA提供了一种名为内存合并（memory coalescing）的技术，用于支持将全局内存中的数据复制到共享内存或寄存器中。

内存合并技术利用了同一warp中的线程在任何时刻都执行相同指令这一特性。当指令为加载操作时，硬件可以检测并行加载操作是否指向全局内存中的相邻内存位置。如果是，硬件会将相邻内存位置的加载操作合并为一次内存访问，这比多次单独的内存访问要快得多。

为了利用内存合并技术实现高效的内存访问，应用程序程序员应在CUDA程序中合理组织数据，使得同一warp中相邻线程标识符的线程访问程序中数组的相邻元素。对于二维数组的访问，具有连续线程标识符的线程应访问行中的相邻元素，因为二维数组按行存储可以使硬件将数据访问合并为一次复制操作。相反，如果同一warp中的线程访问数组列中的相邻元素，则无法实现数据访问的合并，因为相邻列元素存储在内存的不同位置。

为了实现高效的CUDA程序，数据布局应设计为支持内存合并，平铺技术（tiling technique）就是一种常用的编程技术。在平铺技术中，二维数组被分解为大小相同的较小二维数组，称为瓦片（tiles）。使用该数组的算法需要进行修改，以适应瓦片结构，特别是访问二维数据结构的嵌套循环需要修改，以便处理较小的数据集。

以下是一个不使用平铺技术的矩阵乘法CUDA程序示例：

#include <stdio.h>
typedef float * Matrix;
const int N = 32 * 32;
__global__
void MatMulKernel(const Matrix A,const Matrix B,Matrix C){
    float Cval = 0;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    for (int e = 0; e < N; ++e)
        Cval += A[row * N + e] * B[e * N + col];
    C[row * N + col] = Cval;
}
void MatMul (const Matrix A, const Matrix B, Matrix C){
    int size = N * N * sizeof(float);
    Matrix Ad,Bd,Cd;
    cudaMalloc (&Ad, size); cudaMalloc(&Bd, size);
    cudaMemcpy (Ad, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy (Bd, B, size, cudaMemcpyHostToDevice);
    cudaMalloc (&Cd, size);
    dim3 dBlock(32,32);
    dim3 dGrid(N/32,N/32);
    MatMulKernel<<<dGrid,dBlock>>> (Ad,Bd,Cd);
    cudaMemcpy (C,Cd,size, cudaMemcpyDeviceToHost);
}
int main() {
    Matrix A, B, C;
    A = (float *) malloc(N * N * sizeof(float));
    B = (float *) malloc(N * N * sizeof(float));
    C = (float *) malloc(N * N * sizeof(float));
    read_in(A); read_in(B);
    MatMul (A,B,C);
    write_out (C);
    cudaFree (Ad); cudaFree (Bd); cudaFree (Cd);
}

该程序直接在全局内存中访问数据，导致较高的数据访问时间。虽然可以通过大量线程隐藏数据访问时间，但全局内存的带宽可能会限制整体性能。例如，NVIDIA GTX 680的最大内存带宽为192 GB/秒，内核函数 MatMulKernel 中的循环体每次迭代需要两次数据访问和两次算术运算。由于每秒只能加载48 * 10^9个浮点值，因此每秒最多只能执行48 * 10^9次浮点运算，而NVIDIA GTX 680的最大性能为3090 GFLOPS，这意味着仅利用了约1.6%的最大性能。

为了提高CUDA程序的效率，可以先将数据加载到共享内存中，然后从这个更快的内存中访问数据。对于矩阵乘法，由于同一块中的线程会多次使用数据，因此这种方法是有利的。然而，由于共享内存的大小有限，32×N的子数组可能太大而无法放入共享内存。在这种情况下，可以使用平铺技术，将32×N的子数组进一步细分为32×32的瓦片，并将它们依次预加载到共享内存中。

以下是使用平铺技术的矩阵乘法CUDA内核函数示例：

#define TILE_WIDTH 32
__global__
void MatMulTileKernel (Matrix A, Matrix B, Matrix C) {
    __shared__
    float Ads[TILE_WIDTH][TILE_WIDTH];
    __shared__
    float Bds[TILE_WIDTH][TILE_WIDTH];
    int bx = blockIdx.x; int by = blockIdx.y;
    int tx = threadIdx.x; int ty = threadIdx.y;
    int Row = by * TILE_WIDTH + ty;
    int Col = bx * TILE_WIDTH + tx;
    float Cval = 0.0;
    for (int m = 0; m < N/TILE_WIDTH; m++) { /* loop over tiles */
        Ads[ty][tx] = A[Row*N + (m*TILE_WIDTH + tx)];
        Bds[ty][tx] = B[(m*TILE_WIDTH + ty)*N + Col];
        __syncthreads();
        for (int k = 0; k < TILE_WIDTH; k++) /* loop within tile */
            Cval += Ads[ty][k] * Bds[k][tx];
        __syncthreads();
    }
    C[Row * N + Col] = Cval; /* write back to global memory */
}

该内核函数将计算每个元素所需的标量积的循环分解为多个阶段，每个阶段处理A和B的一个瓦片的数据。每个阶段首先将两个瓦片（一个来自每个输入矩阵）加载到共享内存中，然后计算基于这些瓦片数据的部分标量积。在每个阶段结束时，需要进行同步，以确保所有线程都完成了当前瓦片的计算，然后再加载下一个瓦片。

2. OpenCL简介

OpenCL是另一种用于包括GPU在内的计算环境的应用程序编程接口。它是2008年提出的标准化、跨平台编程接口，支持使用由CPU、GPU和其他处理单元组成的计算环境。与CUDA类似，OpenCL基于编程语言C，编程模型也与CUDA非常相似。

OpenCL平台是一个异构平台，由单个主机和一个或多个OpenCL设备（称为计算设备）组成。主机负责与外部环境进行交互，如用户交互或I/O，而设备用于并行计算。一个OpenCL应用程序由一个主机程序和一组用OpenCL-C编程语言实现的内核组成，这些内核由主机程序调用并在设备上执行。

当调用内核时，OpenCL运行时系统会生成一个称为NDRanges的全局索引空间，对于索引空间中的每个点（称为工作项），都会执行一个内核实例。工作项对应于CUDA线程，但不同的是，这些工作项可以通过它们在NDRanges中的全局索引直接寻址。NDRanges是N维索引空间的缩写，目前N可以是1、2或3。

工作项可以分组为工作组，工作组的维度与NDRanges相同。在每个维度上，NDRanges的大小必须能被工作组的数量整除。工作组有一个工作组标识符，工作组中的工作项在其组内还有一个额外的本地标识符，因此一个工作项可以通过其工作组标识符和本地标识符来识别。与CUDA不同的是，工作项有两种识别方式，第一种是通过组标识符和本地标识符的组合，第二种是通过NDRanges中的全局标识符。

OpenCL的并行计算模型是SIMD或SPMD，这意味着所有工作项对不同的数据执行相同的操作。OpenCL程序的数据可以存储在五种不同的内存类型中，与CUDA非常相似。主机内存位于CPU上，只能由主机程序访问。在设备上，有全局内存、常量内存、本地内存和私有内存。全局内存对应于CUDA的全局内存，主机程序可以在GPU的全局内存中动态分配空间，主机和设备程序都可以访问该内存类型。常量内存可以由主机程序读写，由设备程序读取，与CUDA不同的是，主机程序可以在常量内存中动态分配内存空间，并且常量内存的大小不限于64 KB，不同设备的实际大小可能不同。本地内存只能由一个工作组的工作项读写，不能由其他工作组或主机访问，对应于CUDA的共享内存。私有内存仅分配给一个工作项。

OpenCL中的内核对应于CUDA内核函数，但它们使用关键字 __kernel 声明，而不是CUDA中的 __global__ 。以下是一个OpenCL内核 vectoradd() 的示例，用于实现向量a和b的加法，结果存储在向量c中：

__kernel void vectoradd(__global const float *a, __global const float *b, __global float *c) {
    int gid = get_global_id(0);
    c[gid] = a[gid] + b[gid];
}

每个工作项使用全局标识符 get_global_id(0) 寻址，并计算结果向量的一个元素。

执行OpenCL程序的异构计算环境在包含所有设备的上下文中定义，计算任务分配给设备的操作在命令队列中指定。命令队列可以包含内核调用、内存分配操作、复制操作或同步操作，这些操作由设备依次执行。由于命令队列的概念，可以通过指定多个命令队列来实现任务并行性，但需要确保正确的交互。OpenCL还提供了事件的概念，由命令队列中的命令发起。

综上所述，CUDA的内存合并和平铺技术可以提高GPU编程的内存访问效率，而OpenCL作为一种跨平台的编程接口，提供了在异构计算环境中进行并行计算的能力。开发者可以根据具体需求选择合适的编程接口和技术来实现高效的GPU编程。

以下是一个简单的mermaid流程图，展示了使用平铺技术的矩阵乘法的主要步骤：

graph TD;
    A[初始化矩阵A、B、C] --> B[分配GPU内存];
    B --> C[将矩阵A、B复制到GPU];
    C --> D[执行MatMulTileKernel];
    D --> E[将结果矩阵C从GPU复制到主机];
    E --> F[释放GPU内存];

在实际应用中，开发者可以根据具体的硬件环境和计算需求，进一步优化这些技术的使用，以提高程序的性能。例如，在使用OpenCL时，可以根据不同设备的特性调整命令队列和事件的使用，以实现更高效的任务并行性。同时，在使用CUDA的平铺技术时，可以根据共享内存的大小和数据访问模式，合理选择瓦片的大小，以充分利用内存合并的优势。

通用GPU编程：高效内存访问、平铺技术与OpenCL入门

3. 技术对比与应用场景分析

在GPU编程领域，CUDA和OpenCL各有其独特的优势和适用场景。下面我们从多个方面对它们进行对比分析，并探讨不同场景下的技术选择。

3.1 编程复杂度

CUDA ：CUDA的编程相对较为简单，它专门为NVIDIA的GPU设计，提供了一套简洁的API和编程模型。开发者可以专注于GPU的并行计算特性，利用CUDA的线程层次结构（网格、块、线程）来实现高效的并行算法。例如，在矩阵乘法的示例中，CUDA的内核函数可以直接使用线程索引来计算矩阵元素，代码结构清晰易懂。
OpenCL ：OpenCL的编程相对复杂一些，因为它需要考虑异构计算环境中不同设备的特性。开发者需要手动管理设备的选择、内存分配和任务调度等操作。例如，在使用OpenCL时，需要创建上下文、命令队列和事件等对象，以确保程序在不同设备上的正确执行。

3.2 硬件兼容性

CUDA ：CUDA只能在NVIDIA的GPU上运行，这限制了其在其他硬件平台上的应用。但是，由于CUDA是专门为NVIDIA GPU优化的，因此在NVIDIA GPU上可以获得更高的性能。
OpenCL ：OpenCL是跨平台的编程接口，支持多种硬件设备，包括CPU、GPU和其他加速器。这使得开发者可以在不同的硬件平台上运行相同的代码，提高了代码的可移植性。例如，一个OpenCL程序可以在NVIDIA GPU、AMD GPU或Intel CPU上运行，而不需要进行大量的修改。

3.3 内存管理

CUDA ：CUDA的内存管理相对简单，主要包括全局内存、共享内存和寄存器。开发者可以使用 cudaMalloc 和 cudaMemcpy 等函数来管理全局内存的分配和数据传输。共享内存的使用可以提高内存访问效率，但需要注意同步问题。
OpenCL ：OpenCL的内存管理更加灵活，提供了五种不同的内存类型（主机内存、全局内存、常量内存、本地内存和私有内存）。开发者可以根据不同的需求选择合适的内存类型，并使用OpenCL的API来管理内存的分配和访问。例如，常量内存可以用于存储只读数据，本地内存可以用于工作组内的数据共享。

3.4 应用场景

CUDA ：适用于对性能要求较高、且使用NVIDIA GPU的场景。例如，在深度学习、科学计算和图形处理等领域，CUDA可以充分发挥NVIDIA GPU的并行计算能力，实现高效的算法。
OpenCL ：适用于需要跨平台运行、且使用多种硬件设备的场景。例如，在移动设备、嵌入式系统和超级计算机等领域，OpenCL可以提供统一的编程接口，方便开发者在不同的硬件平台上进行并行计算。

以下是一个对比CUDA和OpenCL的表格：
| 特性 | CUDA | OpenCL |
| ---- | ---- | ---- |
| 编程复杂度 | 较低 | 较高 |
| 硬件兼容性 | 仅支持NVIDIA GPU | 支持多种硬件设备 |
| 内存管理 | 相对简单 | 更加灵活 |
| 应用场景 | 高性能计算，NVIDIA GPU | 跨平台计算，多种硬件设备 |

4. 操作步骤总结

为了帮助开发者更好地应用CUDA和OpenCL技术，下面总结了使用这两种技术进行矩阵乘法的操作步骤。

4.1 CUDA矩阵乘法（使用平铺技术）

初始化矩阵 ：在主机上分配内存并初始化输入矩阵A和B，以及结果矩阵C。
分配GPU内存 ：使用 cudaMalloc 函数在GPU的全局内存中分配空间，用于存储矩阵A、B和C。
数据传输 ：使用 cudaMemcpy 函数将矩阵A和B从主机内存复制到GPU的全局内存中。
执行内核函数 ：调用 MatMulTileKernel 内核函数，在GPU上执行矩阵乘法计算。
结果传输 ：使用 cudaMemcpy 函数将结果矩阵C从GPU的全局内存复制到主机内存中。
释放内存 ：使用 cudaFree 函数释放GPU上分配的内存。

以下是一个简单的mermaid流程图，展示了CUDA矩阵乘法的操作步骤：

graph TD;
    A[初始化矩阵A、B、C] --> B[分配GPU内存];
    B --> C[将矩阵A、B复制到GPU];
    C --> D[执行MatMulTileKernel];
    D --> E[将结果矩阵C从GPU复制到主机];
    E --> F[释放GPU内存];

4.2 OpenCL矩阵乘法

创建上下文和命令队列 ：使用OpenCL的API创建上下文和命令队列，用于管理设备和任务调度。
分配内存 ：在GPU的全局内存中分配空间，用于存储矩阵A、B和C。
数据传输 ：将矩阵A和B从主机内存复制到GPU的全局内存中。
编译和执行内核 ：编译OpenCL内核函数，并将其提交到命令队列中执行。
结果传输 ：将结果矩阵C从GPU的全局内存复制到主机内存中。
释放资源 ：释放OpenCL的上下文、命令队列和内存对象。

以下是一个简单的mermaid流程图，展示了OpenCL矩阵乘法的操作步骤：

graph TD;
    A[创建上下文和命令队列] --> B[分配GPU内存];
    B --> C[将矩阵A、B复制到GPU];
    C --> D[编译和执行内核];
    D --> E[将结果矩阵C从GPU复制到主机];
    E --> F[释放资源];

5. 总结与展望

在GPU编程中，高效的内存访问和并行计算是提高程序性能的关键。CUDA的内存合并和平铺技术为我们提供了优化内存访问的有效方法，而OpenCL则为我们提供了跨平台的并行计算解决方案。通过对这两种技术的深入理解和应用，开发者可以根据具体的需求和硬件环境，选择合适的编程接口和技术，实现高效的GPU编程。

未来，随着GPU技术的不断发展和应用场景的不断拓展，GPU编程将面临更多的挑战和机遇。例如，随着深度学习的兴起，对GPU的计算能力和内存带宽提出了更高的要求。开发者需要不断探索新的编程技术和算法，以充分发挥GPU的潜力。同时，随着异构计算环境的普及，OpenCL等跨平台编程接口将变得更加重要，开发者需要掌握多种编程技术，以适应不同的硬件平台和应用需求。

总之，GPU编程是一个充满挑战和机遇的领域，希望本文的介绍能够帮助开发者更好地理解和应用CUDA和OpenCL技术，在实际项目中取得更好的性能和效果。