cuda 学习之内存层次结构

最新推荐文章于 2024-09-29 00:06:28 发布
最新推荐文章于 2024-09-29 00:06:28 发布
阅读量3.1k 收藏 31
点赞数 10
分类专栏： cuda学习 文章标签： cuda 内存
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/xukang95/article/details/102855750
版权

文章目录

cuda 学习之内存层次结构

要想编写高效的程序，那么一定要对内存结构有比较深刻的认识，就像C/C++里面的堆内存，栈内存，全局存储区，静态存储区，常量区等。Cuda是并行计算框架，而GPU的内存有限，那么如果想编写高效的Cuda程序，首先要对其内存结构有一个简单的认识。

首先我们先上一张图，然后通过解释一些名词和代码来进行解释。

在这里插入图片描述
在这里插入图片描述

registers：寄存器

它是GPU片上告诉缓存器，执行单元可以以极低的延迟访问寄存器。寄存器的基本单元是寄存器文件（register file），每个寄存器文件大小为32bit。寄存器文件数量虽然客观，但是平均分给并行执行的线程，每个线程拥有的数量就非常有限了。**编程时，不要为每个线程分配过多的私有变量。**下面程序中，aBegin,aEnd,aStep,a等变量都是寄存器变量，每个线程都会维护这些变量。

__global__ void registerDemo(float *B,float *A ,int wA)
{
    int aBegin = wA*BLOCK_SIZE * blockIdx.y;
	int aEnd = aBegin + wA - 1;
	int aStep = BLOCK_SIZE;
	
	for(int a=aBegin;a<=aEnd;a+=aStep)
	{
	    //...
	}
}

local memory：局部存储器

  • 位于堆栈中，不在寄存器中的所有内容
  • 作用域为特定线程
  • 存储在global内存空间中，速度比寄存器慢很多

对于每个线程，局部存储器也是私有的。如果寄存器被消耗完，数据将被存储在局部存储器中。如果每个线程用了过多的寄存器，或声明了大型结构体或数组，或者编译期无法确定数组的大小，线程的私有数据就有可能被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据将被保存在显存中，而不是片上的寄存器或者缓存中，因此对local memory的访问速度比较慢。

shared memory：共享存储器

  • SM（SM = streaming multiprocessor）中的内存空间
  • 最大48KB
  • 作用域是线程块

共享存储器也是GPU片内的告诉存储器。它是一个块可以被同一block中的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快。是实现线程间通信的延迟最小的方法。共享存储器可用于实现多种功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公共结果（例如规约的结果）。

我们可以动态或者静态的分配shared Memory，其声明即可以在kernel内部也可以作为全局变量。其标识符为：__ shared __。

下面这句话静态的声明了一个2D的浮点型数组：

__shared__ float tile[size_y][size_x];

如果在kernel中声明的话，其作用域就是kernel内，否则是对所有kernel有效。

如果shared Memory的大小在编译器未知的话，可以使用extern关键字修饰，例如下面声明一个未知大小的1D数组：

extern __shared__ int tile[];

由于其大小在编译器未知，我们需要在每个kernel调用时，动态的分配其shared memory，也就是最开始提及的第三个参数：

kernel<<<grid, block, isize * sizeof(int)>>>(...)

应该注意到，只有1D数组才能这样动态使用。

static variable使用shared memory：

#include<iostream>
#include<stdio.h>
 
#if 1
__global__ void example(float *u)
{
	int i = threadIdx.x;
	__shared__ int tmp[4];
	tmp[i] = u[i]; 
	u[i] = tmp[i] * tmp[i] + tmp[3-i] ;
}
#endif
 
#if 1
 
int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u ; 
	size_t size = 4*sizeof(float);
 
	cudaMalloc(&dev_u , size);
	cudaMemcpy(dev_u,host_u,size,cudaMemcpyHostToDevice);
 
	example<<<1,4>>> (dev_u);
 
	cudaMemcpy(host_u , dev_u , size , cudaMemcpyDeviceToHost);
 
	cudaFree(dev_u);
 
	for(int i=0;i<4;i++)
		printf("%f\n",host_u[i]);
	return 0;
}
 
