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

Question

本研究论文在GPU上运行了一系列CUDA微基准测试,以获取全局内存延迟,指令吞吐量等统计信息.此链接是作者在其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.

在同一个链接中,他们解释说:

全局内存测试使用指针追踪代码,其中指针值存储在数组中.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 设计的，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中将 all （用于存储 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;
}

编辑：

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

现在这是我的解释：对于 Kepler GPU，我得到了相同的结果（L1 不可用/禁用）。L2 读取的低于 200 个周期的值是小步幅获得的结果。可以通过增加变量来明确重用L2来提高精度iterations。

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

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

可能的解决方案：

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