《CUDA C编程权威指南》学习笔记 chap4-6

最新推荐文章于 2023-10-21 16:57:26 发布

_ambition_

最新推荐文章于 2023-10-21 16:57:26 发布

阅读量498

点赞数

分类专栏： CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/ambition_zhou/article/details/116522105

版权

CUDA 专栏收录该内容

3 篇文章 2 订阅

订阅专栏

《CUDA C编程权威指南》示例代码下载

第4章全局内存

4.1 CUDA内存模型概述

1. 内存层次结构的优点

在这里插入图片描述
CUDA编程模型能将内存层次结构更好地呈现给用户，能让我们显式地控制它的行为。

2. CUDA内存模型

CUDA内存模型提出了多种可编程内存的类型： ·寄存器 ·共享内存 ·本地内存 ·常量内存 ·纹理内存 ·全局内存
在这里插入图片描述
所有线程都可以访问全局内存。所有线程都能访问的只读内存空间有：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间有不同的用途。纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。对于一个应用程序来说，全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

总结

在这里插入图片描述
+既可以表明标量也可以表示数组

在这里插入图片描述

寄存器
核函数中声明的一个没有其他修饰符的自变量，通常存储在寄存器中。在核函数声明的数组中，如果用于引用该数组的索引是常量且能在编译时确定，那么该数组也存储在寄存器中。
寄存器变量对于每个线程来说都是私有的，一个核函数通常使用寄存器来保存需要频繁访问的线程私有变量。寄存器变量与核函数的生命周期相同。一旦核函数执行完毕，就不能对寄存器变量进行访问了。
寄存器是一个在SM中由活跃线程束划分出的较少资源。
如果一个核函数使用了超过硬件限制数量的寄存器，则会用本地内存替代多占用的寄存器。 nvcc编译器使用启发式策略来最小化寄存器的使用，以避免寄存器溢出。
本地内存
核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。编译器可能存放到本地内存中的变量有：
·在编译时使用未知索引引用的本地数组
·可能会占用大量寄存器空间的较大本地结构体或数组
·任何不满足核函数寄存器限定条件的变量
溢出到本地内存中的变量本质上与全局内存在同一块存储区域，因此本地内存访问的特点是高延迟和低带宽
共享内存
在核函数中使用如下修饰符修饰的变量存放在共享内存中：

__shared__

因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。它的使用类似于CPU一级缓存，但它是可编程的。
共享内存是线程之间相互通信的基本方式，生命周期伴随着整个线程块。当一个线程块执行结束后，其分配的共享内存将被释放并重新分配给其他线程块。访问共享内存必须同步使用如下调用：

void __syncthreads():

SM中的一级缓存和共享内存都使用64KB的片上内存（两者共用64KB），它通过静态划分，但在运行时可以通过如下指令进行动态配置：

cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);

支持的缓存配置如下：
在这里插入图片描述
4. 常量内存
修饰符号：__constant__
常量变量必须在全局空间内和所有核函数之外进行声明。对于所有计算能力的设备，都只可以声明64KB的常量内存。常量内存是静态声明的，并对同一编译单元中的所有核函数可见。
初始化：
将count个字节从src指向的内存复制到symbol指向的内存中，这个变量存放在设备的全局内存或常量内存中

cudaError_t cudaMemcypToSymbol(const void* symbol, const void* src, size_t count);

纹理内存
纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。
全局内存
全局内存是GPU中最大、延迟最高并且最常使用的内存。global指的是其作用域和生命周期。它的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。
修饰符号：__device__
分配过程：主机端cudaMalloc分配全局内存，cudaFree释放全局内存。指向全局内存的指针作为参数传递给核函数
GPU缓存
GPU缓存不可编程，有4种缓存：·一级缓存 ·二级缓存 ·只读常量缓存 ·只读纹理缓存
每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。
静态全局内存示例

#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 * An example of using a statically declared global variable (devData) to store
 * a floating-point value on the device.
 */

__device__ float devData;

__global__ void checkGlobalVariable()
{
    // display the original value
    printf("Device: the value of the global variable is %f\n", devData);

    // alter the value
    devData += 2.0f;
}

