__ldg()内部函数和普通执行有什么区别？

Question

__ldg()内部函数和普通执行有什么区别？

8

我正在尝试探索'__ldg intrinsic'。我已经阅读了NVIDIA关于此的文档，但是没有得到任何令人满意的关于它的使用和实现方面的答案。此外，参考此链接，我尝试在一个简单的1024*1024矩阵乘法示例中实现了__ldg。

#include<stdio.h>
#include<stdlib.h>

__global__ void matrix_mul(float * ad,float * bd,float * cd,int N)
{
        float pvalue=0;
        //find Row and Column corresponding to a data element for each thread
        int Row = blockIdx.y * blockDim.y + threadIdx.y;
        int Col = blockIdx.x * blockDim.x + threadIdx.x;
        //calculate dot product of Row of First Matrix and Column of Second Matrix
        for(int i=0;i< N;++i)
        {
//   I tried with executing this first:
            float m=__ldg(&ad[Row * N+i]);
            float n=__ldg(&bd[i * N + Col]);

//Then I executed this as a normal execution:
//          float m = ad[Row * N+i];
//          float n = bd[i * N + Col];

            pvalue += m * n;
         }
        //store dot product at corresponding position in resultant Matrix
        cd[Row * N + Col] = pvalue;
}

int main()
{
    int N = 1024,i,j;               //N == size of square matrix

    float *a,*b;
    float *ad,*bd,*cd,*c;

    //open a file for outputting the result
    FILE *f;
    f=fopen("Parallel Multiply_ldg.txt","w");

    size_t size=sizeof(float)* N * N;

    //allocate host side memory
    a=(float*)malloc(size);
    b=(float*)malloc(size);
    c=(float*)malloc(size);

    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
        {
            a[i*N+j]=2.0;   //(float)(i*N+j);       //initializing each value with its own index
            b[i*N+j]=1.0;   //(float)(i*N+j);       //random functions can be used alternatively
        }
    }

    //allocate device memory
    cudaMalloc(&ad,size);
    //printf("\nAfter cudaMalloc for ad\n%s\n",cudaGetErrorString(cudaGetLastError()));
    cudaMalloc(&bd,size);
    //printf("\nAfter cudaMalloc bd\n%s\n",cudaGetErrorString(cudaGetLastError()));
    cudaMalloc(&cd,size);
    //printf("\nAfter cudaMalloc cd\n%s\n",cudaGetErrorString(cudaGetLastError()));

    //copy value from host to device
    cudaMemcpy(ad,a,size,cudaMemcpyHostToDevice);
    cudaMemcpy(bd,b,size,cudaMemcpyHostToDevice);

    printf("\nAfter HostToDevice Memcpy\n%s\n",cudaGetErrorString(cudaGetLastError()));

    //calculate execution configuration
    dim3 blocksize(16,16);              //each block contains 16 * 16 (=256) threads
    dim3 gridsize(N/16,N/16);           //creating just sufficient no of blocks

    //GPU timer code
    float time;
    cudaEvent_t start,stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start,0);

    matrix_mul <<< gridsize, blocksize >>> (ad,bd,cd, N);
    cudaDeviceSynchronize();
    cudaEventRecord(stop,0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time,start,stop);         //time taken in kernel call calculated
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    //copy back results
    cudaMemcpy(c,cd,sizeof(float)* N*N,cudaMemcpyDeviceToHost);

    printf("\nAfter DeviceToHost Memcpy\n%s\n",cudaGetErrorString(cudaGetLastError()));

    //output results in output_file
    fprintf(f,"Array A was---\n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            fprintf(f,"%f ",a[i*N+j]);
        fprintf(f,"\n");
    }
    fprintf(f,"\nArray B was---\n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            fprintf(f,"%f ",b[i*N+j]);
        fprintf(f,"\n");
    }
    fprintf(f,"\nMultiplication of A and B gives C----\n");
    for(i=0;i<N;i++)
    {
        for(j=0;j<N;j++)
            fprintf(f,"%f ",c[i*N+j]);              //if correctly computed, then all values must be N
        fprintf(f,"\n");
    }
    printf("\nYou can see output in Parallel Mutiply.txt file in project directory");
    printf("\n\nTime taken is %f (ms)\n",time);
    fprintf(f,"\n\nTime taken is %f (ms)\n",time);
    fclose(f);

    cudaThreadExit();
    //cudaFree(ad); cudaFree(bd); cudaFree (cd);
    free(a);free(b);free(c);
    //_getch();
    return 1;
}

