How to write a chasing pointer using 64-bit pointers in CUDA?

Question

This research document runs a series of several CUDA micro-objects on a GPU to obtain statistics, such as global memory latency, instruction throughput, etc. This link is a link to a set of microobjects that the authors wrote and ran on their GPU.

One of the microobjects, called global.cu , provides code for a pointer test to measure global memory latency.

This is the code for the kernel being launched.

 __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; } 

A line of code chasing a pointer in the case of 32-bit pointers:

 j = *(unsigned int**)j; 

This is the key line, as the rest of the lines of code are used only for measuring time.

I tried to run this on my GPU, but I had a problem. Executing the same microobject without changes gives me a runtime error An illegal memory access was encountered .

In the same link, they explain that:

In global memory tests, the pursuer code is used, in which the pointer values ​​are stored in an array. The pointers to the GT200 are 32 bits. The global memory test will need to be changed if the pointer size changes, for example, 64-bit pointers to Fermi.

It turns out my GPU has a Kepler architecture that has 64 bit pointers.

How do I change this bit of code to check for a pointer that initially deals with 32-bit pointers in order to measure global memory latency with 64-bit pointers?

Edit

From havogt's answer: an important piece of information that I should have included in the question is this piece of code in which an array of memory cells is created, where each record points to a record for the next pointer.

 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

Introduction

Before I explain what you need to do to get the code to work, let me emphasize the following: you must have a very good understanding of the equipment under test and the design of your micro lens. Why is it important? The original code was developed for the GT200, which did not have a cache for normal loads of global memory . If you just fix the pointer problem now, you'll basically measure L2 latency (on Kepler, where L1 is not used by default), because the source code uses very little memory, which fits nicely into the cache.

Disclaimer: for me, this is also the first time we are studying such benchmarking. Therefore, before using the code below, carefully check it. I can not guarantee that I was not mistaken when converting the source code.

Simple solution (mainly measures cache latency)

First, you did not include all relevant parts of the code in your question. The most important part is

 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); } 

where an array of memory cells is created, where each record points to a record for the next pointer. Now all you have to do is replace all unsigned int (which is used to store 32-bit pointers) with unsigned long long int , both in the installation code and in the kernel.

I will not publish the code, since I cannot recommend running such code if you do not understand it, see the Introduction. If you understand this, then it is simple.

My decision

I mainly used as much memory as needed to evaluate all pointers or a maximum memory capacity of 1 GB. In both cases, I wrapped the last record in the first record. Note that depending on the step, many array entries may be uninitialized (because they are never used).

The following code is basically the source code after a little cleanup (but it is still not very clean, sorry ...) and a change in memory. I entered typedef

 typedef unsigned long long int ptrsize_type; 

to highlight where unsigned int from the source code should be replaced with unsigned long long int . I used the repeat1024 macro (from source), which simply copies the line j=*(ptrsize_type **)j; 1,024 times.

Steps can be adjusted in measure_global_latency() . At the output, the step is indicated in bytes.

I leave for you an interpretation of the delay for the different steps. The steps must be configured so that you do not reuse the cache!

 #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; } 

Edit:

A few comments on the comments: I did not include the interpretation of the result, because I do not consider myself an expert in such tests. I was not going to make an interpretation of the exercise for the reader.

Now here is my interpretation: I get the same results for the Kepler GPU (with L1 unavailable / disabled). Something below 200 cycles for reading L2 is what you get in small increments. Accuracy can be improved by increasing the variable iterations for unambiguous reuse of L2.

Now the challenge is to find a step that will not reuse the L2 cache. In my approach, I just blindly try many different (large) steps and hope that L2 will not be reused. There, I also get something around ~ 500 cycles. Of course, the best approach would be to think more about the structure of the cache and take the right step, reasoning, and not through trial and error. This is the main reason why I do not want to interpret the result myself.

Why is the latency again reduced for steps> 1 MB? . The reason for this behavior is because I used a fixed size of 1 GB for maximum memory usage. When using 1024 pointers ( repeat1024 ), the 1MB step just fits into memory. Big steps will wrap around and reuse data from the L2 cache. The main problem with the current code is that the pointer 1024 (1024 * 64 bit) still fits perfectly in the L2 cache. This adds another trap . If you set the number of iterations to something> 1 and exceed the memory limit with 1024*iterations*stride*sizeof(ptrsize_type) , you will use the L2 cache again.

Possible Solution:

• Instead of wrapping the last entry in the first element, you should implement smarter packaging in a (unused!) Location that is between the size of the cache line and the step. But you have to be very careful not to overwrite memory locations, especially if you wrap files several times.