int main(void)
{
    // initialize the global variable
    float value = 3.14f;
    CHECK(cudaMemcpyToSymbol(devData, &value, sizeof(float)));
    printf("Host:   copied %f to the global variable\n", value);

    // invoke the kernel
    checkGlobalVariable<<<1, 1>>>();

    // copy the global variable back to the host
    CHECK(cudaMemcpyFromSymbol(&value, devData, sizeof(float)));
    printf("Host:   the value changed by the kernel to %f\n", value);

    CHECK(cudaDeviceReset());
    return EXIT_SUCCESS;
}

在这里插入图片描述

4.2 内存管理

1. 内存分配和释放

//在主机上使用下列函数分配全局内存：
cudaEroor_t cudaMalloc(void ** devPtr, size_t count); 
//使用value中的值填充从设备内存地址devPtr处开始的count字节
cudaError_t cudaMemset(void *devPtr, int value, size_t count); 
// 释放内存空间
cudaError_t cudaFree(void* devPtr);

2.内存传输

// 从主机向设备传输数据,从内存位置src复制了count字节到内存位置dst
// kind取值：cudaMemcpuHostToHost,cudaMemcpyHostToDevice,cudaMemcpyDeviceTohost,cudMemcpyDeviceToDevice
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, enum cudaMemcpyKind kind); // 同步操作

GPU芯片和GPU板载内存之间的理论峰值带宽很高(Fermi C2050GPU 144GB/s)，但是CPU和GPU之间的带宽较低（CPIe Gen3总线 8GB/s)

3.固定内存

分配的主机内存默认是pageable（可分页），GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存
在这里插入图片描述
CUDA运行时允许你使用如下指令直接分配固定主机内存：
固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。

//分配了count字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的
cudaError_t cudaMallocHost(void** devPtr, size_t count);
// 固定内存的释放
cudaErro_t cudaFreeHost(void* ptr);

当传输超过10MB的数据时，在Fermi设备上使用固定内存通常是更好的选择。

4.零拷贝内存

据说用不到，暂时不看
通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。但有一个例外：零拷贝内存。主机和设备都可以访问零拷贝内存。
零拷贝内存是固定（不可分页）内存，该内存映射到设备地址空间中。

//创建一个到固定内存的映射
cudaError_t cudaHostAlloc(void** pHost, size_t count, unsigned int flags);
// 释放
cudaErro_t cudaFreeHost(void* ptr);
// 获取到固定内存的设备指针,
// 返回一个在pDevice中的设备指针，该指针可以在设备上被引用以访问映射得到的固定主机内存。
cudaError_t cudaHostGetDevicePointer(void** pDevice, void* pHost, unsigned int flags);

flag参数：

cudaHostAllocDefault函数使cudaHostAlloc函数的行为与cudaMallocHost函数一致（固定内存）。
cudaHostAllocPortable函数可以返回能被所有CUDA上下文使用的固定内存，而不仅是执行内存分配的那一个。
cudaHostAllocWriteCombined返回写结合内存，该内存可以在某些系统配置上通过PCIe总线上更快地传输，但是它在大多数主机上不能被有效地读取。因此，写结合内存对缓冲区来说是一个很好的选择，该内存通过设备使用映射的固定内存或主机到设备的传输。
cudaHostAllocMapped是零拷贝内存的最明显的标志，该标志返回，可以实现主机写入和设备读取被映射到设备地址空间中的主机内存。

   // part 2: using zerocopy memory for array A and B
    // allocate zerocpy memory
    CHECK(cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped));
    CHECK(cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped));

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(hostRef, 0, nBytes);
    memset(gpuRef,  0, nBytes);

    // pass the pointer to device
    CHECK(cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0));
    CHECK(cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0));

    // add at host side for result checks
    sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // execute kernel with zero copy memory
    sumArraysZeroCopy<<<grid, block>>>(d_A, d_B, d_C, nElem);

    // copy kernel result back to host side
    CHECK(cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost));

    // check device results
    checkResult(hostRef, gpuRef, nElem);

    // free  memory
    CHECK(cudaFree(d_C));
    CHECK(cudaFreeHost(h_A));
    CHECK(cudaFreeHost(h_B));

    free(hostRef);
    free(gpuRef);

    // reset device
    CHECK(cudaDeviceReset());

// 获取映射到固定内存的设备指针
cudaError_t cudaHostGetDevicePointer(void** pDevice, void* pHost, unsigned int flags);

5.统一虚拟寻址（UVA）