#endif

dynamic variable使用shared memory：

#include<iostream>
#include<stdio.h>
 
__global__ void example(float *u)
{
	int i = threadIdx.x;
  extern __shared__ int tmp[];
	tmp[i] = u[i];
	u[i] = tmp[i] * tmp[i] + tmp[3-i];
}
 
int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u;
	size_t size = 4*sizeof(float);
 
	cudaMalloc(&dev_u,size);
	cudaMemcpy(dev_u , host_u ,size , cudaMemcpyHostToDevice);
	example<<<1,4,  size  >>>>(dev_u);
 
	cudaMemcpy(host_u, dev_u,size,cudaMemcpyDeviceToHost);
	cudaFree(dev_u);
	for(int i=0;i<4;i++)
		printf("%f ",host_u[i]);
	return 0;
}
共享内存-Bank Conflict

为了获得较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，可以被同时访问。因此任何跨越b个不同的内存bank的对n个地址进行读取和写入的操作可以被同时进行，这样就大大提高了整体带宽 ——可达到单独一个bank带宽的b倍。但是很多情况下，我们无法充分发挥bank的功能，以致于shared memory的带宽非常的小，这可能是因为我们遇到了bank冲突。

当一个warp中的不同线程访问一个bank中的不同的字地址时，就会发生bank冲突。
如果没有bank冲突的话，共享内存的访存速度将会非常的快，大约比全局内存的访问延迟低100多倍，但是速度没有寄存器快。然而，如果在使用共享内存时发生了bank冲突的话，性能将会降低很多很多。在最坏的情况下，即一个warp中的所有线程访问了相同bank的32个不同字地址的话，那么这32个访问操作将会全部被序列化，大大降低了内存带宽。

NOTE:不同warp中的线程之间不存在什么bank冲突。

共享内存的地址映射方式

要解决bank冲突，首先我们要了解一下共享内存的地址映射方式。
在共享内存中，连续的32-bits字被分配到连续的32个bank中，这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的：

在这里插入图片描述

上图中数字为bank编号。这样的话，如果你将申请一个共享内存数组(假设是int类型)的话，那么你的每个元素所对应的bank编号就是地址偏移量(也就是数组下标)对32取余所得的结果，比如大小为1024的一维数组myShMem：

  • myShMem[4]: 对应的bank id为#4 (相应的行偏移量为0)
  • myShMem[31]: 对应的bank id为#31 (相应的行偏移量为0)
  • myShMem[50]: 对应的bank id为#18 (相应的行偏移量为1)
  • myShMem[128]: 对应的bank id为#0 (相应的行偏移量为4)
  • myShMem[178]: 对应的bank id为#18 (相应的行偏移量为5)

下面我介绍几种典型的bank访问的形式。

下面这这种访问方式是典型的线性访问方式(访问步长(stride)为1)，由于每个warp中的线程ID与每个bank的ID一一对应，因此不会产生bank冲突。

在这里插入图片描述

下面这种访问虽然是交叉的访问，每个线程并没有与bank一一对应，但每个线程都会对应一个唯一的bank，所以也不会产生bank冲突。

在这里插入图片描述

下面这种虽然也是线性的访问bank，但这种访问方式与第一种的区别在于访问的步长(stride)变为2，这就造成了线程0与线程28都访问到了bank 0，线程1与线程29都访问到了bank 2…，于是就造成了2路的bank冲突。我在后面会对以不同的步长(stride)访问bank的情况做进一步讨论。

在这里插入图片描述

下面这种访问造成了8路的bank冲突。

在这里插入图片描述

这里我们需要注意，下面这两种情况是两种特殊情况：

在这里插入图片描述

上图中，所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制(当一个warp中的所有线程访问一个bank中的同一个字(word)地址时，就会向所有的线程广播这个字(word))，这种情况并不会发生bank冲突。

同样，这种访问方式也不会产生bank冲突：

在这里插入图片描述

这就是所谓的多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。

NOTE:这里的多播机制(multicast)只适用于计算能力2.0及以上的设备

​ 详细请见 共享内存之bank冲突

