__ldg hace más lento el tiempo de ejecución en una situación determinada

Question

He publicado este tema ya ayer, pero no era bien recibido, a pesar de que he sólido repro ahora, por favor tengan paciencia conmigo.Aquí están las especificaciones del sistema:

• Tesla K20m con 331.67 conductor,
• CUDA 6.0,
• Máquina Linux.

Ahora tengo una memoria global de lectura pesada de la aplicación, por tanto, traté de optimizar el uso de __ldg la instrucción en cada lugar donde me estoy leyendo la memoria global.Sin embargo, __ldg no mejorar el rendimiento en todo, el tiempo de funcionamiento de la disminución de aproximadamente 4x.Así que mi pregunta es, ¿cómo viene que la sustitución de glob_mem[index] con __ldg(glob_mem + index) pueden resultar en una disminución del rendimiento?Aquí es una versión primitiva de mi problema para reproducir:

HACER

CPP=g++
CPPFLAGS=-Wall -O4 -std=c++0x -lcudart -lcurand
LIBDIRS=/usr/local/cuda/lib64
NVCC=nvcc
NVCCINCLUDE=/usr/local/cuda/include
NVCC_COMPILER_FLAGS=-Iinclude/ -O4 -arch compute_35 -code sm_35 -c
TARGET=example

.PHONY: all clear clean purge

all: $(TARGET)

$(TARGET): kernel.o main.cpp
    @echo Linking executable "$(TARGET)" ...
    @$(CPP) $(CPPFLAGS) $(addprefix -I,$(NVCCINCLUDE)) $(addprefix -L,$(LIBDIRS)) -o $@ $^

kernel.o: kernel.cu
    @echo Compiling "$@" ...
    $(NVCC) $(addprefix -I,$(NVCCINCLUDE)) $(NVCC_COMPILER_FLAGS) $< -o $@

clean: clear

clear:
    @echo Removing object files ...
    -@rm -f *.o

purge: clear
    @echo Removing executable ...
    -@rm -f $(TARGET)

main.cpp

#include <chrono>
#include <cstdio>

#include "kernel.cuh"

using namespace std;

int main()
{
    auto start = chrono::high_resolution_clock::now();
    double result = GetResult();
    auto elapsed = chrono::high_resolution_clock::now() - start;

    printf("%.3f, elapsed time: %.3f \n", result, (double)chrono::duration_cast<std::chrono::microseconds>(elapsed).count());
    return 0;
}

kernel.cuh

#ifndef kernel_cuh
#define kernel_cuh

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

double GetResult();

#endif

kernel.cu

#include "kernel.cuh"

class DeviceClass
{
    double* d_a;
public:
    __device__ DeviceClass(double* a)
        : d_a(a) {}

    __device__ void foo(double* b, const int count)
    {
        int tid = threadIdx.x + (blockDim.x * blockIdx.x);
        double result = 0.0;
        for (int i = 0; i < count; ++i)
        {
            result += d_a[i];
            //result += __ldg(d_a + i);
        }

        b[tid] = result;
    }
};

__global__ void naive_kernel(double* c, const int count, DeviceClass** deviceClass)
{
    (*deviceClass)->foo(c, count);
}

__global__ void create_device_class(double* a, DeviceClass** deviceClass)
{
    (*deviceClass) = new DeviceClass(a);
}

double GetResult()
{
    const int aSize = 8388608;
    const int gridSize = 8;
    const int blockSize = 1024;

    double* h_a = new double[aSize];
    for (int i = 0; i <aSize; ++i)
    {
        h_a[i] = aSize - i;
    }

    double* d_a;
    cudaMalloc((void**)&d_a, aSize * sizeof(double));
    cudaMemcpy(d_a, h_a, aSize * sizeof(double), cudaMemcpyHostToDevice);

    double* d_b;
    cudaMalloc((void**)&d_b, gridSize * blockSize * sizeof(double));

    DeviceClass** d_devicesClasses;
    cudaMalloc(&d_devicesClasses, sizeof(DeviceClass**));
    create_device_class<<<1,1>>>(d_a, d_devicesClasses);

    naive_kernel<<<gridSize, blockSize>>>(d_b, aSize, d_devicesClasses);
    cudaDeviceSynchronize();

    double h_b;
    cudaMemcpy(&h_b, d_b, sizeof(double), cudaMemcpyDeviceToHost);

    cudaFree(d_a);
    cudaFree(d_b);
    return h_b;
}

Entonces, ¿qué es todo acerca de...En mi aplicación tengo algunos datos globales apuntado por la variable miembro de la clase DeviceClass que se crea en el dispositivo, exactamente como nuevo/eliminar CUDA demostración se muestra.

• Construir esta usando realizar y, a continuación, en ejecutar ./ejemplo,
• Ejecutar este ejemplo, como es la producción:"35184376283136.000, tiempo transcurrido:2054676.000".
• Después de que me quite el comentario de la línea 17 en el kernel.cu y un comentario fuera de línea a la derecha arriba, el resultado se convierte en:"35184376283136.000, tiempo transcurrido:3288975.000"
• así que el uso de __ldg el rendimiento disminuye muy significativamente a pesar de que yo estaba usando hasta ahora sin ningún tipo de problemas en diferentes ocasiones.Lo que podría ser la causa?

Answer 1

La razón para la versión utilizando __ldg siendo más lento es el hecho de que el NVCC compilador no es capaz de realizar el bucle desenrollado optimizaciones correctamente en este escenario en particular.El asunto fue sometido a NVIDIA con ID 1605303.La más reciente respuesta de NVIDIA equipo es como sigue:

Aunque, no nos hemos comunicado esto, hemos hecho antes de las investigaciones de su problema.La solución a tu problema es la mejora de nuestro bucle desenrollado heurística en el back-end del compilador el compilador integrado dentro de ptxas.Se evaluó la posibilidad de resolver este problema en CUDA 8.0, pero una solución inicial para resolver el problema causado inaceptable regresiones.Debido a otras restricciones, no fuimos capaces de desarrollar una solución adecuada en el momento para que sea hecho por CUDA 8.0.

Estamos trabajando activamente para resolver su problema en el futuro CUDA de liberación(que, a raíz de CUDA 8.0).Nos aseguraremos de que le mantendremos informado de nuestros avances en movimiento hacia adelante.