有了UVA，主机内存和设备内存可以共享同一个虚拟地址空间
在这里插入图片描述
在UVA之前，你需要管理哪些指针指向主机内存和哪些指针指向设备内存。有了 UVA，由指针指向的内存空间对应用程序代码来说是透明的。
通过UVA，由cudaHostAlloc分配的固定主机内存具有相同的主机和设备指针，可以将返回的指针直接传递给核函数。

 CHECK(cudaHostAlloc((void **)&h_A, nBytes, cudaHostAllocMapped));
 CHECK(cudaHostAlloc((void **)&h_B, nBytes, cudaHostAllocMapped));

 // initialize data at host side
 initialData(h_A, nElem);
 initialData(h_B, nElem);

 // 下面这一部分不再需要了
 // pass the pointer to device
 // CHECK(cudaHostGetDevicePointer((void **)&d_A, (void *)h_A, 0));
 // CHECK(cudaHostGetDevicePointer((void **)&d_B, (void *)h_B, 0));

 // 直接使用指针h_A h_B
 // execute kernel with zero copy memory
 sumArraysZeroCopy<<<grid, block>>>(h_A, h_B, d_C, nElem);

6.统一内存寻址

在CUDA 6.0中，引入了“统一内存寻址”这一新特性，它用于简化CUDA编程模型中的内存管理。统一内存中创建了一个托管内存池，内存池中已分配的空间可以用相同的内存地址（即指针）在CPU和GPU上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。
统一内存寻址依赖于UVA的支持，但它们是完全不同的技术。UVA为系统中的所有处理器提供了一个单一的虚拟内存地址空间。但是，UVA不会自动将数据从一个物理位置转移到另一个位置，这是统一内存寻址的一个特有功能。
统一内存寻址提供了一个**“单指针到数据”模型**，在概念上它类似于零拷贝内存。

// 静态声明一个设备变量作为托管变量
__managed__
// 动态分配托管内存
// 分配size字节的托管内存，并用devPtr返回一个指针。该指针在所有设备和主机上都是有效的
cudaError_t cudaMallocManaged(void** devPtr, size_t size, unsigned int flag = 0);

4.3 内存访问模式

1. 对齐与合并访问

核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现的。这几句话不懂
设备内存访问的两个特性：-对齐内存访问 -合并内存访问

在这里插入图片描述

优化内存事务效率：用最少的事务次数满足最多的内存请求

2. 全局内存读取

数据通过以下3种缓存/缓冲路径进行传输：·一级和二级缓存 ·常量缓存 ·只读缓存
禁用一级缓存标志： -Xptxas -dlcm=cg
启用一级缓存标志： -Xptxas -dlcm=ca

内存加载：

缓存加载（启用一级缓存
缓存加载操作经过一级缓存，在粒度为128字节的一级缓存行上由设备内存事务进行传输。缓存加载可以分为对齐/非对齐及合并/非合并。