constant memory：常数存储器

  • 属于全局内存，大小64KB
  • 线程请求同一个数据时很快，请求不同的数据时性能下降
  • 在运行中不变，所有constant变量的值必须在kernel启动之前从host设置
  • __ global __函数参数通过 constant memory穿的到device端， 限定4 KB，即kernel参数通过常量内存传递

它是只读的地址空间。常熟存储器中的数据位于显存，但拥有缓存加速。常数存储器的空间较小，在Cuda程序中用于存储需要频繁访问的只读参数。当来自同一half-warp的线程访问常数存储器中的同一数据时，如果发生缓存命中，那么只需要一个周期就可以获得数据。常数存储器有缓存机制，用以节约带宽，加快访问速度。每个SM拥有8KB的常数存储器缓存。常数存储器是只读的，因此不存在缓存一致性问题。

constant memory的使用：

#include<iostream>
 
using namespace std;
 
__constant__ int devVar = 100;
 
__global__ void xminus(int *a)
{
	int i = threadIdx.x;
	a[i] = devVar+i;
}
 
int main()
{
	int *h_a = (int*)malloc(4*10) ;
	int *d_a ;
	cudaMalloc(&d_a, 4*10) ;
	cudaMemset(d_a, 0, 40) ;
 
	xminus<<<1,4>>>(d_a);
	
	cudaMemcpy(h_a, d_a, 4*10, cudaMemcpyDeviceToHost) ;
 
	for(int i = 0; i < 4 ; i++)
		cout << h_a[i] << " " ;
	cout << endl ;
}

texture memory：纹理存储器

类似constant memory，是只读内存，以某种形式访问的时候可以提升性能。原本是用在OpenGL和DirectX渲染管线中的。
有用的特点：

  • 不需考虑要聚合coalescing访问的问题
  • 通过“CUDA Array”进行缓存的2D或3D空间的数据位置
  • 在1D，2D或3D数组上进行快速插值
  • 将整数转换为“unitized”浮点数

global memory：全局存储器

  • 独立于GPU核心的硬件RAM
  • GPU绝大多数内存空间都是全局内存
  • 全局内存的IO是GPU上最慢的IO形式（除了访问host端内存）

全局存储器位于显存（占据了显存的绝大部分），CPU、GPU都可以进行读写访问。整个网格中的任意线程都能读写全局存储器的任意位置由于全局存储器是可写的。全局存储器能够提供很高的带宽，但同时也具有较高的访存延迟。显存中的全局存储器也称为线性内存。线性内存通常使用cudaMalloc()函数分配，cudaFree()函数释放，并由cudaMemcpy()进行主机端与设备端的数据传输。

此外，也可以使用__device__关键字定义的变量分配全局存储器，这个变量应该在所有函数外定义，必须对使用这个变量的host端和device端函数都可见才能成功编译。在定义__device__变量的同时可以对其赋值。

static variable使用global memory：

global_mem_static.cu:

#include<stdio.h>
#include<iostream>
 
__device__ float devU[4];
__device__ float devV[4];
 
//__global__ function 
__global__ void addUV()
{
	int i = threadIdx.x;
	devU[i] += devV[i];
}
 
int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};
	
	int size = 4* sizeof(float);
 
	//cudaMemcpyToSymbol:将数据复制到__constant__或者__device__变量中
	//cudaMemcpyFromSymbol:同上相反
	//cudaMalloc:在设备端分配内存
	//cudaMemcpy:数据拷贝
	//cudaFree():内存释放
	//cudaMemset():内存初始化
	cudaMemcpyToSymbol(devU,hostU,size,0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(devV,hostV,size,0,cudaMemcpyHostToDevice);
 
	addUV<<<1,4>>>();
 
	cudaMemcpyFromSymbol( hostU,devU,size,0,cudaMemcpyDeviceToHost );
 
	for(int i=0;i<4;i++)
		printf("hostU[%d] = %f\n",i,hostU[i]);
	return 0;
}

dynamic variable使用global memory：

global_mem_dynamic.cu:

#include<iostream>
#include<stdio.h>
 
