在我们编写CUDA程序时我们经常要对CPU和GPU上分配内存和管理,这样就增加了编写程序发复杂度。在cuda6后引入了统一寻址(Unified Memory)技术,该技术使得CPU和GPU使用同一个指针对同一块内存进行处理,省去了原始的在CPU和CPU分配内存然后来回拷贝的过程,简化了程序的编写。
统一寻址创建了一块托管内存(managed memory),这是在CPU和GPU共享的内存,它在CPU和GPU之间架起了桥梁。这块托管内存CPU和GPU都能访问通过单一指针访问得到,最关键的一点是,系统会自动地在host端和device端转移在Unified Memory内存分配的数据,这样就好像看起来CPU内存好像运行在GPU,GPU内存好像运行在CPU。
下面写了一个简单的向量相加程序来比较使用统一寻址和不使用统一寻址程序的复杂度和计算时间。
1.传统的向量相加程序:
#include <cuda_runtime.h> #include <stdio.h> #define NUM 1024*2 __global__ void add(int *a,int *b,int *c){ int tid=threadIdx.x+blockIdx.x*blockDim.x; if(tid<NUM){ c[tid]=a[tid]+b[tid]; } } int main(){ int i; int *a,*b,*c; int *a_device,*b_device,*c_device; cudaEvent_t start,end; float elapsedTime; cudaEventCreate(&start); cudaEventCreate(&end); cudaEventRecord(start,0); a=(int *)malloc(NUM*sizeof(int)); b=(int *)malloc(NUM*sizeof(int)); c=(int *)malloc(NUM*sizeof(int)); cudaMalloc((void **)&a_device,NUM*sizeof(int)); cudaMalloc((void **)&b_device,NUM*sizeof(int)); cudaMalloc((void **)&c_device,NUM*sizeof(int)); for(i=0;i<NUM;i++){ a[i]=i; b[i]=i; } cudaMemcpy(a_device,a,sizeof(int)*NUM,cudaMemcpyHostToDevice); cudaMemcpy(b_device,b,sizeof(int)*NUM,cudaMemcpyHostToDevice); add<<<NUM/1024+1,1024>>>(a_device,b_device,c_device); cudaDeviceSynchronize(); cudaMemcpy(c,c_device,sizeof(int)*NUM,cudaMemcpyDeviceToHost); cudaEventRecord(end,0); cudaEventSynchronize(end); cudaEventElapsedTime(&elapsedTime,start,end); printf("tie===%3.1f ms\n",elapsedTime); cudaFree(a_device); cudaFree(b_device); cudaFree(c_device); free(a); free(b); free(c); return 0; }
运行结果:tie===0.3 ms
2.使用统一寻址方式:
#include <cuda_runtime.h> #include <stdio.h> #define NUM 1024*2 __global__ void add(int *a,int *b,int *c){ int tid=threadIdx.x+blockIdx.x*blockDim.x; if(tid<NUM){ c[tid]=a[tid]+b[tid]; } } int main(){ int i,j; int *a,*b,*c; cudaEvent_t start,end; float elapsedTime; cudaEventCreate(&start); cudaEventCreate(&end); cudaEventRecord(start,0); cudaMallocManaged(&a,NUM*sizeof(int)); cudaMallocManaged(&b,NUM*sizeof(int)); cudaMallocManaged(&c,NUM*sizeof(int)); for(i=0;i<NUM;i++){ a[i]=i; b[i]=i; } add<<<NUM/1024+1,1024>>>(a,b,c); cudaDeviceSynchronize(); cudaEventRecord(end,0); cudaEventSynchronize(end); cudaEventElapsedTime(&elapsedTime,start,end); printf("tie===%3.1f ms\n",elapsedTime); cudaFree(a); cudaFree(b); cudaFree(c); return 0; }
运行结果:tie===10.2 ms
从以上代码的比较可以看出,使用统一寻址方式可以减少代码编写的复杂度,但不会减少代码运行的时间复杂度,反而还增加了时间复杂度。因此在以后编写代码时根据自己的需求使用统一寻址编码方式。为什么会增加这么大的时间复杂度?,目前还不太清楚,希望有理解的朋友能赐教一下。
参考:http://www.ouccloud.com/356.html
相关推荐
CUDA02 - 访存优化和Unified Memory.doc
Unified memory GPGPU 2015: High Performance Computing with CUDAUniversity of Cape Town (South Africa), April, 20th-24th, 2015Manuel Ujaldón Associate Professor @ Univ. of Malaga (Spain) Conjoint ...
Bob Crovella, 6/18/2020CUDA UNIFIED MEMORY2AGENDA• Managed Memory - basic idea, objectives, benefits• Demand-Paging, Oversubscription, Concurrency, Atomics• Use Cases: Deep Copies, Linked Lists, ...
Compute unified device architecture (CUDA) is a software development platform that allows us to run C-like programs on the nVIDIA graphics processing unit (GPU). This paper presents an acceleration ...
MadisonSummary / Objective• Premise: Managing and optimizing host-device data transfers has been challenging• Key point: Unified Memory (UM) support in CUDA 6 simplifies the programmer’s job• The ...
Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0 Compute Mode: (multiple host threads can use ::cudaSetDevice() with device simultaneously) > ...
Device supports Unified Addressing (UVA): Yes Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0 Compute Mode: (multiple host threads can use ::cudaSetDevice() with device simultaneously) > ...
目标是确定哪种内存放置策略最适合使用CUDA内核的给定数据对象(数组)。 我们使用不同的内存放置策略准备内核的代码变体,输入数据大小并全部运行。 收集性能和配置文件信息,并将其制作为带标签的训练数据集。 ...
CUDA(Compute Unified Device Architecture)是基于 CPU 和 GPU 异构的常用编程模型,CUDA 编程模型直接使 用类 C 语言进行开发,大大减少了编程工作,通过简洁的代码就可以调用 GPU 进行并行计算,从而提高大规模...
本文总结了基于GPU平台的字典压缩与解压缩技术,讨论了使用CUDA(Compute Unified Device Architecture)实现高效的压缩和解压缩算法。该技术基于GPU的并行计算能力,通过 memory access 和parallel assembling技术...