(上图情况利用率只有50%：启用了一级缓存，SM执行的物理家在操作必须128B对齐）

CPU一级缓存和GPU一级缓存之间的差异： CPU一级缓存优化了时间和空间局部性。GPU一级缓存是专为空间局部性而不是为时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。
没有缓存的加载（禁用一级缓存
更细粒度的加载（32B），可以为非对齐或非合并的内存访问带来更好的总线利用率

3. 全局内存写入

一级缓存不能用在Fermi或Kepler GPU上进行存储操作，在发送到设备内存之前存储操作**只通过二级缓存。**存储操作在32个字节段的粒度上被执行。
在这里插入图片描述

4. 结构体数组与数组结构体

// 数组结构体AoS
struct innerStruct{
	float x;
	float y;
}
struct innerStruct myAos[N]
// 结构体数组SoA
struct innerArray{
	float x[N];
	float y[N];
}

__global__ void testInnerStruct(innerStruct *data, innerStruct * result,
                                const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

对于AOS数据布局，加载请求和内存存储请求是重复的。因为字段x和y在内存中是被相邻存储的，并且有相同的大小，每当执行内存事务时都要加载特定字段的值，被加载的字节数的一半也必须属于其他字段。因此，请求加载和存储的50%带宽是未使用的。

5. 性能调整

优化设备内存带宽利用率有两个目标：
· 对齐及合并内存访问，以减少带宽的浪费
· 足够的并发内存操作，以隐藏内存延迟

实现并发内存访问最大化是通过以下方式获得的：
· 增加每个线程中执行独立内存操作的数量
· 对核函数启动的执行配置进行实验，以充分体现每个SM的并行性

4.4 核函数可达到的带宽

内存延迟，即完成一次独立内存请求的时间；
内存带宽，即SM访问设备内存的速度，

1. 内存带宽

理论带宽是当前硬件可以实现的绝对最大带宽。
有效带宽是核函数实际达到的带宽，它是测量带宽
在这里插入图片描述
即使是较低的加载效率，一级缓存中的缓存加载也可以限制交叉加载对性能的负面影响

2. 矩阵装置问题

矩阵访问模式

// case 2 transpose kernel: read in rows and write in columns
// 按行读取按列存储
__global__ void transposeNaiveRow(float *out, float *in, const int nx,
                                  const int ny)
{
    unsigned int ix = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;

    if (ix < nx && iy < ny)
    {
        out[ix * ny + iy] = in[iy * nx + ix];
    }
}

// case 3 transpose kernel: read in columns and write in rows
// 按列读取按行存储
__global__ void transposeNaiveCol(float *out, float *in, const int nx,
                                  const int ny)
{
    unsigned int ix = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;

    if (ix < nx && iy < ny)
    {
        out[iy * nx + ix] = in[ix * ny + iy];
    }
}

在这里插入图片描述
通过缓存交叉读取能够获得最高的加载吞吐量
2. 展开技术

// case 4 transpose kernel: read in rows and write in columns + unroll 4 blocks
__global__ void transposeUnroll4Row(float *out, float *in, const int nx,
                                    const int ny)
{
    unsigned int ix = blockDim.x * blockIdx.x * 4 + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;

    unsigned int ti = iy * nx + ix; // access in rows
    unsigned int to = ix * ny + iy; // access in columns

    if (ix + 3 * blockDim.x < nx && iy < ny)
    {
        out[to]                   = in[ti];
        out[to + ny * blockDim.x]   = in[ti + blockDim.x];
        out[to + ny * 2 * blockDim.x] = in[ti + 2 * blockDim.x];
        out[to + ny * 3 * blockDim.x] = in[ti + 3 * blockDim.x];
    }
}

调整块大小
增加并行性最简单的方式是调整块的大小。

4.5 使用统一内存的矩阵加法


    // malloc host memory
    float *A, *B, *hostRef, *gpuRef;
    CHECK(cudaMallocManaged((void **)&A, nBytes));
    CHECK(cudaMallocManaged((void **)&B, nBytes));
    CHECK(cudaMallocManaged((void **)&gpuRef,  nBytes);  );
    CHECK(cudaMallocManaged((void **)&hostRef, nBytes););

    // initialize data at host side
    initialData(A, nxy);
    initialData(B, nxy);

    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);

    // add matrix at host side for result checks
    sumMatrixOnHost(A, B, hostRef, nx, ny);

    // invoke kernel at host side
    int dimx = 32;
    int dimy = 32;
    dim3 block(dimx, dimy);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);

    // warm-up kernel, with unified memory all pages will migrate from host to
    // device
    sumMatrixGPU<<<grid, block>>>(A, B, gpuRef, 1, 1);

    // after warm-up, time with unified memory
    sumMatrixGPU<<<grid, block>>>(A, B, gpuRef, nx, ny);
    
    CHECK(cudaDeviceSynchronize());
    
    // check kernel error
    CHECK(cudaGetLastError());

    // check device results
    checkResult(hostRef, gpuRef, nxy);

    // free device global memory
    CHECK(cudaFree(A));
    CHECK(cudaFree(B));
    CHECK(cudaFree(hostRef));
    CHECK(cudaFree(gpuRef));

在CUDA应用中手动优化数据移动的性能比使用统一内存的性能要更优

第5章共享内存和常量内存

5.1 CUDA共享内存概述

两种内存：板载内存、片上内存
全局内存是较大的板载内存，具有相对较高的延迟。共享内存是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽。

1. 共享内存（shared memory，SMEM）

共享内存相较于全局内存而言，延迟要低大约20～30倍，而带宽高其大约10倍
在这里插入图片描述
共享内存的地址空间被线程块中所有的线程共享，

2. 共享内存分配

声明：__shared__

可以静态或动态地分配共享内存变量。

__shared__ float tile[size_y][xize_x]; // 静态分配
__shared__ int tile[]; // 动态分配，通过<<<>>>参数设置，但是只能动态声明一维动态数组

共享内存可以被声明为一个本地或全局的 CUDA核函数
声明位置在核函数外还是核函数内