__global__ void add4f(float *u , float *v)
{
	int i = threadIdx.x;
	u[i] += v[i];
}
 
void print(float * U ,int size)
{
	for(int i=0;i<4;i++)
	{
		printf("U[%d] = %f\n",i,U[i]);
	}
}
 
int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};
 
	float * devU ;
	float * devV ; 
	int size = sizeof(float) * 4;
	
	//在设备内存上分配空间
	cudaMalloc( &devU,size );
	cudaMalloc( &devV,size );
	//数据拷贝
	cudaMemcpy( devU ,hostU ,size ,cudaMemcpyHostToDevice );
	cudaMemcpy( devV ,hostV ,size ,cudaMemcpyHostToDevice );
 
	add4f<<<1,4>>> (devU,devV);
	//数据返回
	cudaMemcpy(hostU,devU,size,cudaMemcpyDeviceToHost);
 
	print(hostU,size);
	//释放空间
	cudaFree(devV);
	cudaFree(devU);
	return 0;
}
全局内存-对齐与合并

CUDA全局内存的访问是通过”内存事务“实现的，其分类128字节（L1/L2缓存均参与）和32字节（L2缓存参与）两种。

缓存加载（L1+L2）:这种情况下，”内存事务“中加载的缓存粒度是128字节。

  1. 对齐合并访问，线程束首地址对齐128字节，且连续访问128字节内存。只需一次”128字节内存事务“即可完成内存请求。效率100%。

    在这里插入图片描述

  2. 访问对齐，随机访问128字节内存。由于所请求的内存地址仍然在一个缓存行中，因此，也只需一次”128字节内存事务“即可完成内存请求。效率100%。

在这里插入图片描述

  1. 非对齐的连续访问，线程束请求的字节未对齐128，而是分布在两个128字节段范围内。由于启用L1缓存，因此加载必须从首地址128的倍数开始，因此需要加载0-127和127-255两个”128字节内存事务“才能完成内存请求。其中一半数据是请求之外的，因此效率=请求加载的全局内存/所需加载的全局内存=50%。 这主要是由于缓存的加载模式导致，因为缓存不是一次性只加载一个数据，而是一批数据，如128字节，而且又必须保证对齐操作，这就导致了加载的浪费。

在这里插入图片描述
4. 同一warp中线程只访问一个地址。只需要一个”128内存事务“就可以完成请求，但是效率却非常低。因为所需4字节，而加载了128字节。效率为4/128=3.125%。

在这里插入图片描述

  1. 最坏的情况，warp中32线程所请求的内存全部分散，因此加载的”内存事务“的可能在0-32之间不等。完成一次内存请求最差需要进行32次”内存事务“，而加载的128字节中，却只有4字节是warp所需的。

在这里插入图片描述
非缓存加载（L2）:这种情况下，”内存事务“中加载的缓存粒度是32字节，这是比128字节更细粒度的加载，会对非对齐或非合并的访问带来好处。类似地，对比上述情况，逐一分析。

  1. 对齐合并访问，一个warp的128字节的请求，需要4个32字节”内存事务“完成。效率100%。

在这里插入图片描述

  1. 对齐，但访问是不连续。所需地址将占用4个内存段，不会产生加载浪费。效率100%。

在这里插入图片描述
3. 非对齐的连续访问。由于请求内存首地址没有对齐128字节，请求的地址最多落在5个内存段中，总线利用率至少为80%，相比缓存加载50%有了很大的提供，这主要是由于加载了更少的未请求字节。

在这里插入图片描述

  1. warp中线程束访问一个内存地址。效率为4/32=12.5。非缓存加载要优于缓存加载。

    在这里插入图片描述

  2. warp中线程的请求内存地址全部分散，则需要的”内存事务“的也是0-32个，但每个内存事务是32字节，而不再是128字节。这也是非缓存加载优于缓存加载的地方。

在这里插入图片描述

host memory :主机端内存

