CUDA存储:
- 是GUP中最快的memory,每个线程的寄存器变量是私有的,当线程执行结束时变量就会失效,不同的设备架构,数量是不一样的。
- Shared Memory:片上,比Registers稍微慢一点点但是无线接近于它,同一个Block中的线程可以共享同一块Shared Memory。线程同步,保证写完了再读。
- Local Memory:当Registers不够了的时候会自动写入进Local Memory中,不用自动去写。会消耗较多的寄存器空间的大型结构或数组。
- Constant Memory:其范围是全局的,针对所有的Kernel,而Kernel也只能从constantMemory中去读取数据,而当一个warp中所有的线程都从同一个Memory地址读取数据时,constantMemory最优。
常量内存:
- 渲染图可以想象为无限远的线从后穿过打在屏幕上通过读取常量内存的信息去渲染。
代码部分:
struct Sphere
{
float r,g,b;
float radius;
float x,y,z;
__device__ float hit(float ox,float oy,float *n)
{
float dx = ox -x;
float dy = oy -y;
if(dx*dx+dy*dy<radius*radius)
{
float dz=sqrt(radius*radius - dx*dx -dy*dy);
*n=dz/sqrt(radius*radius);
return dz+z;
}
retrun -INF;
}
};
- 其中Hit表示相交
- 如果没有相交返回负无穷,有相交就返回距离
CUDA存储
- Texture Memory:在device memory中并且使用一个只读cache,实际上也和global memory在一起但是有其独有的只读cache。可以参考热传导模型举例分析。
- Global Memory:最常见的,是空间最大的,在device memory中并且和memory transaction对齐,合并式访存。
其中在连续线程中读取连续数据时速度最快
- Shared Memory:不能超过L1缓存,可以多线程访问存储器所以存储器又可以被划分为banks,其中每个bank周期可以相应一个地址。
在SharedMemory中可能存在Bank冲突即:Bank confilct
总结:
矩阵和存储让我们进一步了解了CUDA。