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

下载链接

  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 }
时间: 2024-10-10 08:53:45

零内存拷贝和普通拷贝对比的相关文章

CUDA零内存拷贝 疑问考证

今天思考了一下CUDA零内存拷贝的问题,感觉在即将设计的程序中会派上用场,于是就查了一下相关信息. 以下是一些有帮助的链接: cuda中的零拷贝用法--针对二维指针 cuda中的零拷贝用法--针对一维指针 cuda零拷贝用法-二维结构体指针 浅谈CUDA零拷贝内存 经过调查发现,零拷贝技术适用于集中计算.较少内存拷贝次数的问题.比如向量点积.求和运算等问题. 既然零拷贝技术是在CPU上开辟内存空间,GPU可以直接访问该空间,那么我就产生了一个疑问:"如果CPU上开辟的空间大于GPU的可用空间的时

C语言中的位拷贝与值拷贝浅谈(转载)

注:C语言实现的PHP变量的赋值过程中,就涉及到了 深拷贝和浅拷贝 位拷贝拷贝的是地址(也叫浅拷贝),而值拷贝则拷贝的是内容(深拷贝).深拷贝和浅拷贝可以简单理解为:如果一个类拥有资源,当这个类的对象发生复制过程的时候,资源重新分配,这个过程就是深拷贝,反之,没有重新分配资源,就是浅拷贝. 位拷贝,及"bitwise assignment"是指将一个对象的内存映像按位原封不动的复制给另一个对象,所谓值拷贝就是指,将原对象的值复制一份给新对象. 在用"bitwise assig

C++之拷贝构造与拷贝赋值

拷贝构造和拷贝赋值------一个有点难的问题 介绍之前,我们需要首先了解深拷贝与浅拷贝的差异: 何为深拷贝,深拷贝不会复制指针,而是令目标对象拥有独立的资源,该资源是从元对象中复制,即先找到对象的指针,在通过指针拷贝其内容: 何为浅拷贝,即之赋值指针的地址,不会赋值指针的目标,容易引发double free异常,即多个目标指向同一个内存: 缺省拷贝构造函数和缺省拷贝赋值函数 如果一个类没有显示的定义一个拷贝构造函数和拷贝赋值运算符,则编译器会为其默认提供一个,但是这个函数只能进行浅拷贝: 如果

【转载】C++中的位拷贝和值拷贝

---恢复内容开始--- 原文:C++中的位拷贝和值拷贝 原文:http://blog.csdn.net/liam1122/article/details/1966617 为了便于说明我们以String类为例: 首先定义String类,而并不实现其成员函数. Class String { public: String(const char *ch=NULL);//默认构造函数 String(const String &str);//拷贝构造函数 ~String(void); String &

.net中String是引用类型还是值类型 以及 C#深层拷贝浅层拷贝

http://www.cnblogs.com/yank/archive/2011/10/24/2204145.html http://www.cnblogs.com/zwq194/archive/2012/08/06/2625403.html 关于String为值类型还是引用类型的讨论一直没有平息,最近一直在研究性能方面的问题,今天再次将此问题进行一次明确.希望能给大家带来点帮助. 如果有错误请指出. 来看下面例子: //值类型 int a = 1; int b = a; a = 2; Cons

拷贝控制2(拷贝控制和资源管理/交换操作)

为了定义拷贝构造函数和拷贝赋值运算符,我们首先必须确认此类型对象的拷贝语义.通常可以定义拷贝操作,使类的行为看起来像一个值或者像一个指针(即所谓的深拷贝和浅拷贝) 类的行为像一个值,意味着它应该也有自己的状态.当我们拷贝一个像值的对象时,副本和原对象是完全独立的.改变副本不会对原对象有任何影响,反之亦然 行为像指针的类则共享状态.当我们拷贝一个这种类的对象时,副本和原对象使用相同的底层数据.改变副本也会改变原对象,反之亦然 在我们使用过的标准库类中,标准库容器和 string 类的行为像一个值.

含有指针变量的类需要重写拷贝构造函数,拷贝赋值函数,析构函数

编译器自带拷贝构造(ctor)和拷贝赋值函数(operator =), 但是对于成员变量含有指针的类,其不能使用默认的拷贝赋值函数.因为使用默认的,会直接将指针指向的地址进行赋值 (浅拷贝,共享内存,共指一个对象),而不是分配一块内存,具有相同的数值 (深拷贝,独立,两个对象). 浅拷贝容易造成dangling pointer. 用一个例子来展示: 1 #ifndef __MYSTRING__ 2 #define __MYSTRING__ 3 4 class String{ 5 public:

C++笔记(11):拷贝控制(拷贝移动,构造赋值,析构)

控制对象拷贝,赋值,析构 拷贝构造函数,移动构造函数 拷贝赋值运算符,移动赋值运算符 析构函数 ------------------------------------------------------------------------------------------------------------------------------------- 1. 拷贝构造函数:参数必须是引用类型&,一般是const的 拷贝构造函数的第1个参数指的是对于自身类类型的引用 2.拷贝赋值运算符:本

Item 12:完整地拷贝对象(拷贝构造函数、复制运算符) Effective C++笔记

Item 12: Copy all parts of an object 在一个成熟的面向对象的C++系统中,只有两种拷贝对象的方式:复制构造函数和赋值运算符, 不妨称他们为拷贝函数. 拷贝函数属于编译器默认生成的函数(参考:Item 5:那些被C++默默地声明和调用的函数), 默认的拷贝函数确实会完整地拷贝对象,但有时我们选择重载拷贝函数,问题就出在这里! 一个正确拷贝函数的实现是这样的: class Customer{ string name; public: Customer::Custo

c++类的拷贝、赋值与销毁(拷贝构造函数、拷贝赋值运算符析构函数)

拷贝构造函数     如果一个构造函数的第一个参数是自身类类型的引用,且任何额外参数都有默认值,则此构造函数是拷贝构造函数. 拷贝构造函数第一个参数必须是一个引用类型.此参数几乎总是一个const的引用.拷贝构造函数在几种情况下都会被隐式地使用.因此,拷贝构造函数通常不应该是explicit的. 合成拷贝构造函数 与合成默认构造函数不同,即使我们定义了其他构造函数,编译器也会为我们合成一个拷贝构造函数. 对某些类来说,合成拷贝构造函数用来阻止我们拷贝该类类型的对象.而一般情况,合成的拷贝构造函数