内存以及优化方法
内存读写速度
线程寄存器读写:1个时钟周期延迟
线程本地内存读写:非常慢
块的共享内存读写:1个时钟周期延迟,但是可能冲突
网格全局内存读写:500个时钟周期延迟,联合访问时会有隐含延迟。
网格的常量内存和纹理内存读取:500个时钟周期延迟。但是有缓存
内存模型
寄存器
修饰符:核函数中声明,没有修饰符。
共享范围:线程私有
速度:块。一个时钟周期延迟。当R/W依赖或寄存器内存冲突时会有延迟,但是活动线程数>=192时,可以忽略此延迟。
读取模式:读写。
生命周期:生命周期与线程一致。
优化方法:当每个块的线程数为64倍时,获得最佳效果。
本地内存
修饰符:无
范围:线程私有
速度:慢
读取模式:读写。
当寄存器用完时,将需要存储的值放在本地内存中。
共享内存
修饰符:__shared__
共享范围:块内所有线程。
速度:快,1个是时钟周期延迟。当组冲突时会有延迟,没有冲突时,会和寄存器一样快。
生命周期:整个线程块。
读取模式:读写。
缺点:有库冲突,半个warp中,多个线程访问相同的库。必须序列化访问才能解决冲突。
常量内存
修饰符:__constant__,必须在全局空间内和所有核函数之外声明。
共享范围:网格内,可以从主机访问
速度:速度与芯片上的缓存有关;慢速时,1次缓存未命中,所以从全局内存读取;快速时,都命中缓存。
读取模式:只读。
生命周期:随着应用程序结束而结束。
空间:64KB;
初始化:在host端使用 cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src,size_t count);初始化
与常量内存进行数据交换,使用案例:
__constant__ float const_data[256];
float data[256];
cudaMemcpyToSymbol(const_data, data, sizeof(data));
cudaMemcpyFromSymbol(data, const_data, sizeof(data));
纹理内存
共享范围:网格内,可以从主机访问。
速度与常量内存一致。
读取模式:只读。
生命周期:随着应用程序结束而结束。
块共享内存,块内所有线程都可访问,生命周期与块一致。
全局内存,所有线程都可以访问,生命周期随着应用程序结束而结束。
全局内存
修饰符:__device__或者在host端使用cudaMalloc函数动态声明。
共享范围:全局。
速度:慢,500个时钟周期。联合访问时会有冲突
读取模式:读写。
生命周期:随着应用程序结束而结束。
初始化:cudaMalloc
内存操作函数
将host拷贝到global
cudaMemcpyToSymbol(target,&src,sizeof(src));
将host变量拷贝到device,字节数为源变量的大小,host端变量传递的是地址。
将global拷贝到host
cudaMemcpyFromSymbol(&target,src,sizeof(src));
将device变量拷贝到host,字节数为源变量的大小,host端变量传递的是地址。
重置关联资源
cudaDeviceReset()
重置当前线程所关联过的当前设备的所有资源
标签:__,管理,读写,host,线程,内存,data From: https://www.cnblogs.com/RedNoseBo/p/17076786.html