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

Posted 2023-03-23

【中文标题】__ldg 在某些情况下会导致执行时间变慢

【英文标题】：__ldg causes slower execution time in certain situation

【发布时间】：2014-09-21 08:51:30

【问题描述】：

我昨天已经发布了这个问题，但没有受到好评，虽然我现在有可靠的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 会显着降低性能，即使我在不​​同场合使用它直到现在都没有任何问题。可能是什么原因？

【问题讨论】：

在您的 foo 函数中，循环展开有所不同，原因我还不完全清楚。要消除这种对结果的影响，请在该函数的 for 循环之前添加 #pragma unroll 8，然后重新运行您的测试。

据我所知，__ldg 只是使用纹理缓存的一种更方便的方式，它针对半随机访问进行了调整。您的特定内核具有纯线性读取，由于合并已经很快。额外的缓存在这里没有帮助，并且可能会增加延迟。此外，检查 PTX 的两种情况表明编译器为主循环生成了完全不同的代码。

仅供参考，我最初的解决方案是将 DeviceClass 的成员变量 d_a 作为静态大小的数组。我也有 host 构造函数，使用它在主机上创建对象，然后使用 cudaMemcpy 在设备上创建它的副本，该副本与静态大小的数组很好地配合使用。使用这种方法产生了很好的结果，但显然我不想使用静态大小的数组，因为大小取决于驱动我计算的外部配置文件。

我相信编译器无法确定循环的预期行程计数，因此如果您通过count = 3，避免意外是相当保守的，尽管我不能很好地解释了两种情况之间展开策略的差异。然而，当我指示展开大于 3 的级别（我尝试过 4、8 和 64）时，我看到了预期的结果，即 __ldg 代码更快。

L1 缓存、共享内存或常量内存都更适合这种工作负载，因为所有线程都在同时访问 d_a[i] 的相同元素。要使用 L1 缓存，请阻止您的内部循环以最大限度地重用。共享内存和常量内存需要按照 N 体实现的方式进行更多重构。 github.com/ArchaeaSoftware/cudahandbook/tree/master/nbody

【参考方案1】：

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

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

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

【讨论】：

在 CUDA 9 中修复？

不知道，自从我完成论文后就没有使用过 CUDA，请随意尝试一下 :)