零内存拷贝和普通拷贝对比

  1 /*
  2 * Copyright 1993-2010 NVIDIA Corporation.  All rights reserved.
  3 *
  4 * NVIDIA Corporation and its licensors retain all intellectual property and
  5 * proprietary rights in and to this software and related documentation.
  6 * Any use, reproduction, disclosure, or distribution of this software
  7 * and related documentation without an express license agreement from
  8 * NVIDIA Corporation is strictly prohibited.
  9 *
 10 * Please refer to the applicable NVIDIA end user license agreement (EULA)
 11 * associated with this source code for terms and conditions that govern
 12 * your use of this NVIDIA software.
 13 *
 14 */
 15 
 16 
 17 #include "../common/book.h"
 18 #include "cuda.h"
 19 #include "cuda_runtime.h"
 20 #include "device_launch_parameters.h"
 21 #include "device_functions.h"
 22 #define imin(a,b) (a<b?a:b)
 23 
 24 const int N = 33 * 1024 * 1024;
 25 const int threadsPerBlock = 256;
 26 const int blocksPerGrid =
 27 imin(32, (N + threadsPerBlock - 1) / threadsPerBlock);
 28 
 29 
 30 __global__ void dot(int size, float *a, float *b, float *c) {
 31     __shared__ float cache[threadsPerBlock];
 32     int tid = threadIdx.x + blockIdx.x * blockDim.x;
 33     int cacheIndex = threadIdx.x;
 34 
 35     float   temp = 0;
 36     while (tid < size) {
 37         temp += a[tid] * b[tid];
 38         tid += blockDim.x * gridDim.x;
 39     }
 40 
 41     // set the cache values
 42     cache[cacheIndex] = temp;
 43 
 44     // synchronize threads in this block
 45     __syncthreads();
 46 
 47     // for reductions, threadsPerBlock must be a power of 2
 48     // because of the following code
 49     int i = blockDim.x / 2;
 50     while (i != 0) {
 51         if (cacheIndex < i)
 52             cache[cacheIndex] += cache[cacheIndex + i];
 53         __syncthreads();
 54         i /= 2;
 55     }
 56 
 57     if (cacheIndex == 0)
 58         c[blockIdx.x] = cache[0];
 59 }
 60 
 61 
 62 float malloc_test(int size) {
 63     cudaEvent_t     start, stop;
 64     float           *a, *b, c, *partial_c;
 65     float           *dev_a, *dev_b, *dev_partial_c;
 66     float           elapsedTime;
 67 
 68     HANDLE_ERROR(cudaEventCreate(&start));
 69     HANDLE_ERROR(cudaEventCreate(&stop));
 70 
 71     // allocate memory on the CPU side
 72     a = (float*)malloc(size*sizeof(float));
 73     b = (float*)malloc(size*sizeof(float));
 74     partial_c = (float*)malloc(blocksPerGrid*sizeof(float));
 75 
 76     // allocate the memory on the GPU
 77     HANDLE_ERROR(cudaMalloc((void**)&dev_a,
 78         size*sizeof(float)));
 79     HANDLE_ERROR(cudaMalloc((void**)&dev_b,
 80         size*sizeof(float)));
 81     HANDLE_ERROR(cudaMalloc((void**)&dev_partial_c,
 82         blocksPerGrid*sizeof(float)));
 83 
 84     // fill in the host memory with data
 85     for (int i = 0; i<size; i++) {
 86         a[i] = i;
 87         b[i] = i * 2;
 88     }
 89 
 90     HANDLE_ERROR(cudaEventRecord(start, 0));
 91     // copy the arrays ‘a‘ and ‘b‘ to the GPU
 92     HANDLE_ERROR(cudaMemcpy(dev_a, a, size*sizeof(float),
 93         cudaMemcpyHostToDevice));
 94     HANDLE_ERROR(cudaMemcpy(dev_b, b, size*sizeof(float),
 95         cudaMemcpyHostToDevice));
 96 
 97     dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b,
 98         dev_partial_c);
 99     // copy the array ‘c‘ back from the GPU to the CPU
