1.global memory
在 CUDA 中，一般的数据复制到的显卡内存的部份，称为 global memory 。这些内存是没有 cache 的，而且，存取 global memory 所需要的时间（即 latency）是非常长的，通常是数百个 cycles。由于 global memory 并没有 cache ，所以要避开巨大的 latency 的方法，就是要利用大量的 threads。假设现在有大量的threads 在同时执行，那么当一个 thread 读取内存，开始等待结果的时候，GPU 就可以立刻切换到下一个 thread，并读取下一个内存位置。因此，理想上当 thread 的数目够多的时候，就可以完全把 global memory 的巨大 latency 隐藏起来了。
2.内存存取模式
显卡上的内存是 DRAM，因此最有效率的存取方式，是以连续的方式存取。当一个 thread 在等待内存的数据时，GPU 会切换到下一个 thread。也就是说，实际上执行的顺序是类似 thread 0 -> thread 1 -> thread 2 -> ...
理论上 256 个 threads 最多只能隐藏 256 cycles 的 latency。但是 GPU 存取 global memory 时的 latency 可能高达 500 cycles 以上。
3.块
一个 block 中的 thread ，具有一个共享的 shared memory ，也可以进行同步工作。不同 block 之间的 thread 则不行。
一个 block 内的 thread 可以有共享的内存，也可以进行同步。
利用 __shared__ 声明的变量表示这是 shared memory ，是一个 block 中每个 thread 都共享的内存。它会使用在 GPU 上的内存，所以存取的速度相当快，不需要担心 latency 的问题。
__syncthreads() 是一个 CUDA 的内部函数，表示 block 中所有的 thread 都要同步到这个点，才能继续执行。
4.体系结构(以下皆以G80为准)
G80多核流处理器(Streaming Multiprocessor SM)只有一个。每个SM中含8个流处理器(Streaming Processor SP)。
在SM中，warp是线程调度的单位，每个warp包含32个线程。
G80中，每个SM中可驻留8个块，每个SM中最多可驻留768个线程，于是每个SM最多可驻留768/32=24个warp。
每个SM中寄存器的大小为8KB，每个SM的共享存储器大小为16KB，共享存储器是分配给块使用的。