如何在 CUDA 中使用 64 位指针编写指针追踪基准？

Question

This research paper 在 GPU 上运行一系列多个 CUDA 微基准，以获得全局内存延迟、指令吞吐量等统计数据。This link 是作者编写并在其 GPU 上运行的微基准集的链接。

其中一个名为 global.cu 的微基准测试提供了用于测量全局内存延迟的指针追踪基准测试的代码。

这是运行的内核代码。

__global__ void global_latency (unsigned int ** my_array, int array_length, int iterations, int ignore_iterations, unsigned long long * duration) {

    unsigned int start_time, end_time;
    unsigned int *j = (unsigned int*)my_array; 
    volatile unsigned long long sum_time;

    sum_time = 0;
    duration[0] = 0;

    for (int k = -ignore_iterations; k < iterations; k++) {
        if (k==0) {
            sum_time = 0; // ignore some iterations: cold icache misses
        }

        start_time = clock();
        repeat256(j=*(unsigned int **)j;) // unroll macro, simply creates an unrolled loop of 256 instructions, nothing more
        end_time = clock();

        sum_time += (end_time - start_time);
    }

    ((unsigned int*)my_array)[array_length] = (unsigned int)j;
    ((unsigned int*)my_array)[array_length+1] = (unsigned int) sum_time;
    duration[0] = sum_time;
}

在 32 位指针情况下执行指针追逐的代码行是：

j = *(unsigned int**)j;

这是关键行，因为剩下的代码行只用于时间测量。

我尝试在我的 GPU 上运行它，但遇到了一个问题。在没有任何更改的情况下运行相同的微基准测试会给我一个运行时错误An illegal memory access was encountered。

In the same link 他们解释说：

全局内存测试使用指针追踪代码，其中指针值存储在数组中。 GT200 上的指针是 32 位的。如果指针大小发生变化，则需要更改全局内存测试，例如 Fermi 上的 64 位指针。

原来我的 GPU 是 Kepler 架构的，它有 64 位指针。

如何修改最初处理 32 位指针的指针追踪代码，以便使用 64 位指针测量全局内存延迟？

编辑：

来自 havogt 的回答：我应该在问题中包含的一条重要信息是这部分代码，其中构建了一个内存位置数组，每个入口指向下一个指针的条目。

for (i = 0; i < N; i += step) {
    // Device pointers are 32-bit on GT200.
    h_a[i] = ((unsigned int)(uintptr_t)d_a) + ((i + stride) % N)*sizeof(unsigned int);
}

Answer 1

简介

在我解释如何使代码正常工作之前，让我强调以下几点：您应该非常了解您正在测试的硬件和您的微基准测试的设计。它为什么如此重要？ 原始代码是为 GT200 设计的，它没有用于普通全局内存加载的缓存。如果您现在只修复指针问题，您将基本上测量 L2 延迟（在 Kepler 上，默认情况下不使用 L1），因为原始代码使用非常小的内存，非常适合缓存。

免责声明：对我来说这也是第一次研究这样的基准代码。因此，在使用下面的代码之前，请仔细检查。我不保证我在转换原始代码时没有出错。

简单的解决方案（基本上测量缓存延迟）

首先，您没有在问题中包含代码的所有相关部分。最重要的部分是

for (i = 0; i < N; i += step) {
    // Device pointers are 32-bit on GT200.
    h_a[i] = ((unsigned int)(uintptr_t)d_a) + ((i + stride) % N)*sizeof(unsigned int);
}

其中构建了一个内存位置数组，其中每个条目指向下一个指针的条目。 现在您需要做的就是在设置代码和内核中将所有unsigned int（用于存储32 位指针）替换为unsigned long long int。

我不会发布代码，因为如果您不理解，我不建议您运行此类代码，请参阅简介。懂了就简单了。

我的解决方案

基本上我所做的是使用尽可能多的内存来评估所有指针或最大内存量为 1GB。在这两种情况下，我都将最后一个条目包装到第一个条目。请注意，根据步幅，许多数组条目可能未初始化（因为它们从未使用过）。

下面的代码基本上是原代码经过一番清理（但还是不是很干净，不好意思……）和内存的变化。我介绍了一个 typedef

typedef unsigned long long int ptrsize_type;

突出显示原始代码中的unsigned int 在哪些位置必须替换为unsigned long long int。我使用了 repeat1024 宏（来自原始代码），它只复制了 j=*(ptrsize_type **)j; 行 1024 次。

可以在measure_global_latency() 中调整步幅。在输出中，步幅以字节为单位。

我将不同步幅的延迟解释留给您。需要调整步幅，以免重复使用缓存！

#include <stdio.h> 
#include <stdint.h>

#include "repeat.h"

typedef unsigned long long int ptrsize_type;

