How to caching Global data in on-chip (level 1) cache in Morden GPU

原创 2017年01月03日 23:21:10

1.Fermi arch

因为在CC 2.x(Compute Capability NVIDIA 计算能力)时,L1 Data Cache 还是可用的,我们可以缓存 local 和 global 的数据,不管ld(load 读)或者st (store 写),其默认的操作参数都是cache all 的。ld.ca 和 st.wb 是 其默认指令。

但是这样的话,SM之间会出现 cache coherency 问题!

英伟达的策略是 对于global的写操作 不往 L1 cache 进行的,因此提了一个 local write back global write evict 的策略。。如果有写操作命中 时,就将命中的data 置失效,并且通知 其他 SM 同样失效这个data(但是还没有找的其真正是怎么实现的!)

2.Kepler arch

CC 3.X 中,新引入了 Read Only Data Cache,其他与上一代类似。

1)重点说一下这个 Read Only Data Cache,之前我们说,写操作已经完全不往 L1这层 caching 了,那么load 指令呢?还是要放的,但是还是会有cc问题 ,那么为了应对这个,就只caching 那些只读数据,因此加了这个 Read Only Data Cache!

2)但是往 Read Only Data Cache 缓存数据,还要在 cuda 代码中,用一个 ldg()函数来告诉编译器。
这个 LDG() 只支持 CC 3.5 以上的GPU

Note:这个__ldg()函数对应的ptx指令是 ld.global.nc

eg

T __ldg(const T* address);

returns the data of type T located at address address, where T is char, short, int, long longunsigned char, unsigned short, unsigned int, unsigned long long, int2, int4, uint2, uint4, float, float2, float4, double, or double2. 

The operation is cached in the read-only data cache

Read more at: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4UiAE527J 
Follow us: @GPUComputing on Twitter | NVIDIA on Facebook

3.Maxwell

架构真正巨变的一代,从这一代开始,没有了 L1 Data cache 这个概念,只剩下:
a unified L1/texture cache of 24 KB used to cache reads from global memory

L1 / Texture 统一了,这就不好办了,,因为texture cache 是通过特殊 的 tex 模块来进行访问的,那该如何是好。
从这一代起,
Global memory accesses are only cached in L2

但是只读的数据还是可以通过 ldg()来缓存到L1/texture cache 中的,并且还需要通过关键字进行 限制 const 和 _restrict 。

Data that is not read-only for the entire lifetime of the kernel cannot be cached in the unified L1/texture cache for devices of compute capability 5.0. For devices of compute capability 5.2, it is, by default, not cached in the unified L1/texture cache, but caching may be enabled using the following mechanisms:
Perform the read using inline assembly with the appropriate modifier as described in the PTX reference manual;

Compile with the -Xptxas -dlcm=ca compilation flag, in which case all reads are cached, except reads that are performed using inline assembly with a modifier that disables caching;

Compile with the -Xptxas -fscm=ca compilation flag, in which case all reads are cached, including reads that are performed using inline assembly regardless of the modifier used.

When caching is enabled using some the three mechanisms listed above, devices of compute capability 5.2 will cache global memory reads in the unified L1/texture cache for all kernel launches except for the kernel launches for which thread blocks consume too much of the multiprocessor’s resources. These exceptions are reported by the profiler.

版权声明:本文为博主原创文章,未经博主允许不得转载。

CPU的高速缓存(cache)处理

带有高速缓存的CPU设计
  • hyhop150
  • hyhop150
  • 2016年05月18日 00:21
  • 684

深入System.Web.Caching命名空间 教你Hold住缓存管理(三)

本文分三篇,从缓存所在命名空间System.Web.Caching开始,详细的介绍.NET框架提供的缓存类和操作方法。看完之后你将学会: 第一篇-如何实现简单的数据缓存第二篇-缓存从文件中读取的...
  • lwj0310
  • lwj0310
  • 2014年06月21日 15:17
  • 563

cache 是什么意思 它包括的L1,L2,L3分别是什么东西?

CPU缓存   缓存大小也是CPU的重要指标之一,而且缓存的结构和大小对CPU速度的影响非常大,CPU内缓存的运行频率极高,一般是和处理器同频运作,工作效率远远大于系统内存和硬盘。实际工作时,CPU...
  • hemmingway
  • hemmingway
  • 2014年08月22日 09:30
  • 6891

cache目录没有操作权限

今天部署新服务器发现了cache文件夹木有权限读写 App 26772 stderr: Started GET "/investors" for 103.238.226.130 at 2015-08...
  • shiralwz
  • shiralwz
  • 2015年08月25日 17:07
  • 933

哪些图形算法是不能用GPU来计算的?

哪些图形算法是不能用 GPU 来计算的? 比如软阴影,全局光照,Final Gathering,或者某种渲染效果是不能用GPU计算的。我的意思是,GPU渲染真的能完全替代CPU渲染吗? 罗登,...
  • dj0379
  • dj0379
  • 2014年09月13日 00:37
  • 923

nginx源码分析1———进程间的通信机制六(UNIX域协议)

相关介绍 Unix域协议并不是一个实际的协议族,而是在单机上客户端与服务器通信的一种方法。...
  • sina_yangyang
  • sina_yangyang
  • 2015年08月06日 23:58
  • 375

cuda 函数前缀:device/global/host 相关问题

在深度学习caffe框架等多处,用到了CUDA函数,使代码加速,其文件类型为XXX.cu。在运行这些CUDA函数时,会遇到一下C++里没有的问题,比如函数前缀问题,如果在修改编写.cu文件时不注意,会...
  • u012905422
  • u012905422
  • 2016年10月23日 22:08
  • 3528

完整的拆分nginx访问日志

filter { grok { match =>[ "message","%{IPORHOST:clientip} \[%{HTTPDATE:tim...
  • zhaoyangjian724
  • zhaoyangjian724
  • 2016年09月07日 11:38
  • 756

GPU上缘何没有大量的cache

近年来,GPU广泛应用与高性能领域,其通用计算能力也得到了进一步的利用。与传统的CPU相比,GPU在处理能力和储存器带宽上有着明显的优势,在成本和功耗上也不需要付出太大的代价。         在当...
  • gaoxiang__
  • gaoxiang__
  • 2015年05月29日 09:48
  • 1251

Java中使用ResponseCache类构建本地缓存

Java中使用ResponseCache类构建本地缓存 CacheRequest CacheResponse 只是对GET请求进行缓存 缓存失效策略...
  • u013905744
  • u013905744
  • 2016年08月26日 10:45
  • 910
内容举报
返回顶部
收藏助手
不良信息举报
您举报文章:How to caching Global data in on-chip (level 1) cache in Morden GPU
举报原因:
原因补充:

(最多只允许输入30个字)