主机端内存，即CPU对应的我们普通意义上的内存。主机端内存分为两种：可分页内存（pageable memory）和页锁定（page-locked 或pinned）内存。可分页内存即为通过操作系统API（malloc()，new()）分配的存储器空间：而页锁定内存始终不会被分配到低俗的虚拟内存中，能够保证存在于屋里内存中，并且能够通过DMA加速与设备端的通信。

总结

下面一张图可以总结这些内存与线程的关系

在这里插入图片描述

参考文献：
https://blog.csdn.net/qq_17239003/article/details/79038333
https://blog.csdn.net/lavorange/article/details/20465869
https://segmentfault.com/a/1190000007533157?utm_source=tag-newest

xukang95
关注
专栏目录
GPU深度学习性能的三驾马车：Tensor Core、内存带宽与内存层次结构
shengyin714959的博客
08-05 51
以一个简单的 A*B=C 矩阵乘法为例（其中每个矩阵的大小都是 32×32），我们会将反复访问的数据加载到共享内存（shared memory）中，这样做的主要原因是共享内存的延迟约为全局内存的六分之一（200 个周期 vs 34 个周期）。例如，在训练规模为GPT-3级别的大型神经网络时使用了大矩阵（由于矩阵越大，对张量计算核心进行运算越有利），即便是这种情况下，张量计算核心的利用率约为45-65％，这说明即使是训练大型神经网络，张量计算核心也有约 50% 的时间处于闲置状态。
深度学习运算：CUDA 编程简介
gongdiwudu的专栏
04-26 2661
如今，当我们谈论深度学习时，通常会将其实现与利用 GPU 来提高性能联系起来。GPU（图形处理单元）最初设计用于加速图像、2D 和 3D 图形的渲染。然而，由于它们能够执行许多并行操作，因此它们的实用性超出了深度学习等应用程序。
CUDA - 内存分配函数
weixin_67051070的博客
07-25 3405
主要区别:总结:
CUDA的存储器结构
XYSYZX的博客
10-26 356
CUDA的存储器结构 1. shared memory CUDA C有共享内存（shared memory），它的关键字是__share__。添加了该关键字的变量表明该变量是一个共享变量。 共享内存被一个线程块内的所有线程共享，块内线程可以修改它的值，它的生命周期和它所在的线程块一样。 因为共享内存的缓冲区在物理GPU上，而不是GPU外的系统内存上，所以访问共享内存的延迟要低于访问普通缓冲区的延...
cudaMalloc与cudaMemcpy详解
最新发布
qq_43527718的博客
09-29 141
【代码】cudaMalloc与cudaMemcpy详解。
【CUDA编程】CUDA内存模型
非晚非晚的博客
11-04 2568
thread：每个线程都用自己的registers（寄存器）和local memory（局部内存） block：每个线程块(block)内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源 grid：每个grid都有自己的global memory（全局内存），constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用。其中常量内存和纹理内存为只读内存空间。
CUDA从入门到放弃（六）：CUDA内存结构（Memory Hierarchy）
mingo_敏
03-26 1405
CUDA从入门到放弃（六）：CUDA内存结构（Memory Hierarchy）CUDA线程在执行过程中可以从多个内存空间访问数据。每个线程都有私有的局部内存。每个线程块具有共享内存，该内存对所有线程块内的线程可见，并且与线程块具有相同的生命周期。线程块集群中的线程块可以相互执行对共享内存的读取、写入和原子操作。所有线程都可以访问相同的全局内存。此外，还有两个所有线程都可以访问的只读内存空间：常量内存空间和纹理内存空间。对于同一应用程序来说，全局内存、常量内存和纹理内存空间在内核启动之间是持久的。
cuda存储结构
Tulip_Alice的博客
01-17 1667
cuda
【CUDA 基础】4.1 内存模型概述
weixin_30699235的博客
06-26 271
title: 【CUDA 基础】4.1 内存模型概述 categories: - CUDA - Freshman tags: - CUDA内存模型 - CUDA内存层次结构 - 寄存器 - 共享内存 - 本地内存 - 常量内存 - 纹理内存 - 全局内存 toc: true date: 2018-04-28 22:28:08 Abstract: ...
cuda.rar_cuda_cuda学习_cuda编程
09-24
这个“cuda.rar_cuda_cuda学习_cuda编程”压缩包显然包含了针对CUDA编程的资源，特别是对于初学者的学习教程。 CUDA编程的核心在于利用GPU的并行计算能力，它允许开发者直接编写C/C++代码，通过特殊的CUDA C/C++...
CUDA.rar_CUDA ppt_GPU_cuda_cuda学习_cuda编程，PPT
09-23
2. **CUDA编程模型**：讲解如何通过CUDA C/C++编程接口来访问GPU资源，如定义__global__函数，理解共享内存、全局内存、常量内存和纹理内存等不同内存层次。 3. **CUDA编程环境设置**：指导如何安装CUDA Toolkit，...
CUDA编程：线程层次结构与矩阵加法
"CUDA编程要点与线程层次结构详解" CUDA编程是一种用于利用NVIDIA GPU进行并行计算的技术。CUDA的核心是CUDA C/C++，它扩展了标准C/C++，允许开发者直接控制GPU的硬件资源来加速计算密集型任务。本文将深入探讨CUDA...
CUDA——内存
UCAS_HMM的博客
10-18 688
CUDA——内存原理
[CUDA编程] --- CUDA内存模型
M2嵌入式
08-25 1157
用表格进行总结：设备存储器的重要特征下图表达的更加直观一些。Local, Global, Contant, Texture为片外DRAM，其中Global, Constant, texture内存在Host端代码声明，所有线程可见。SM拥有私有的Registers和Shared Memory（其实还有SM私有的L1 cache以及共有的L2 cache，下一篇会详细说明），Constant和Texture内存有专有的Caches， 这些在片上。
CUDA GPU内存结构
weixin_42730667的博客
03-01 2089
了解CUDA编程首先必不可少的，需要熟悉英伟达GPU的内存分布，如下图： 整体内存结构按照层次划分，和OpenCL类似： SM：流处理器簇，为GPU的对小硬件调度单位，一次性可以调度多个SP SP:GPU 硬件内核执行单元，每个SP相当于一个CPU核。 Register File:寄存器文件，访问速度最快的存储单元，用来存储SP上运行的线程内部活跃的寄存器。 shared memor...
CUDA精讲（3）-- GPU体系结构之内存层次结构
EdiosnMa的博客
10-10 2316
内存层次结构 内存层次结构是GPU当中非常重要的一部分，对于性能的优化至关重要，内存的层次结构可以分成一下三个部分来讲： Register & Local memory Shared memory Global & constant memory GPU内存层次概览 每个thread的local variable对应：register、local memory 每个block中的shared variable对应一个shared memory 每个grid中的globel var
CUDA内存介绍
weixin_41694059的博客
12-07 955
CUDA内存介绍 Global Memory Global Memory就是我们在书写CUDA程序时最常使用的内存，cudaMemcpy也是从CPU 内存拷贝到Global Memory。 Global Mem能被所有thread访问。 Device可读/写, host可读/写。 Local Memory 每个thread都拥有自己私有的local memory，负责存储一些局部变量（automatic variable）。 对于局部变量而说，一些小型的局部变量会被放到register里，当register
内存对齐指令
chubaisheng8627的博客
06-20 169
通常来说内存对齐很能提高速度的，使用如下指令在两个操作系统下面令内存对齐 #ifdef_WIN32 typedef__declspec(align(16))struct #elif_LINUX typedef__attribute__((aligned(16)))st...
CUDA——内存管理
xiaohu的博客
05-23 2910
内存管理 CUDA是C语言的扩展，内存方面基本集成了C语言的方式，由程序员控制CUDA内存，当然，这些内存的物理设备是在GPU上的，而且与CPU内存分配不同，CPU内存分配完就完事了，GPU还涉及到数据传输，主机和设备之间的传输。 为达到最优性能，CUDA提供了在主机端准备设备内存的函数，并且显式地向设备传递数据，显式的从设备取回数据。 统一内存Unified Memory 统一内存是可从系统中的任何处理器访问的单个内存地址空间（参见上图）。 这种硬件/软件技术允许应用程序分配可以从 CPU 或 GPU 上
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值