80 likes | 184 Views
Explore experimental results of using constant memory in CUDA programming. Compare with global memory performance for adding two vectors. Understand limitations and benefits of constant memory. Analyze speedup achieved with constant memory usage.
E N D
Measuring Performance of Constant Memory These notes will introduce: Results of an experiment using constant memory ITCS 6/8010 CUDA Programming, UNC-Charlotte, B. Wilkinson, March 3, 2011 ConstantMemTiming.ppt
Program The test program simply adds two vectors A and B together to produce a third vector, C One version uses constant memory for A and B Another version uses regular global memory for A and B Note maximum available for constant memory on the GPU (all compute capabilities so far) is 64 Kbytes total.
Code Array declarations #define N 8192 // max size allowed for two vectors in const. mem // Constants held in constant memory __device__ __constant__ int dev_a_Cont[N]; __device__ __constant__ int dev_b_Cont[N]; // regular global memory for comparison __device__ int dev_a[N]; __device__ int dev_b[N]; // result in device global memory __device__ int dev_c[N];
// kernel routines __global__ void add_Cont() { // using constant memory int tid = blockIdx.x * blockDim.x + threadIdx.x; if(tid < N){ dev_c[tid] = dev_a_Cont[tid] + dev_b_Cont[tid]; } } __global__ void add() { //not using constant memory int tid = blockIdx.x * blockDim.x + threadIdx.x; if(tid < N){ dev_c[tid] = dev_a[tid] + dev_b[tid]; } }
/*----------- GPU using constant memory ------------------------*/ printf("GPU using constant memory\n"); for(int i=0;i<N;i++) { // load arrays with some numbers a[i] = i; b[i] = i*2; } // copy vectors to constant memory cudaMemcpyToSymbol(dev_a_Cont,a,N*sizeof(int),0,cudaMemcpyHostToDevice); cudaMemcpyToSymbol(dev_b_Cont,b,N*sizeof(int),0,cudaMemcpyHostToDevice); cudaEventRecord(start, 0); // start time add_Cont<<<B,T>>>(); // does not need array ptrs cudaThreadSynchronize(); // wait for all threads to complete cudaEventRecord(stop, 0); // end time cudaMemcpyFromSymbol(a,"dev_a_Cont",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaMemcpyFromSymbol(b,"dev_b_Cont",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaMemcpyFromSymbol(c,"dev_c",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsed_time_Cont, start, stop); Watch for this zero. I missed it off and it took some time to spot Missed originally
/*----------- GPU not using constant memory ------------------------*/ printf("GPU using constant memory\n"); for(int i=0;i<N;i++) { // load arrays with some numbers a[i] = i; b[i] = i*2; } // copy vectors to constant memory cudaMemcpyToSymbol(dev_a_Cont,a,N*sizeof(int),0,cudaMemcpyHostToDevice); cudaMemcpyToSymbol(dev_b_Cont,b,N*sizeof(int),0,cudaMemcpyHostToDevice); cudaEventRecord(start, 0); // start time add<<<B,T>>>(); // does not need array ptrs cudaThreadSynchronize(); // wait for all threads to complete cudaEventRecord(stop, 0); // end time cudaMemcpyFromSymbol(a,"dev_a_Cont",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaMemcpyFromSymbol(b,"dev_b_Cont",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaMemcpyFromSymbol(c,"dev_c",N*sizeof(int),0,cudaMemcpyDeviceToHost); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsed_time, start, stop);
Speedup around 1.2 after first launch (20%) 1st launch, 1.6 2nd run, 1.217 3rd run, 1.225