【cuda学习日记】4.2 内存访问模式

4.2.1 缓存加载
如图，全局内存通过缓存来实现加载/存储。所有对全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存。如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现的。一行一级缓存是128个字节，它映射到设备内存中一个128字节的对齐段。如果线程束中的每个线程请求一个4字节的值，那么每次请求就会获取128字节的数据，这恰好与缓存行和设备内存段的大小相契合。

因此在优化应用程序时，你需要注意设备内存访问的两个特性：
·对齐内存访问
·合并内存访问
当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时，就会出现对齐内存访问。
当一个线程束中全部的32个线程访问一个连续的内存块时，就会出现合并内存访问。

如图是对齐与合并内存的加载操作：

如图是非对齐和未合并的内存访问。在这种情况下，可能需要3个128字节的内存事务来从设备内存中读取数据：

4.2.2 没有缓存的加载
没有缓存的加载不经过一级缓存，它在内存段的粒度上（32个字节）而非缓存池的粒度（128个字节）执行。这是更细粒度的加载，可以为非对齐或非合并的内存访问带来更好的总线利用率。

如图， 线程束请求32个连续的4字节元素但加载没有对齐到128个字节的边界。请求的地址最多落在5个内存段内，总线利用率至少为80%。

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

void initialData(float *ip, int size)
{
    time_t t;
    srand((unsigned int) time(&t));

    for (int i = 0; i < size; i++) {
        ip[i] = (float) (rand() & 0xff) / 10.0f;
    }
}

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("different on %dth element: host %f gpu %f\n", i, hostRef[i],
                    gpuRef[i]);
            break;
        }
    }

    if (!match)  printf("Arrays do not match.\n\n");
}

__global__ void readOffset(float *A , float *B, float *C, const int N, int offset){
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < N) C[i] = A[k] + B[k];
}

__global__ void warmup(float *A , float *B, float *C, const int N, int offset){
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < N) C[i] = A[k] + B[k];
}


void sumArraysOnHost(float *A, float *B, float *C, const int N, int offset)
{
    for (int idx = offset , k = 0; idx< N; idx ++, k++)
    {
        C[k] = A[idx] + B[idx];
    }
}

int main(int argc, char **argv){
    int dev = 0;
    cudaSetDevice(dev);

    unsigned int isize = 1<< 20;
    unsigned int bytes = isize * sizeof(float);

    cudaDeviceProp deviceprop;
    CHECK(cudaGetDeviceProperties(&deviceprop,dev));
    //printf("device %d: %s memory size %d bytes %5.2fMB\n",dev,deviceprop.name, isize, bytes/(1024.0f * 1024.0f) );


    int blocksize = 512;
    int offset = 0;
    if (argc > 1) offset = atoi(argv[1]);
    if (argc > 2) blocksize = atoi(argv[2]);

    dim3 block(blocksize,1);
    dim3 grid((isize + block.x - 1)/ block.x, 1);

    float *h_A = (float *)malloc(bytes);
    float *h_B = (float *)malloc(bytes);
    float *hostRef = (float *)malloc(bytes);
    float *gpuRef = (float *)malloc(bytes);

    initialData(h_A, isize);
    memcpy(h_B, h_A, bytes);
    
    sumArraysOnHost(h_A,h_B, hostRef, isize, offset);

    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, bytes);
    cudaMalloc((float**)&d_B, bytes);
    cudaMalloc((float**)&d_C, bytes);


    cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_A, bytes, cudaMemcpyHostToDevice);


    Timer timer;
    timer.start();
    warmup<<<grid,block>>>(d_A, d_B, d_C, isize, offset);
    cudaDeviceSynchronize();
    timer.stop();
    float elapsedTime = timer.elapsedms();
    //printf("warmup <<<%4d, %4d>>> offset %4d elapsed %f ms \n", grid.x, block.x, offset, elapsedTime);
    

    //
    timer.start();
    readOffset<<<grid,block>>>(d_A, d_B, d_C, isize, offset);
    cudaDeviceSynchronize();
    timer.stop();
    elapsedTime = timer.elapsedms();
    printf("readOffset <<<%4d, %4d>>> offset %4d elapsed %f ms \n", grid.x, block.x, offset, elapsedTime);


    cudaMemcpy(gpuRef, d_C, bytes, cudaMemcpyDeviceToHost);
    checkResult(hostRef, gpuRef, isize - offset);
    
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    cudaDeviceReset();
    return 0;

}

offset不同的值的时候，exeution时间没有看到明显趋势。
C:\Users\Administrator\Desktop\edward_temp\test\chapter4>readSegment.exe 0
readOffset <<<2048, 512>>> offset 0 elapsed 0.090112 ms

C:\Users\Administrator\Desktop\edward_temp\test\chapter4>readSegment.exe 11
readOffset <<<2048, 512>>> offset 11 elapsed 0.083968 ms

C:\Users\Administrator\Desktop\edward_temp\test\chapter4>readSegment.exe 128
readOffset <<<2048, 512>>> offset 128 elapsed 0.083200 ms

用NCU生成报告看：
ncu -o profile2 --metrics smsp__sass_average_data_bytes_per_sector_mem_global_op_ld readSegment.exe [0/11/128]

0 – smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct [%] 100
11 – smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct [%] 80.00
128 – smsp__sass_average_data_bytes_per_sector_mem_global_op_ld.pct [%] 100
看起来像是颗粒度32个字节。