如何在数据自动化系统中使用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位。如果指针大小发生变化，例如费米上的64位指针，则需要更改全局内存测试。

原来我的GPU是开普勒结构的，它有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延迟(在开普勒，默认情况下不使用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 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;
}

编辑：

关于评论的更多细节:我没有包括对结果的解释，因为我不认为自己是这类基准的专家。我不打算把这一解释作为读者的练习。

下面是我的解释:对于开普勒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缓存。

可能的解决办法：

• 与其将最后一个条目包装到第一个元素，还不如实现一个更聪明的包装到一个(未使用的！)位于缓存行大小和步长之间的位置。但是，您需要非常小心，不要覆盖内存位置，特别是在多次包装的情况下。