From 70ec6875c56f72f5ed4346a10fa76cf94259c157 Mon Sep 17 00:00:00 2001 From: unknown Date: Sun, 21 Nov 2021 17:32:28 +0800 Subject: [PATCH] add capter7 --- ReadMe.md | 3 + capter7/ReadMe.md | 56 ++++++++++++++ capter7/error.cuh | 26 +++++++ capter7/global.cu | 193 ++++++++++++++++++++++++++++++++++++++++++++++ capter7/matrix.cu | 134 ++++++++++++++++++++++++++++++++ 5 files changed, 412 insertions(+) create mode 100644 capter7/ReadMe.md create mode 100644 capter7/error.cuh create mode 100644 capter7/global.cu create mode 100644 capter7/matrix.cu diff --git a/ReadMe.md b/ReadMe.md index f9487f8..ac6b38a 100644 --- a/ReadMe.md +++ b/ReadMe.md @@ -9,6 +9,9 @@ CUDA gpu 编程学习，基于 《CUDA 编程——基础与实践》（樊哲 3. [简单 CUDA 程序的基本框架](./capter3/ReadMe.md) 4. [CUDA 程序的错误检测](./capter4/ReadMe.md) 5. [GPU 加速的关键](./capter5/ReadMe.md) +6. [CUDA 内存组织](./capter6/ReadMe.md) +7. [全局内存的合理使用](./capter7/ReadMe.md) + CUDA 官方文档： [CUDA c++编程指南](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html) diff --git a/capter7/ReadMe.md b/capter7/ReadMe.md new file mode 100644 index 0000000..c8e474a --- /dev/null +++ b/capter7/ReadMe.md @@ -0,0 +1,56 @@ +# 全局内存的合理使用 + +在各种设备内存中，全局内存具有最低的访问速度，往往是一个 CUDA 程序的性能瓶颈。 + +------ + +## 全局内存的合并与非合并访问 + +对全局内存的访问将触发内存事务，即数据传输。 +在启用了 L1 缓存的情况下，对全局内存的读取将首先尝试经过 L1 缓存；如果未命中， +则尝试经过 L2 缓存；如果再次未命中，则直接从 DRAM 读取。 + +一次 **数据传输处理** 的数据量在默认情况下是 32 字节。 +一次数据传输中，从全局内存转移到 L2 缓存的一片内存的首地址一定是 32 的整数倍。 +也就是说，一次数据传输只能从全局内存读取地址为 0-31 字节、32-63 字节等片段的数据。 + +**合并度**，即线程束请求的字节数与由此导致的所有内存事务中所传输的字节数之比。 +如果所有数据传输处理的数据都是线程束所需要的，则合并度为 100%，即 **合并访问**； +否则，即为 **非合并访问**。 + +以仅使用 L2 缓存的情况为例，一次数据传输指的就是将 32 字节数据从全局内存（DRAM） +通过 32 字节的 L2 缓存片段（cache sector）传输到 SM。 +考虑一个线程束访问单精度浮点数类型的全局内存变量的场景， +一个单精度浮点数占有 4 个字节，故一次访问需要 32*4 个字节的数据。在理想情况下， +即合并度为 100% 时，将仅触发 128/32=4 次调用 L2 缓存的数据传输。 +如果线程束请求的全局内存地址刚好为 0-127 字节或 128-255 字节，就能与 4 次数据 +传输所处理的数据完全吻合，这种情况下就是合并访问。 + +64 位系统中基本数据类型的内存长度（字节）： + + int size: 4 + short size: 2 + float size: 4 + double size: 8 + char size: 1 + bool size: 1 + long size: 4 + int pointer size: 8 + float pointer size: 8 + double pointer size: 8 + char pointer size: 8 + +------ + +## 矩阵转置 + +在核函数中，如果读取操作是非合并访问，则可以采用 *只读数据缓存技术*，通过加载函数 +`__ldg()` 读取全局内存，从而对数据的读取进行缓存、缓解非合并访问的影响。 + +从帕斯卡架构开始，编译器会自动判断并调用 `__ldg()` 函数提升性能；对于开普勒架构、 +麦克斯韦架构，默认情况下不会使用 `__ldg()` 函数，需要手动配置。 + +对于核函数中全局内存的写入，则没有类似函数可用。所以若不能满足读取和写入都是合并的， +一般应该尽量做到合并写入。 + +------ diff --git a/capter7/error.cuh b/capter7/error.cuh new file mode 100644 index 0000000..e9d4cef --- /dev/null +++ b/capter7/error.cuh @@ -0,0 +1,26 @@ +#pragma once +#include +#include +#include +#include + +#define CHECK(call) \ +do { \ + const cudaError_t error_code = call; \ + if (error_code != cudaSuccess) \ + { \ + printf("CUDA ERROR: \n"); \ + printf(" FILE: %s\n", __FILE__); \ + printf(" LINE: %d\n", __LINE__); \ + printf(" ERROR CODE: %d\n", error_code); \ + printf(" ERROR TEXT: %s\n", cudaGetErrorString(error_code)); \ + exit(1); \ + } \ +}while(0); \ + + + + + + + diff --git a/capter7/global.cu b/capter7/global.cu new file mode 100644 index 0000000..8477305 --- /dev/null +++ b/capter7/global.cu @@ -0,0 +1,193 @@ +#include "error.cuh" +#include +#include +#include + +using namespace std::chrono; + + +__global__ void add(float *x, float *y, float *z, int N) +{ + // 实现顺序的合并访问。 + int n = threadIdx.y * blockDim.x + threadIdx.x; + if (n >= N) return; + + for (int i = 0; i < 1000 ; ++i) + { + z[n] = sqrt(x[n] + y[n]); + } +} + +__global__ void add_permuted(float *x, float *y, float *z, int N) +{ + // 实现乱序的合并访问(相较顺序模式，耗时增加)。 + int tid = threadIdx.x^0x1; + int n = threadIdx.y * blockDim.x + tid; + if (n >= N) return; + + for (int i = 0; i < 1000 ; ++i) + { + z[n] = sqrt(x[n] + y[n]); + } +} + +__global__ void add_offset(float *x, float *y, float *z, int N) +{ + // 实现不对齐的非合并访问(相较顺序模式，耗时增加)。 + int n = threadIdx.y * blockDim.x + threadIdx.x + 1; + if (n >= N) return; + + for (int i = 0; i < 1000 ; ++i) + { + z[n] = sqrt(x[n] + y[n]); + } +} + +__global__ void add_stride(float *x, float *y, float *z, int N) +{ + // 实现跨越式的非合并访问(相较顺序模式，耗时增加)。 + int n = blockIdx.x + threadIdx.x*gridDim.x; + if (n >= N) return; + + for (int i = 0; i < 1000 ; ++i) + { + z[n] = sqrt(x[n] + y[n]); + } +} + +__global__ void add_broadcast(float *x, float *y, float *z, int N) +{ + // 实现广播式的非合并访问(相较顺序模式，耗时增加)。 + int n = threadIdx.x + blockIdx.x*gridDim.x; + if (n >= N) return; + + for (int i = 0; i < 1000 ; ++i) + { + z[n] = sqrt(x[n] + y[n]); + } +} + +void add_cpu(float *x, float *y, float *z, int N) +{ + for (int k = 0; k < N; ++k) + { + for (int i = 0; i < 1000 ; ++i) + { + z[k] = sqrt(x[k] + y[k]); + } + } +} + + + +int main() +{ + int N = 1.0e6; + int M = N * sizeof(float); + + float *h_x, *h_y, *h_z; + h_x = new float[N]; + h_y = new float[N]; + h_z = new float[N]; + for (int i =0 ; i < N; ++i) + { + h_x[i] = 1.0; + h_y[i] = 2.0; + } + + auto t1 = system_clock::now(); + + // cpu 调用，测试加速比。 + add_cpu(h_x, h_y, h_z, N); + + auto t2 = system_clock::now(); + double time = duration(t2 - t1).count(); + std::cout << "cpu time cost: " << time << " ms" << std::endl; + + float *d_x, *d_y, *d_z; + CHECK(cudaMalloc(&d_x, M)); + CHECK(cudaMalloc(&d_y, M)); + CHECK(cudaMalloc(&d_z, M)); + CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDefault)); + CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyDefault)); + + float elapsed_time = 0; + float curr_time = 0; + cudaEvent_t start, stop; + CHECK(cudaEventCreate(&start)); + CHECK(cudaEventCreate(&stop)); + CHECK(cudaEventRecord(start)); + cudaEventQuery(start); + + // 顺序合并访问模式（一个线程块有一个线程束，一次请求 128 字节，如 d_x 中 0-31 个元素）。 + // 若 d_x 的首地址为 0，则 0-31 元素的内存分别为 0-3 字节、4-7 字节、... 124-127 字节； + // 对应 4 次数据传输 0-31 字节、32-63 字节、64-95 字节、96-127 字节，合并度 100%。 + add<<<128, 32>>>(d_x, d_y, d_z, N); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("add time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + // 乱序合并访问模式。 + add_permuted<<<128, 32>>>(d_x, d_y, d_z, N); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("add_permuted time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + // 不对齐的非合并访问模式（一个线程块依然有一个线程束，一次请求 128 字节，如 d_x 中 1-32 个元素）。 + // 若 d_x 的首地址为 0，则 1-32 元素的内存分别为 4-7 字节、... 124-127 字节、128-131 字节； + // 对应 5 次数据传输 0-31 字节、32-63 字节、64-95 字节、96-127 字节、128-159 字节， + // 合并度 4*32/(5*32) * 100% = 80%。 + add_offset<<<128, 32>>>(d_x, d_y, d_z, N); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("add_offset time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + // 跨越式非合并访问模式（一个线程块依然有一个线程束，一次请求 128 字节、32个元素）。 + // 对于第一个线程块，线程束将访问 d_x 中 0、128、256、... 等元素。 + // 因为每个元素都不不在一个 32 字节连续内存中，所以将导致 32 次数据传输， + // 合并度 4*32/(32*32) * 100% = 12.5%。 + add_stride<<<128, 32>>>(d_x, d_y, d_z, N); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("add_stride time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + // 广播式非合并访问模式（一个线程块依然有一个线程束，一次请求 128 字节、32个元素）。 + // 对于第一个线程块，线程束将一致地访问 d_x 中第 0 元素；所以只产生一次数据传输； + // 但是线程束只使用了 4 个字节，合并度 4/32 * 100% = 12.5%。 + // （这种访问更适合使用常量内存变量。） + add_broadcast<<<128, 32>>>(d_x, d_y, d_z, N); + CHECK(cudaDeviceSynchronize()); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("add_broadcast time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDefault)); + + delete[] h_x; + delete[] h_y; + delete[] h_z; + CHECK(cudaFree(d_x)); + CHECK(cudaFree(d_y)); + CHECK(cudaFree(d_z)); + + return 0; +} diff --git a/capter7/matrix.cu b/capter7/matrix.cu new file mode 100644 index 0000000..0301965 --- /dev/null +++ b/capter7/matrix.cu @@ -0,0 +1,134 @@ + +#include "error.cuh" +#include + +#ifdef USE_DP + typedef double real; + const real EPSILON = 1.0e-15; +#else + typedef float real; + const real EPSILON = 1.0e-6f; +#endif + +// using namespace std; // 不能使用std，会导致 `copy()` 不能使用（命名冲突）。 + + +__constant__ int TILE_DIM = 32; // 设备内存中线程块中矩阵维度（线程块大小，最大1024）。 + +__global__ void copy(const real *src, real *dst, const int N); +__global__ void transpose1(const real *src, real *dst, const int N); +__global__ void transpose2(const real *src, real *dst, const int N); + + +int main() +{ + const int N = 10000; + const int M = N * N * sizeof(real); + + int SIZE = 0; + CHECK(cudaMemcpyFromSymbol(&SIZE, TILE_DIM, sizeof(int))); + + const int grid_size_x = (N + SIZE - 1)/SIZE; // 获取网格大小。 + const int grid_size_y = grid_size_x; + + const dim3 block_size(SIZE, SIZE); + const dim3 grid_size(grid_size_x, grid_size_y); + + real *h_matrix_org, *h_matrix_res; + h_matrix_org = new real[N*N]; + h_matrix_res = new real[N*N]; + for (int i = 0; i < N; ++i) + { + for (int j = 0; j < N; ++j) + { + h_matrix_org[j] = i; + } + } + + float elapsed_time = 0; + float curr_time = 0; + cudaEvent_t start, stop; + CHECK(cudaEventCreate(&start)); + CHECK(cudaEventCreate(&stop)); + CHECK(cudaEventRecord(start)); + cudaEventQuery(start); + + real *d_matrix_org, *d_matrix_res; + CHECK(cudaMalloc(&d_matrix_org, M)); + CHECK(cudaMalloc(&d_matrix_res, M)); + CHECK(cudaMemcpy(d_matrix_org, h_matrix_org, M, cudaMemcpyDefault)); + + copy<<>>(d_matrix_org, d_matrix_res, N); + CHECK(cudaMemcpy(h_matrix_res, d_matrix_res, M, cudaMemcpyDefault)); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("matrix copy time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + transpose1<<>>(d_matrix_org, d_matrix_res, N); + CHECK(cudaMemcpy(h_matrix_res, d_matrix_res, M, cudaMemcpyDefault)); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("matrix transpose1 time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + transpose2<<>>(d_matrix_org, d_matrix_res, N); + CHECK(cudaMemcpy(h_matrix_res, d_matrix_res, M, cudaMemcpyDefault)); + + CHECK(cudaEventRecord(stop)); + CHECK(cudaEventSynchronize(stop)); + CHECK(cudaEventElapsedTime(&curr_time, start, stop)); + printf("matrix transpose2 time cost: %f ms.\n", curr_time - elapsed_time); + elapsed_time = curr_time; + + delete[] h_matrix_res; + delete[] h_matrix_org; + CHECK(cudaFree(d_matrix_org)); + CHECK(cudaFree(d_matrix_res)); + + return 0; +} + + +__global__ void copy(const real *src, real *dst, const int N) +{ + // TILE_DIM = blockDim.x = blockDim.y + const int nx = blockIdx.x * TILE_DIM + threadIdx.x; // 矩阵列索引。 + const int ny = blockIdx.y * TILE_DIM + threadIdx.y; // 矩阵行索引。 + const int index = ny * N + nx; + + if (nx >= N || ny >= N) + { + return; + } + + dst[index] = src[index]; // 全局内存中数组也是线性存放的。 +} + +__global__ void transpose1(const real *src, real *dst, const int N) +{ + const int nx = threadIdx.x + blockIdx.x * TILE_DIM; + const int ny = threadIdx.y + blockIdx.y * TILE_DIM; + + if (nx < N && ny < N) + { + // 矩阵转置（合并读取、非合并写入）。 + dst[nx*N + ny] = src[ny*N + nx]; + } +} + +__global__ void transpose2(const real *src, real *dst, const int N) +{ + const int nx = threadIdx.x + blockIdx.x * TILE_DIM; + const int ny = threadIdx.y + blockIdx.y * TILE_DIM; + + if (nx < N && ny < N) + { + // 矩阵转置（非合并读取、合并写入）。 + dst[ny*N + nx] = __ldg(&src[nx*N + ny]); // 显示调用 `__ldg()` 函数缓存全局内存。 + } +} \ No newline at end of file