que - ¿Cómo escribir un punto de referencia de persecución de puntero usando punteros de 64 bits en CUDA?

nvidia cuda

Question

Este documento de investigación ejecuta una serie de varias microbenchmarks CUDA en una GPU para obtener estadísticas como latencia de memoria global, rendimiento de la instrucción, etc. Este enlace es el enlace al conjunto de microbenchmarks que los autores escribieron y ejecutaron en su GPU.

Una de las microbenchmarks llamada global.cu proporciona el código para un punto de referencia de persecución de puntero para medir la latencia de memoria global.

Este es el código del kernel que se ejecuta.

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

La línea de código que realiza la búsqueda del puntero en el caso de los punteros de 32 bits es:

j = *(unsigned int**)j;

Esta es la línea clave, porque las líneas de código restantes solo se usan para medir el tiempo.

Traté de ejecutar esto en mi GPU, pero me enfrenté a un problema. Ejecutar el mismo microbenchmark sin cambios me da un error en el tiempo de ejecución. Se An illegal memory access was encountered .

En el mismo enlace explican que:

Las pruebas de memoria global utilizan un código de búsqueda de puntero en el que los valores del puntero se almacenan en una matriz. Los punteros en GT200 son 32 bits. La prueba de memoria global deberá cambiarse si cambia el tamaño del puntero, por ejemplo, los punteros de 64 bits en Fermi.

Resulta que mi GPU es de arquitectura Kepler, que tiene punteros de 64 bits.

¿Cómo modifico ese bit de código de persecución de puntero que originalmente trata con punteros de 32 bits, para medir la latencia de memoria global utilizando punteros de 64 bits?

Editar :

De la respuesta de havogt : Una parte importante de la información que debería haber incluido en la pregunta es esta parte del código, donde se crea una matriz de ubicaciones de memoria donde cada entrada apunta a la entrada para el siguiente puntero.

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

Introducción

Antes de explicar lo que debe hacer para que el código funcione, permítame enfatizar lo siguiente: Debe comprender muy bien el hardware que está probando y el diseño de su microbenchmark. ¿Por qué es importante? El código original fue diseñado para la GT200 que no tenía una memoria caché para cargas de memoria globales ordinarias . Si ahora solo arregla el problema del puntero, medirá básicamente la latencia L2 (en Kepler, donde por defecto no se usa L1) porque el código original usa una memoria muy pequeña que encaja perfectamente en el caché.

Descargo de responsabilidad: para mí también es la primera vez que estudio un código de evaluación comparativa. Por lo tanto, verifique cuidadosamente antes de usar el siguiente código. No garantizo que no cometí errores al transformar el código original.

La solución simple (mide básicamente la latencia del caché)

Primero, no incluyó todas las partes relevantes del código en su pregunta. La parte más importante es

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

donde se construye una matriz de ubicaciones de memoria donde cada entrada apunta a la entrada para el siguiente puntero. Ahora todo lo que tiene que hacer es reemplazar todo unsigned int (que se usa para almacenar los punteros de 32 bits) por unsigned long long int , tanto en el código de configuración como en el kernel.

No publicaré el código ya que no puedo recomendar ejecutar dicho código si no lo entiendes, ver Introducción . Si lo entiendes, entonces es simple.

Mi solución

Básicamente, lo que hice fue usar tanta memoria como fuera necesaria para evaluar todos los punteros o una cantidad máxima de memoria de 1GB. En ambos casos envolví la última entrada a la primera entrada. Tenga en cuenta que, dependiendo de la zancada, es posible que muchas entradas de matriz no se inicialicen (porque nunca se utilizan).

El siguiente código es básicamente el código original después de un poco de limpieza (pero aún no está muy limpio, lo siento ...) y el cambio en la memoria. Introduje un typedef

typedef unsigned long long int ptrsize_type;

para resaltar en qué ubicaciones el unsigned int del código original tiene que ser reemplazado con unsigned long long int . repeat1024 macro repeat1024 (del código original) que simplemente copia la línea j=*(ptrsize_type **)j; 1024 veces.

Los avances se pueden ajustar en measure_global_latency() . En la salida, la zancada se da en bytes.

Dejo la interpretación de la latencia para los diferentes pasos hacia ti. ¡Hay que ajustar los pasos para que no reutilices el caché!

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

Editar:

Algunos detalles más de los comentarios: no incluí la interpretación del resultado porque no me considero un experto en dichos puntos de referencia. No era mi intención hacer de la interpretación un ejercicio para el lector.

Ahora aquí está mi interpretación: obtengo los mismos resultados para las GPU Kepler (con L1 no disponible / deshabilitado). Algo por debajo de 200 ciclos para una lectura en L2 es lo que obtienes con un pequeño zancada. La precisión se puede mejorar aumentando la variable de iterations para reutilizar definitivamente L2.

La tarea difícil ahora es encontrar un paso que no reutilice el caché L2. En mi enfoque, simplemente intento a ciegas muchos pasos grandes (grandes) y espero que L2 no se reutilice. Allí, también obtengo algo alrededor de ~ 500 ciclos. Por supuesto, el mejor enfoque sería pensar más sobre la estructura del caché y deducir el paso correcto mediante el razonamiento y no por ensayo y error. Esa es la razón principal por la que no quería interpretar el resultado yo mismo.

¿Por qué la latencia está disminuyendo de nuevo para zancadas> 1 MB? La razón de este comportamiento es que utilicé un tamaño fijo de 1GB para el uso máximo de memoria. Con las búsquedas de 1024 punteros ( repeat1024 ), una zancada de 1MB solo cabe en la memoria. Avances más grandes envolverán y volverán a utilizar los datos de la memoria caché L2. El problema principal con el código actual es que el puntero 1024 (1024 * 64 bit) todavía se ajusta perfectamente en la caché L2. Esto introduce otra trampa : si establece el número de iterations en algo> 1 y excede el límite de memoria con 1024*iterations*stride*sizeof(ptrsize_type) volverá a utilizar la caché L2.

Solución posible:

• En lugar de envolver la última entrada al primer elemento, uno debe implementar una envoltura más inteligente en una ubicación (¡sin usar!) Que se encuentra entre el tamaño de la línea de caché y la zancada. Pero debe tener mucho cuidado de no sobreescribir las ubicaciones de memoria, especialmente si está envolviendo varias veces.