如何在 CUDA 中使用 64 位指针编写指针追踪基准测试？
How to write a pointer-chasing benchmark using 64-bit pointers in CUDA?
This research paper runs a series of several CUDA microbenchmarks on a GPU to obtain statistics like global memory latency, instruction throughput, etc. This link 是作者编写的微基准测试集的 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 他们解释说：
The global memory tests use pointer chasing code where the pointer values are stored in an array. Pointers on GT200 are 32 bits. The global memory test will need to be changed if the pointer size changes, e.g., 64-bit pointers on Fermi.
原来我的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);
}
简介
在我解释您必须做什么才能使代码正常工作之前，让我强调以下几点：您应该非常了解您正在测试的硬件和微基准测试的设计。它为什么如此重要？ 原始代码是为 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。
我不会post代码因为我不能推荐运行这样的代码如果你不理解它，请看介绍。懂了就简单了
我的解决方案
基本上我所做的是使用尽可能多的内存来评估所有指针或最大内存量为 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;
}
编辑：
评论的更多细节：我没有包括对结果的解释，因为我不认为自己是此类基准测试的专家。 我不打算将解释作为 reader 的练习。
现在这是我的解释：我对开普勒 GPU 得到了相同的结果（L1 不是 available/disabled）。 L2 读取低于 200 个周期是您迈出的一小步。可以通过增加 iterations 变量以明确重用 L2 来提高准确性。
现在棘手的任务是找到一个不重用 L2 缓存的步幅。在我的方法中，我只是盲目地尝试许多不同的（大的）步幅，并希望 L2 不被重用。在那里，我也得到了大约 500 个周期的东西。当然，更好的方法是更多地考虑缓存的结构，并通过推理而不是反复试验来推断出正确的步幅。这就是我不想自己解释结果的主要原因。
为什么步长 > 1MB 时延迟会再次降低？ 出现此行为的原因是我使用了 1GB 的固定大小来实现最大内存使用。使用 1024 次指针查找 (repeat1024)，1MB 的步幅刚好适合内存。较大的步幅将环绕并再次使用 L2 缓存中的数据。当前代码的主要问题是 1024 指针（1024*64 位）仍然完全适合二级缓存。
这引入了另一个 陷阱 ：如果您将 iterations 的数量设置为大于 1 并超过 1024*iterations*stride*sizeof(ptrsize_type) 的内存限制，您将再次使用 L2 缓存。
可能的解决方案：
- 与其将最后一个条目包装到第一个元素，不如将更智能的包装实施到缓存行大小和跨度之间的（未使用！）位置。但是您需要非常小心，不要覆盖内存位置，尤其是当您多次回绕时。
This research paper runs a series of several CUDA microbenchmarks on a GPU to obtain statistics like global memory latency, instruction throughput, etc. This link 是作者编写的微基准测试集的 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 他们解释说：
The global memory tests use pointer chasing code where the pointer values are stored in an array. Pointers on GT200 are 32 bits. The global memory test will need to be changed if the pointer size changes, e.g., 64-bit pointers on Fermi.
原来我的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);
}
简介
在我解释您必须做什么才能使代码正常工作之前，让我强调以下几点：您应该非常了解您正在测试的硬件和微基准测试的设计。它为什么如此重要？ 原始代码是为 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。
我不会post代码因为我不能推荐运行这样的代码如果你不理解它，请看介绍。懂了就简单了
我的解决方案
基本上我所做的是使用尽可能多的内存来评估所有指针或最大内存量为 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;
}
编辑：
评论的更多细节：我没有包括对结果的解释，因为我不认为自己是此类基准测试的专家。 我不打算将解释作为 reader 的练习。
现在这是我的解释：我对开普勒 GPU 得到了相同的结果（L1 不是 available/disabled）。 L2 读取低于 200 个周期是您迈出的一小步。可以通过增加 iterations 变量以明确重用 L2 来提高准确性。
现在棘手的任务是找到一个不重用 L2 缓存的步幅。在我的方法中，我只是盲目地尝试许多不同的（大的）步幅，并希望 L2 不被重用。在那里，我也得到了大约 500 个周期的东西。当然，更好的方法是更多地考虑缓存的结构，并通过推理而不是反复试验来推断出正确的步幅。这就是我不想自己解释结果的主要原因。
为什么步长 > 1MB 时延迟会再次降低？ 出现此行为的原因是我使用了 1GB 的固定大小来实现最大内存使用。使用 1024 次指针查找 (repeat1024)，1MB 的步幅刚好适合内存。较大的步幅将环绕并再次使用 L2 缓存中的数据。当前代码的主要问题是 1024 指针（1024*64 位）仍然完全适合二级缓存。
这引入了另一个 陷阱 ：如果您将 iterations 的数量设置为大于 1 并超过 1024*iterations*stride*sizeof(ptrsize_type) 的内存限制，您将再次使用 L2 缓存。
可能的解决方案：
- 与其将最后一个条目包装到第一个元素，不如将更智能的包装实施到缓存行大小和跨度之间的（未使用！）位置。但是您需要非常小心，不要覆盖内存位置，尤其是当您多次回绕时。