// 可在某个核函数内部或所有核函数外部进行声明的例子
extern __shared__ int tile[];

3. 共性内存存储体和访问模式

内存存储体
共享内存被分为32个同样大小的内存模型，它们被称为存储体，可以被同时访问。
存储体冲突
在共享内存中当多个地址请求落在相同的内存存储体中时，就会发生存储体冲突，这会导致请求被重复执行。硬件会将存储体冲突的请求分割到尽可能多的独立的无冲突事务中

并行访问：多个地址访问多个存储体
串行访问：多个地址访问同一个存储体（最坏模式），必须以串行方式进行请求
广播访问：单一地址读取单一存储体，线程束中所有的线程都读取同一存储体中相同的地址若一个内存事务被执行，那么被访问的字就会被广播到所有请求的线程中。

两种模式（每个存储体有多少位）：64位模式、32位模式
存储体索引＝（字节地址÷4字节/存储体）%32存储体（32位模式例子）
字节地址除以4转换为一个4字节字索引，然后进行模32操作，将4字节字索引转换为存储体索引。

内存填充
内存填充是避免存储体冲突的一种方法，填充的内存不能用于数据存储。其唯一的作用是移动数据元素，以便将原来属于同一个存储体中的数据分散到不同存储体中。
访问模式配置
Kepler设备支持4字节和8字节的共享内存访问模式，默认是4字节模式。

// 访问模式查询：
cudaError_t cudaDeviceGetSharedMemConfig(cudaSharedMemConfig * pConfig);
// 设置存储体大小
cudaError_t cudaDeviceSetSharedMemConfig(cudaSharedMemConfig config);
// 支持的存储体配置：cudaSharedMemBankSizeDefault/FourByte/EightByte

4. 配置共享内存量

每个SM都有64 KB的片上内存（L1Cache和共享内存共用，但不重叠）。
配置方法：按设备、按核函数

// 1. 按核函数
cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);

// 2.由核函数指针func指定配置
cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cheConfig);

支持的缓存配置如下：
在这里插入图片描述
对Kepler设备而言，一级缓存用于寄存器溢出，指定-Xptxas- v选项给nvcc，可以知道核函数使用了多少寄存器。当内核使用的寄存器数量超过了硬件限制所允许的数量时，应该为寄存器溢出配置一个更大的一级缓存。
对Fermi设备而言，本地内存用于溢出寄存器，但本地内存的加载可能被缓存在一级缓存中。

5. 同步

同步基本方法：障碍、内存栅栏。
为了显式地强制程序以一个确切的顺序执行，必须在应用程序代码中插入内存栅栏和障碍。这是保证与其他线程共享资源的核函数行为正确的唯一途径。

显式障碍

void __syncthreads();

在障碍中，所有调用的线程等待其余调用的线程到达障碍点。
__syncthreads还确保在障碍点之前，被这些线程访问的所有全局和共享内存对同一块中的所有线程都可见。
错误示范：

// 块中的所有线程没有达到相同的障碍点。
if(threadID % 2 == 0) 
	__syncthreads();
else 
	__syncthreads();

内存栅栏
在内存栅栏中，所有调用的线程必须等到全部内存修改对其余调用线程可见时才能继续执行。
内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。
有3种内存栅栏：块、网格或系统

// 1. 线程块内创建内存栅栏
// 保证了栅栏前被调用线程产生的对共享内存和全局内存的所有写操作
// 对栅栏后同一块中的其他线程都是可见的
// 内存栅栏不执行任何线程同 步，所以对于一个块中的所有线程来说，没有必要实际执行这个指令。
void __threadfence_block();

// 2. 创建网格级内存栅栏
// 挂起调用的线程，直到全局内存中的所有写操作对相同网格内的所有线 程都是可见的
vlid __threadfence();

//3. 跨系统（包括主机和设备）设置内存栅栏
// 挂起调用的线程，以确保该线程对全局内存、锁页主机内存和其他设备内存中的所有写操作
// 对全部设备中的线程和主机线程是可见的。
void __threadfence_system();

volatile修饰符：防止编译器优化，编译器假定任何其他线程在任何时间都可以更改或使用该变量的值。因此，这个变量的任何引用都会直接被编译到全局内存读指令或全局内存写指令中，它们都会忽略缓存。
共享内存&全局内存：
延迟和带宽：DRAM比共享内存高20～30倍的延迟 ·SMEM比DRAM大10倍的带宽
访问粒度：DRAM32B或128B；SMEM：4B或8B

