__ldg 在某些情况下会导致执行时间变慢

Question

我昨天已经发布了这个问题，但没有受到好评，虽然我现在有可靠的repro，请多多包涵。以下是系统规格：

• 配备 331.67 驱动程序的特斯拉 K20m，
• CUDA 6.0，
• Linux 机器。

现在我有一个全局内存读取繁重的应用程序，因此我尝试在我正在读取全局内存的每个地方使用__ldg 指令对其进行优化。但是，__ldg 根本没有提高性能，运行时间减少了大约 4 倍。所以我的问题是，用__ldg(glob_mem + index) 替换glob_mem[index] 怎么可能导致性能下降？这是我的问题的原始版本供您重现：

制作

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)

ma​​in.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;
}

那么这到底是怎么一回事...在我的应用程序中，我有一些全局数据由在设备上创建的类 DeviceClass 的成员变量指向，正如 new/delete CUDA 演示所示。

• 使用 make 构建它，然后执行 ./example，
• 按原样运行此示例会产生：“35184376283136.000, elapsed time: 2054676.000”。
• 在 kernel.cu 中取消注释第 17 行并将其正上方的行注释掉后，结果变为：“35184376283136.000, elapsed time: 3288975.000”
• 所以使用 __ldg 会显着降低性能，即使我在不​​同场合使用它直到现在都没有任何问题。可能是什么原因？

Answer 1

使用__ldg 的版本变慢的原因是NVCC 编译器在这种特定情况下无法正确执行循环展开优化。该问题已提交给 NVIDIA，ID 为 1605303。NVIDIA 团队的最新回复如下：

尽管我们尚未将此事告知您，但我们已经对您的问题进行了事先调查。您的问题的解决方案是改进我们在后端编译器中的循环展开启发式 - 嵌入在 ptxas 中的编译器。我们评估了在 CUDA 8.0 中解决此问题的可能性，但解决您的问题的初始解决方案导致了不可接受的回归。由于其他限制，我们无法及时开发出适合 CUDA 8.0 的解决方案。

我们正在积极努力在未来的 CUDA 版本（CUDA 8.0 的后续版本）中解决您的问题。我们将确保让您随时了解我们的进展。