title: cuda global and shared memory access time
body: <p>Recently I am studying about CUDA. I want to know about CUDA memory access times.</p>
  
  <p>In, CUDA Programming Guide written memory access times:</p>
  
  <ul>
  <li>Global memory access time is 400 ~ 600 Cycle</li>
  <li>Shared memory( L1 Cache ) access time is 20 ~ 40 Cycle</li>
  </ul>
  
  <p>I think that Cycle is same as clock. Is this correct ? If that is Correct, so I examined memory access times. The host is fixed but the kernel code has 3 versions. This is My Code:</p>
  
  <hr>
  
  <blockquote>
    <p>host Code</p>
  </blockquote>
  
  <pre><code>float* H1  = (float*)malloc(sizeof(float)*100000);
  float* D1;
  
  for( int i = 0 ; i &lt; 100000 ; i++ ){
      H1[i]  = i;
  }
  
  cudaMalloc( (void**)&amp;D1,   sizeof(float)*100000);
  cudaMemcpy( D1, H1,    sizeof(float)*100000, cudaMemcpyHostToDevice );
  
  
  cudaPrintfInit();
  test&lt;&lt;&lt;1,1&gt;&gt;&gt;( D1 );
  cudaPrintfDisplay(stdout, true);
  
  cudaPrintfEnd();
  </code></pre>
  
  <hr>
  
  <blockquote>
    <p>kernel version 1</p>
  </blockquote>
  
  <pre><code>float Global1;
  float Global2;
  float Global3;
  
  int Clock;
  
  Clock = clock();
  Global1 = Dev_In1[1];
  Clock = clock() - Clock;
  cuPrintf("Global Memory Access #1 : %d\n", Clock );
  
  Clock = clock();
  Global2 = Dev_In1[2];
  Clock = clock() - Clock;
  cuPrintf("Global Memory Access #2 : %d\n", Clock );
  
  Clock = clock();
  Global3 = Dev_In1[3];
  Clock = clock() - Clock;
  cuPrintf("Global Memory Access #3 : %d\n", Clock );
  </code></pre>
  
  <p>It`s result</p>
  
  <p>Global Memory Access #1 : 882<br>
  Global Memory Access #2 : 312<br>
  Global Memory Access #3 : 312       </p>
  
  <p>I think that first access not cache so took 800 Cycle
  but 2nd access 3rd access took 312 Cycle because, Dev_In[2], Dev_In[3] is cached..</p>
  
  <hr>
  
  <blockquote>
    <p>kernel version 2</p>
  </blockquote>
  
  <pre><code>int Global1, Global2, Global3;              
  int Clock;              
  
  Clock = clock();                
  Global1 = Dev_In1[1];               
  Clock = clock() - Clock;                
  cuPrintf("Global Memory Access #1 : %d\n", Clock );             
  
  Clock = clock();                
  Global2 = Dev_In1[50000];               
  Clock = clock() - Clock;                
  cuPrintf("Global Memory Access #2 : %d\n", Clock );             
  
  Clock = clock();                
  Global3 = Dev_In1[99999];               
  Clock = clock() - Clock;                
  cuPrintf("Global Memory Access #3 : %d\n", Clock );             
  </code></pre>
  
  <p>it`s result</p>
  
  <p>Global Memory Access #1 : 872<br>
  Global Memory Access #2 : 776<br>
  Global Memory Access #3 : 782       </p>
  
  <p>i think that not cached Dev_In1[50000] and Dev_In2[99999] at 1st access time</p>
  
  <p>so... #1,#2,#3 is late... </p>
  
  <hr>
  
  <blockquote>
    <p>kernel version 3</p>
  </blockquote>
  
  <pre><code>int Global1, Global2, Global3;                  
  int Clock;                  
  
  Clock = clock();                    
  Global1 = Dev_In1[1];                   
  Clock = clock() - Clock;                    
  cuPrintf("Global Memory Access #1 : %d\n", Clock );                 
  
  Clock = clock();                    
  Global1 = Dev_In1[50000];                   
  Clock = clock() - Clock;                    
  cuPrintf("Global Memory Access #2 : %d\n", Clock );                 
  
  Clock = clock();                    
  Global1 = Dev_In1[99999];                   
  Clock = clock() - Clock;                    
  cuPrintf("Global Memory Access #3 : %d\n", Clock );                 
  </code></pre>
  
  <p>result</p>
  
  <p>Global Memory Access #1 : 168<br>
  Global Memory Access #2 : 168<br>
  Global Memory Access #3 : 168       </p>
  
  <p>I don`t understand this result </p>
  
  <p>Dev_In[50000], Dev_In[99999] is not cached, but access time is very fast!!
  just, i used 1 variable.... </p>
  
  <p>SO.. My question is that gpu cycle == gpu clock ?</p>
  
  <p>and In result1, result2, result3 why memory access time is very fast in result3?</p>