5.2 共享内存的数据布局

1. 方形共享内存（row==col)

在这里插入图片描述
访问方法：

// 最好是有访问共享内存连续位置的线程，且该线程带有 连续的threadIdx.x值
// 所以第一种方法更好
tile[threadIdx.y][threadIdx.x]
tile[threadIdx.x][threadIdx.y]

行主序列访问和列主序列访问：
核函数有两个简单操作：
·将全局线程索引按行主序写入到一个二维共享内存数组中
·从共享内存中按行主序读取这些值并将它们存储到全局内存中

按行主序索引

__global__ void setRowReadRow (int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.y][threadIdx.x] ;
}

三个内存操作：·共享内存的存储操作 ·共享内存的加载操作 ·全局内存的存储操作

在将数据分配给共享内存块时交换threadIdx.y和threadIdx.x，线程束的内存将会按列主序访问
每个共享内存的加载和存储将导致Fermi装置中有32路存储体冲突，导致Kepler装置中有16路存储体冲突。

__global__ void setColReadCol (int *out)
{
    // static shared memory
    __shared__ int tile[BDIMX][BDIMY];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

动态共享内存
动态共享内存必须被声明为一个未定大小的一维数组，因此，需要基于二维线程索引来计算内存访问索引。
为了确保合并存储，需要通过线程坐标按行主序对out数组写入

__global__ void setRowReadColDyn(int *out)
{
    // dynamic shared memory
    extern  __shared__ int tile[];

    // mapping from thread index to global memory index
    unsigned int row_idx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int col_idx = threadIdx.x * blockDim.y + threadIdx.y;

    // shared memory store operation
    tile[row_idx] = row_idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[row_idx] = tile[col_idx];
}

// 启动内核时，必须指定共享内存的大小
setRowReadColDyn<<<grid, block, BDIMX*BDIMY*sizeof(int)>>>(d_C);

使用nvprof检查setRowReadColDyn核函数的内存事务
在这里插入图片描述
写操作无冲突，读操作有16路冲突。
4. 填充静态声明的共享内存
只需简单地将一列添加到二维共享内存分配中。
setRowReadCol按列主序读取时报告了16路冲突。通过在每行添加一个元素，列元素便分布在了不同的存储体中，因此读和写操作都是无冲突的。

__shared__ int tile[BDIMY][BDIMX+1];

__global__ void setRowReadColPad(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX + IPAD];

    // mapping from thread index to global memory offset
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

填充动态声明的共享内存
当执行从二维线程索引到一维内存索引的索引转换时，对于每一行必须跳过一个填充的内存空间

unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
    unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;

在这里插入图片描述

2. 矩形共享内存(row != col)

按行主序写入共享内存，并按列主序读取共享内存
内核有3个内存操作：
·写入每个线程束的共享内存行，以避免存储体冲突
·读取每个线程束中的共享内存列，以完成矩阵转置
·使用合并访问写入每个线程束的全局内存行

__global__ void setRowReadCol(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from 2D thread index to linear memory
    // 将当前线程的二维线程索 引转换为一维全局线程ID：
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // convert idx to transposed coordinate (row, col)
    // 输出的全局内存中的数 据元素是转置过的，所以需要计算转置矩阵中的新坐标
    unsigned int irow = idx / blockDim.y;
    unsigned int icol = idx % blockDim.y;

    // shared memory store operation
    // 行主序写入 所以没有写入冲突
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[icol][irow];
}

5.3 减少全局内存访问

1. 使用共享内存的并行归约

原始函数：

__global__ void reduceGmem(int *g_idata, int *g_odata, unsigned int n)
{
    // set thread ID
    unsigned int tid = threadIdx.x;
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // boundary check
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= n) return;

    // in-place reduction in global memory
    if (blockDim.x >= 1024 && tid < 512) idata[tid] += idata[tid + 512];

    __syncthreads();

    if (blockDim.x >= 512 && tid < 256) idata[tid] += idata[tid + 256];

    __syncthreads();

    if (blockDim.x >= 256 && tid < 128) idata[tid] += idata[tid + 128];

    __syncthreads();

    if (blockDim.x >= 128 && tid < 64) idata[tid] += idata[tid + 64];

    __syncthreads();

    // unrolling warp
    if (tid < 32)
    {
        volatile int *vsmem = idata;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid +  8];
        vsmem[tid] += vsmem[tid +  4];
        vsmem[tid] += vsmem[tid +  2];
        vsmem[tid] += vsmem[tid +  1];
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

使用共享内存：

__global__ void reduceSmem(int *g_idata, int *g_odata, unsigned int n)
{
    __shared__ int smem[DIM];

    // set thread ID
    unsigned int tid = threadIdx.x;

    // boundary check
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= n) return;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // set to smem by each threads
    smem[tid] = idata[tid];
    __syncthreads();

    // in-place reduction in shared memory
    if (blockDim.x >= 1024 && tid < 512) smem[tid] += smem[tid + 512];

    __syncthreads();

    if (blockDim.x >= 512 && tid < 256) smem[tid] += smem[tid + 256];

    __syncthreads();

    if (blockDim.x >= 256 && tid < 128) smem[tid] += smem[tid + 128];

    __syncthreads();

    if (blockDim.x >= 128 && tid < 64)  smem[tid] += smem[tid + 64];

    __syncthreads();

    // unrolling warp
    if (tid < 32)
    {
        volatile int *vsmem = smem;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid +  8];
        vsmem[tid] += vsmem[tid +  4];
        vsmem[tid] += vsmem[tid +  2];
        vsmem[tid] += vsmem[tid +  1];
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = smem[0];
}