我在我的内核中评论了__ldg部分并通过正常执行运行，反之亦然。

在这两种情况下，它都给出了正确的乘法结果。我对这些执行之间的时间差异感到困惑，因为它的差别很大，几乎超过100倍！

使用__ldg时，它会给出： 所需时间为0.014432（毫秒）

而在没有__ldg的正常执行情况下，它会给出： 所需时间为36.858398（毫秒）

这是使用__ldg内置函数的确切方法吗？ __ldg内置函数的意义是什么？如何正确地使用它？显然，我在上面的代码中所做的是错误和幼稚的。我正在寻找解释和示例。提前致谢。

- sandeep.ganage

1

只是为了我理解一下 - 结果是正确的，你获得了很大的加速，你的问题是“我做错了什么？”？... - talonmies

1

@talonmies 速度如此之快，以至于我怀疑是否正确。如果是正确的，那么这个__ldg做了什么魔法，使我得到了加速？如果这是不正确的，那么正确使用__ldg的方法是什么？总体而言，我正在寻找关于这个__ldg概念及其实现的更多解释。 - sandeep.ganage

@AviGinsburg 不是的！这两个都是GPU结果。一个使用了__ldg，另一个没有使用__ldg。 - sandeep.ganage

1

@Cicada，这将导致向编译器建议缓存ad和bd，而不是显式的ldg。 - Avi Ginsburg

2

结果不正确。当您在您提到的任何一个GPU上运行此代码时，它将打印出“无效设备函数”，但您似乎忽略了这一点。而且，您声称结果是正确的也不是真的。可能您感到困惑，并且在重新编译代码时没有更改文件名。__ldg需要计算能力3.5。 - Robert Crovella

显示剩余9条评论

2个回答

2

在NVIDIA GPU中，有一种被称作纹理的图像，这些图像具有特殊的逻辑，可用于处理图像。

这种纹理内存是GPU中另一种可用的内存类型。尤其是常量、全局和寄存器文件内存与此纹理内存没有任何关系。

Kepler GPU及更高版本增加了使用“GPU纹理管道”中的该内存的能力。

但让我们明确常量缓存和只读缓存之间的区别。

常量缓存

通过常量缓存加载的数据必须相对较小，并且必须以所有warp线程在任何给定时间都访问同一位置的方式进行访问。

只读缓存或纹理内存缓存

缓存可以更大，并且可以以非均匀模式进行访问。只读缓存的粒度为32字节。

您可以将其用作CUDA内核的“只读缓存”。

1. Data stored in global memory can be cached in that place GPU Texture Memory
2. With doing that you give promise to the compiler that data is read-only for the 
   duration of a kernel execution in GPU.

有两种方法可以实现这个。 A. 使用内置函数 __ldg。

Example: output[i] += __ldg(&input[j]);

B. Qualifying pointers to global memory

const float* __restrict__ input
output[idx] += input[idx];

比较：

由于编译器的原因，内在的__ldg是更好的选择。

- Konstantin Burlachenko

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Avi Ginsburg · Accepted Answer

来自CUDA C编程指南

计算能力为3.x的设备的全局内存访问会缓存在L2中，而计算能力为3.5的设备还可以在只读数据缓存中进行缓存（详见前面的章节）；它们不会被缓存到L1中。

...

对于整个内核生命周期都是只读的数据，也可以使用__ldg()函数（请参阅只读数据缓存加载函数）将其缓存在只读数据缓存中。当编译器检测到某些数据满足只读条件时，它将使用__ldg()来读取它。编译器可能无法始终检测到某些数据的只读条件。将用于加载这些数据的指针标记为const和__restrict__限定符可以增加编译器检测到只读条件的可能性。

只读缓存访问的延迟比全局内存访问要低得多。由于矩阵乘法多次从内存中访问相同的值，缓存在只读缓存中可以大大提高速度（在内存限制应用程序中）。