title: What is the difference between __ldg() intrinsic and a normal execution?
body: <p>I am trying to explore '__ldg intrinsic'. I have gone through NVIDIA's documentation for this but didn't get any satisfactory answer over its use and implementations. Moreover with reference to <a href="https://stackoverflow.com/questions/24069524/error-using-ldg-in-cuda-kernel-at-compile-time/24073775#24073775">THIS</a> I tried implementing __ldg in a simple 1024*1024 matrix multiplication example.</p>
  
  <pre><code>#include&lt;stdio.h&gt;
  #include&lt;stdlib.h&gt;
  
  __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&lt; N;++i)
          {
  //   I tried with executing this first:
              float m=__ldg(&amp;ad[Row * N+i]);
              float n=__ldg(&amp;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&lt;N;i++)
      {
          for(j=0;j&lt;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(&amp;ad,size);
      //printf("\nAfter cudaMalloc for ad\n%s\n",cudaGetErrorString(cudaGetLastError()));
      cudaMalloc(&amp;bd,size);
      //printf("\nAfter cudaMalloc bd\n%s\n",cudaGetErrorString(cudaGetLastError()));
      cudaMalloc(&amp;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(&amp;start);
      cudaEventCreate(&amp;stop);
      cudaEventRecord(start,0);
  
      matrix_mul &lt;&lt;&lt; gridsize, blocksize &gt;&gt;&gt; (ad,bd,cd, N);
      cudaDeviceSynchronize();
      cudaEventRecord(stop,0);
      cudaEventSynchronize(stop);
      cudaEventElapsedTime(&amp;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&lt;N;i++)
      {
          for(j=0;j&lt;N;j++)
              fprintf(f,"%f ",a[i*N+j]);
          fprintf(f,"\n");
      }
      fprintf(f,"\nArray B was---\n");
      for(i=0;i&lt;N;i++)
      {
          for(j=0;j&lt;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&lt;N;i++)
      {
          for(j=0;j&lt;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;
  }
  </code></pre>
  
  <p>I commented that __ldg part in my kernel and executed by normal execution, and vice versa.
  In both cases it gives me correct multiplication result. I am confused with the time difference I am getting between these executions, because its huge almost more than 100X! </p>
  
  <p>In case of __ldg it gives me: <code>Time taken is 0.014432 (ms)</code></p>
  
  <p>And in case of normal execution without __ldg it gives me : <code>Time taken is 36.858398 (ms)</code></p>
  
  <p>Is this the exact way of using __ldg intrisic? What is the significance of __ldg intrinsic and what is the proper way of using it? Apparently what I did above in my code is wrong and naive. I am looking for explanation and example. Thanks in advance.</p>