2. 使用展开的并行归约

__global__ void reduceGmemUnroll(int *g_idata, int *g_odata, unsigned int n)
{
    // set thread ID
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;

    // convert global data pointer to the local pointer of this block
    int *idata = g_idata + blockIdx.x * blockDim.x * 4;

    // unrolling 4
    if (idx < n)
    {
        int a1, a2, a3, a4;
        a1 = a2 = a3 = a4 = 0;
        a1 = g_idata[idx];
        if (idx + blockDim.x < n) a2 = g_idata[idx + blockDim.x];
        if (idx + 2 * blockDim.x < n) a3 = g_idata[idx + 2 * blockDim.x];
        if (idx + 3 * blockDim.x < n) a4 = g_idata[idx + 3 * blockDim.x];
        g_idata[idx] = a1 + a2 + a3 + a4;
    }

    __syncthreads();

    // in-place reduction in global memory
    if (blockDim.x >= 1024 && tid < 512) idata[tid] += idata[tid + 512];

    __syncthreads();

    if (blockDim.x >= 512 && tid < 256) idata[tid] += idata[tid + 256];

    __syncthreads();

    if (blockDim.x >= 256 && tid < 128) idata[tid] += idata[tid + 128];

    __syncthreads();

    if (blockDim.x >= 128 && tid < 64) idata[tid] += idata[tid + 64];

    __syncthreads();

    // unrolling warp
    if (tid < 32)
    {
        volatile int *vsmem = idata;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid +  8];
        vsmem[tid] += vsmem[tid +  4];
        vsmem[tid] += vsmem[tid +  2];
        vsmem[tid] += vsmem[tid +  1];
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

3. 使用动态共享内内存的并行归约

// 声明
extern __shared__ int smem[];
// 启动 必须指定待动态分配的共享内存数量
reduceSmemUnrollDyn<<<grid.x / 4, block, DIM*sizeof(int)>>>(d_idata, d_odata, size);
// 代码其他部分和静态分配一样

用nvprof计算核函数的运行时间，那么会发现用动态分配共享内存实现的核函数和用静态分配共享内存实现的核函数之间没有显著的差异。

4. 有效带宽

有效带宽＝（读字节＋写字节）÷（运行时间×109）GB/s

5.4 合并的全局内存访问

使用共享内存也能帮助避免对未合并的全局内存的访问。

1. 基准转置内核

#define INDEX(ROW, COL, INNER) ((ROW) * (INNER) + (COL))

__global__ void copyGmem(float *out, float *in, const int nrows, const int ncols)
{
    // matrix coordinate (ix,iy)
    unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int col = blockIdx.x * blockDim.x + threadIdx.x;

    // transpose with boundary test
    if (row < nrows && col < ncols)
    {
		// NOTE this is a transpose, not a copy
        out[INDEX(col, row, nrows)] = in[INDEX(row, col, ncols)];
    }
}

2. 使用共享内存的矩阵转置

为了避免交叉全局内存访问，可以使用二维共享内存来缓存原始矩阵的数据。
这个又不懂了，难受
在这里插入图片描述
1.线程束执行合并读取一行，该行存储在全局内存中的原始矩阵块中。
2.然后，该线程束按行主序将该数据写入共享内存中，因此，这个写操作没有存储体冲突。
3.因为线程块的读/写操作是同步的，所以会有一个填满全局内存数据的二维共享内存数组。
4.该线程束从二维共享内存数组中读取一列。由于共享内存没有被填充，所以会发生存储体冲突。
5.然后该线程束执行数据的合并写入操作，将其写入到全局内存的转置矩阵中的某行。

__global__ void transposeSmem(float *out, float *in, int nrows, int ncols)
{
    // static shared memory
    __shared__ float tile[BDIMY][BDIMX];

    // coordinate in original matrix
    unsigned int row = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;

    // linear global memory index for original matrix
    unsigned int offset = INDEX(row, col, ncols);

    if (row < nrows && col < ncols)
    {
      // load data from global memory to shared memory
      tile[threadIdx.y][threadIdx.x] = in[offset];
    }

    // thread index in transposed block
    unsigned int bidx, irow, icol;
    bidx = threadIdx.y * blockDim.x + threadIdx.x;
    irow = bidx / blockDim.y;
    icol = bidx % blockDim.y;

	  // NOTE - need to transpose row and col on block and thread-block level:
	  // 1. swap blocks x-y
	  // 2. swap thread x-y assignment (irow and icol calculations above)
	  // note col still has continuous threadIdx.x -> coalesced gst
	  col = blockIdx.y * blockDim.y + icol;
	  row = blockIdx.x * blockDim.x + irow;

    // linear global memory index for transposed matrix
	  // NOTE nrows is stride of result, row and col are transposed
    unsigned int transposed_offset = INDEX(row, col, nrows);
    // thread synchronization
    __syncthreads();

	  // NOTE invert sizes for write check
    if (row < ncols && col < nrows)
    {
        // store data to global memory from shared memory
        out[transposed_offset] = tile[icol][irow]; // NOTE icol,irow not irow,icol
    }
}

3. 使用填充共享内存的矩阵转置

4. 使用展开的矩阵转置

5. 增大并行性

5.5 常量内存

常量内存是一种专用的内存，它用于只读数据和统一访问线程束中线程的数据

// 声明
__constant__ float coef[SIZE];

// 初始化
// 设备只能读取常 量内存，所以常量内存中的值必须使用以下运行时函数进行初始化
cudaErrot_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count, size_t offset, cudaMemcpyKind kind);