__global__ void global_latency (ptrsize_type** my_array, int array_length, int iterations, unsigned long long * duration) {

    unsigned long long int start_time, end_time;
    ptrsize_type *j = (ptrsize_type*)my_array;
    volatile unsigned long long int sum_time;

    sum_time = 0;

    for (int k = 0; k < iterations; k++)
    {

        start_time = clock64();
        repeat1024(j=*(ptrsize_type **)j;)
        end_time = clock64();

        sum_time += (end_time - start_time);
    }

    ((ptrsize_type*)my_array)[array_length] = (ptrsize_type)j;
    ((ptrsize_type*)my_array)[array_length+1] = (ptrsize_type) sum_time;
    duration[0] = sum_time;
}

void parametric_measure_global(int N, int iterations, unsigned long long int maxMem, int stride)
{
    unsigned long long int maxMemToArraySize = maxMem / sizeof( ptrsize_type );
    unsigned long long int maxArraySizeNeeded = 1024*iterations*stride;
    unsigned long long int maxArraySize = (maxMemToArraySize<maxArraySizeNeeded)?(maxMemToArraySize):(maxArraySizeNeeded);

    ptrsize_type* h_a = new ptrsize_type[maxArraySize+2];
    ptrsize_type** d_a;
    cudaMalloc ((void **) &d_a, (maxArraySize+2)*sizeof(ptrsize_type));

    unsigned long long int* duration;
    cudaMalloc ((void **) &duration, sizeof(unsigned long long int));

    for ( int i = 0; true; i += stride)
    {
        ptrsize_type nextAddr = ((ptrsize_type)d_a)+(i+stride)*sizeof(ptrsize_type);
        if( i+stride < maxArraySize )
        {
            h_a[i] = nextAddr;
        }
        else
        {
            h_a[i] = (ptrsize_type)d_a; // point back to the first entry
            break;
        }
    }
    cudaMemcpy((void *)d_a, h_a, (maxArraySize+2)*sizeof(ptrsize_type), cudaMemcpyHostToDevice);

    unsigned long long int latency_sum = 0;
    int repeat = 1;
    for (int l=0; l <repeat; l++)
    {
        global_latency<<<1,1>>>(d_a, maxArraySize, iterations, duration);
        cudaThreadSynchronize ();

        cudaError_t error_id = cudaGetLastError();
        if (error_id != cudaSuccess)
        {
            printf("Error is %s\n", cudaGetErrorString(error_id));
        }

        unsigned long long int latency;
        cudaMemcpy( &latency, duration, sizeof(unsigned long long int), cudaMemcpyDeviceToHost);
        latency_sum += latency;
    }

    cudaFree(d_a);
    cudaFree(duration);

    delete[] h_a;
    printf("%f\n", (double)(latency_sum/(repeat*1024.0*iterations)) );
}

void measure_global_latency()
{
    int maxMem = 1024*1024*1024; // 1GB
    int N = 1024;
    int iterations = 1;

    for (int stride = 1; stride <= 1024; stride+=1)
    {
        printf ("  %5d, ", stride*sizeof( ptrsize_type ));
        parametric_measure_global( N, iterations, maxMem, stride );
    }
    for (int stride = 1024; stride <= 1024*1024; stride+=1024)
    {
        printf ("  %5d, ", stride*sizeof( ptrsize_type ));
        parametric_measure_global( N, iterations, maxMem, stride );
    }
}

int main()
{
    measure_global_latency();
    return 0;
}

编辑：

关于 cmets 的更多细节：我没有包括对结果的解释，因为我不认为自己是此类基准测试方面的专家。 我无意将解释作为读者的练习。

现在这是我的解释：对于 Kepler GPU（L1 不可用/禁用），我得到了相同的结果。 L2 读取低于 200 个周期是您迈出一小步就能得到的结果。通过增加iterations 变量以绝对重用L2，可以提高准确性。

现在的棘手任务是找到不重用 L2 缓存的步幅。在我的方法中，我只是盲目地尝试许多不同的（大）步幅，并希望 L2 不会被重用。在那里，我也得到了大约 500 个周期的东西。当然，更好的方法是更多地考虑缓存的结构，并通过推理而不是反复试验来推断出正确的步幅。这就是我不想自己解释结果的主要原因。

为什么步幅 > 1MB 时延迟再次降低？ 出现这种现象的原因是我使用了 1GB 的固定大小来实现最大内存使用量。使用 1024 个指针查找 (repeat1024)，1MB 的步幅正好适合内存。较大的步幅将环绕并再次使用 L2 缓存中的数据。当前代码的主要问题是 1024 指针（1024*64 位）仍然完全适合 L2 缓存。 这引入了另一个陷阱：如果您将iterations 的数量设置为> 1 并超过1024*iterations*stride*sizeof(ptrsize_type) 的内存限制，您将再次使用L2 缓存。

可能的解决方案：

• 不应将最后一个条目包装到第一个元素，而应该对位于缓存行大小和步幅之间的（未使用的！）位置实施更智能的包装。但是您需要非常小心，不要覆盖内存位置，尤其是在您多次回绕的情况下。