1 / 12

ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 31, 2011 MemCoalescing

Memory Coalescing. These notes will demonstrate the effects of memory coalescing. ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 31, 2011 MemCoalescing.ppt.

neka
Download Presentation

ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 31, 2011 MemCoalescing

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript


  1. Memory Coalescing These notes will demonstrate the effects of memory coalescing ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, Jan 31, 2011 MemCoalescing.ppt

  2. Memory coalescing is combining separate memory accesses into one combined access – it is done by the GPU when the locations are sequential locations in global memory banks. Consider setting the elements of two-dimensional array to given data values. This could be done across rows or down columns In the following code, we will demonstrate the effects of each approach

  3. Approach Load numbers into a two-dimensional array Flattened global threadID of thread is loaded into array element being accessed so one can tell which thread accesses which location when one prints out array Access is done across rows and also across column and the time of execution compared. In practice, a problem may dictate the access order GPU structure -- one or more 2-D blocks in a 2-D grid. Each block is 2-D 32x32 threads fixed (max. for compute cap. 2.x)

  4. One way __global__ void gpu_Comput1 (int *h, int N, int T) { int col = threadIdx.x + blockDim.x * blockIdx.x; int row = threadIdx.y + blockDim.y * blockIdx.y; int threadID = col + row * N; // thread ID int index = col + row * N; // array index for (int t = 0; t < T; t++) // loop to reduce other time effects h[index] = threadID; // load array with global thread ID }

  5. Another way __global__ void gpu_Comput2 (int *h, int N, int T) { int col = threadIdx.x + blockDim.x * blockIdx.x; int row = threadIdx.y + blockDim.y * blockIdx.y; int threadID = col + row * N; // thread ID int index = row + col * N; // array index for (int t = 0; t < T; t++) // loop to reduce other time effects h[index] = threadID; // load array with global thread ID }

  6. /* ------------------------- GPU Computation 1 -----------------------------------*/ gpu_Comput1<<< Grid, Block >>>(dev_h, N, T); // launch once kernel outside timing cudaEventRecord( start, 0 ); gpu_Comput1<<< Grid, Block >>>(dev_h, N, T); cudaEventRecord( stop, 0 ); // measure end time cudaEventSynchronize( stop ); // wait for event recording cudaEventElapsedTime( &elapsed_time_ms1, start, stop ); cudaMemcpy(h,dev_h, size ,cudaMemcpyDeviceToHost); //Results to check printf("\nComputation with memory coalescing possible\n"); printArray(h,N); printf("\nTime to calculate results on GPU: %f ms.\n", elapsed_time_ms1); Computation 2 similar

  7. Some results A grid of one block and one iteration Array 32x32 No speedup recorded because time of other operations dominate execution time

  8. A grid of one block and 1000000 iterations Array 32 x 32 Speedup = 17.16

  9. Repeat just to check results are consistent

  10. A grid of 16 x 16 blocks and 10000 iterations Array 512x512 Speedup = 12.08 Different numbers of iterations produce similar results

  11. Different Array Sizes 1000 iterations. Block size 32 x 32. Number of blocks to suit array size

  12. Questions

More Related