常量内存位于设备的DRAM上（和全局内存一样），并且有一个专用的片上缓存。每个SM常量内存缓存大小的限制为64KB。
常量内存有一个不同的 最优访问模式。在常量内存中，如果线程束中的所有线程都访问相同的位置，那么这个访问模式就是最优的。

1. 使用常量内存实现一维模板

2. 与只读缓存的比较

每个Kepler SM都有48KB的只读缓存,只读缓存的粒度为32 个字节。

5.6 线程池洗牌指令

洗牌指令使得线程束中的线程彼此之间可以直接交换数据，而不是通过共享内存或全局内存来进行的。
洗牌指令（shuffle instruction）作为一种机制被加入其中，只要两个线程在相同的线程束中，那么就允许这两个线程直接读取另一个线程的寄存器。洗牌指令比共享内存有更低的延迟，并且该指令在执行数据交换时不消耗额外的内存。

1. 线程束洗牌指令的不同形式

2. 线程束的共享数据

3. 使用线程束洗牌指令的并行归约

暂时停更

先去看SEAL

_ambition_

关注

0
点赞
踩
1

收藏

觉得还不错? 一键收藏
2
评论
《CUDA C编程权威指南》学习笔记 chap4-6

《CUDA C编程权威指南》示例代码下载第4章全局内存4.1 CUDA内存模型概述1. 内存层次结构的优点CUDA编程模型能将内存层次结构更好地呈现给用户，能让我们显式地控制它的行为。2. CUDA内存模型CUDA内存模型提出了多种可编程内存的类型： ·寄存器 ·共享内存 ·本地内存 ·常量内存 ·纹理内存 ·全局内存所有线程都可以访问全局内存。所有线程都能访问的只读内存空间有：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间有不同的用途。纹理内存为各种数据布局提供了不
复制链接

扫一扫