在某些情况下，__ldg会导致较慢的执行时间

Question

我昨天已经发了这篇文章，但没有受到很好的欢迎，尽管我现在有了可靠的复制品，请耐心等待。以下是系统规范：

• 特斯拉K20m有331.67司机，
• 库达6.0，
• Linux机器。

现在，我有了一个全局内存，读取了大量的应用程序，因此我尝试在我正在读取全局内存的每个地方使用__ldg指令来优化它。然而，__ldg根本没有提高性能，运行时间减少了大约4倍。因此，我的问题是，为什么用glob_mem[index]替换__ldg(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)

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

那到底是怎么回事..。在我的应用程序中，DeviceClass类的成员变量指向了一些全局数据，这个类是在设备上创建的，就像新/删除CUDA演示所显示的那样。

• 使用make构建这个程序，然后执行。
• 按照原样运行此示例："35184376283136.000，运行时间: 2054676.000“。
• 在我取消对kernel.cu中的第17行进行注释并在其上方注释行之后，结果是："35184376283136.000，运行时间: 3288975.000“。
• 因此，使用__ldg会显着地降低性能，即使到目前为止我还在使用它，在不同的情况下没有任何问题。可能是什么原因？

Answer 1

使用__ldg的版本之所以速度较慢，是因为NVCC编译器无法在此特定场景中正确执行循环展开优化。该问题已提交给NVIDIA，ID为1605303。NVIDIA小组最近的答复如下：

虽然，我们还没有通知您，我们已经对您的问题做了事先的调查。解决您的问题的方法是改进我们的循环--在后端编译器中展开启发式--嵌入在ptxas中的编译器。我们在CUDA 8.0中评估了解决这一问题的可能性，但是解决问题的初始解决方案会导致不可接受的倒退。由于其他限制因素，我们未能及时制定出适当的解决办法，以便在CUDA 8.0中完成。 我们正在积极努力解决您的问题，在未来的CUDA发布(以下的CUDA 8.0)。我们将确保随时向你通报我们的进展情况。