How to write a pointer-chasing benchmark using 64-bit pointers in CUDA?
<p><a href="http://www.stuffedcow.net/files/gpuarch-ispass2010.pdf" rel="nofollow">This research paper</a> runs a series of several CUDA microbenchmarks on a GPU to obtain statistics like global memory latency, instruction throughput, etc. <a href="http://www.stuffedcow.net/research/cudabmk" rel="nofollow">This link</a> is the link to the set of microbenchmarks that the authors wrote and ran on their GPU. </p>

<p>One of the microbenchmarks called <code>global.cu</code> gives the code for a pointer-chasing benchmark to measure global memory latency. </p>

<p>This is the code of the kernel that is run.</p>

<pre><code>__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 &lt; 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;
}
</code></pre>

<p>The line of code performing the pointer chasing in the case of 32-bit pointers is:</p>

<pre><code>j = *(unsigned int**)j;
</code></pre>

<p>This is the key line, because the remaining lines of code are only used for time measurement.</p>

<p>I tried to run this on my GPU, but I faced an issue. Running the same microbenchmark with no changes gives me a runtime error of <code>An illegal memory access was encountered</code>.</p>

<p><a href="http://www.stuffedcow.net/research/cudabmk" rel="nofollow">In the same link</a> they explain that:</p>

<blockquote>
  <p>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.</p>
</blockquote>

<p>It turns out that my GPU is of Kepler architecture, which has 64-bit pointers.</p>

<p>How do I modify that bit of pointer-chasing code which originally deals with 32-bit pointers, in order to measure global memory latency using 64-bit pointers?</p>

<p><strong>Edit</strong>:</p>

<p>From <em>havogt</em>'s answer: An important piece of information that I should have included in the question is this portion of the code, where an array of memory locations is built where each entry points to the entry for the next pointer.</p>

<pre><code>for (i = 0; i &lt; 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);
}
</code></pre>