100     HANDLE_ERROR(cudaMemcpy(partial_c, dev_partial_c,
101         blocksPerGrid*sizeof(float),
102         cudaMemcpyDeviceToHost));
103 
104     HANDLE_ERROR(cudaEventRecord(stop, 0));
105     HANDLE_ERROR(cudaEventSynchronize(stop));
106     HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
107         start, stop));
108 
109     // finish up on the CPU side
110     c = 0;
111     for (int i = 0; i<blocksPerGrid; i++) {
112         c += partial_c[i];
113     }
114 
115     HANDLE_ERROR(cudaFree(dev_a));
116     HANDLE_ERROR(cudaFree(dev_b));
117     HANDLE_ERROR(cudaFree(dev_partial_c));
118 
119     // free memory on the CPU side
120     free(a);
121     free(b);
122     free(partial_c);
123 
124     // free events
125     HANDLE_ERROR(cudaEventDestroy(start));
126     HANDLE_ERROR(cudaEventDestroy(stop));
127 
128     printf("Value calculated:  %f\n", c);
129 
130     return elapsedTime;
131 }
132 
133 
134 float cuda_host_alloc_test(int size) {
135     cudaEvent_t     start, stop;
136     float           *a, *b, c, *partial_c;
137     float           *dev_a, *dev_b, *dev_partial_c;
138     float           elapsedTime;
139 
140     HANDLE_ERROR(cudaEventCreate(&start));
141     HANDLE_ERROR(cudaEventCreate(&stop));
142 
143     // allocate the memory on the CPU
144     HANDLE_ERROR(cudaHostAlloc((void**)&a,
145         size*sizeof(float),
146         cudaHostAllocWriteCombined |
147         cudaHostAllocMapped));
148     HANDLE_ERROR(cudaHostAlloc((void**)&b,
149         size*sizeof(float),
150         cudaHostAllocWriteCombined |
151         cudaHostAllocMapped));
152     HANDLE_ERROR(cudaHostAlloc((void**)&partial_c,
153         blocksPerGrid*sizeof(float),
154         cudaHostAllocMapped));
155 
156     // find out the GPU pointers
157     HANDLE_ERROR(cudaHostGetDevicePointer(&dev_a, a, 0));
158     HANDLE_ERROR(cudaHostGetDevicePointer(&dev_b, b, 0));
159     HANDLE_ERROR(cudaHostGetDevicePointer(&dev_partial_c,
160         partial_c, 0));
161 
162     // fill in the host memory with data
163     for (int i = 0; i<size; i++) {
164         a[i] = i;
165         b[i] = i * 2;
166     }
167 
168     HANDLE_ERROR(cudaEventRecord(start, 0));
169 
170     dot << <blocksPerGrid, threadsPerBlock >> >(size, dev_a, dev_b,
171         dev_partial_c);
172 
173     HANDLE_ERROR(cudaThreadSynchronize());
174     HANDLE_ERROR(cudaEventRecord(stop, 0));
175     HANDLE_ERROR(cudaEventSynchronize(stop));
176     HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
177         start, stop));
178 
179     // finish up on the CPU side
180     c = 0;
181     for (int i = 0; i<blocksPerGrid; i++) {
182         c += partial_c[i];
183     }
184 
185     HANDLE_ERROR(cudaFreeHost(a));
186     HANDLE_ERROR(cudaFreeHost(b));
187     HANDLE_ERROR(cudaFreeHost(partial_c));
188 
189     // free events
190     HANDLE_ERROR(cudaEventDestroy(start));
191     HANDLE_ERROR(cudaEventDestroy(stop));
192 
193     printf("Value calculated:  %f\n", c);
194 
195     return elapsedTime;
196 }
197 
198 
199 int main(void) {
200     cudaDeviceProp  prop;
201     int whichDevice;
202     HANDLE_ERROR(cudaGetDevice(&whichDevice));
203     HANDLE_ERROR(cudaGetDeviceProperties(&prop, whichDevice));
204     if (prop.canMapHostMemory != 1) {
205         printf("Device can not map memory.\n");
206         return 0;
207     }
208 
209     float           elapsedTime;
210 
211     HANDLE_ERROR(cudaSetDeviceFlags(cudaDeviceMapHost));
212 
213     // try it with malloc
214     elapsedTime = malloc_test(N);
215     printf("Time using cudaMalloc:  %3.1f ms\n",
216         elapsedTime);
217 
218     // now try it with cudaHostAlloc
219     elapsedTime = cuda_host_alloc_test(N);
220     printf("Time using cudaHostAlloc:  %3.1f ms\n",
221         elapsedTime);
222 }