CUDA学习笔记

函数类型限定符

__global__

__host__

__device__

变量类型限定符

__device__

__shared__

__constant__

CUDA线程概念

网格（Grid）、线程块（Block）和线程（Thread）的组织关系

线程索引的计算公式

dim3结构类型

函数类型限定符

global

表示一个内核函数，主机（CPU）上调用，但在设备（GPU）上执行

host

由CPU调用，且在CPU执行的函数，默认的函数限定符

device

设备执行的程序，device前缀定义的函数只能在GPU上执行，所以device修饰的函数里面不能调用一般常见的函数

变量类型限定符

device

shared

constant

CUDA线程概念

CUDA线程的概念与CPU的概念基本一致：进程是系统资源分配的基本单位，而线程是CPU 能调度和独立运行的基本单位（最小单元）

网格（Grid）、线程块（Block）和线程（Thread）的组织关系

多个线程thread组成一个block，多个block组成一个grid.

CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread）组成，相当于把GPU上的计算单元分为若干（2~3）个网格，每个网格内包含若干（65535）个线程块，每个线程块包含若干（512）个线程，三者的关系如下图：

Thread，block，grid是CUDA编程上的概念，为了方便程序员软件设计，组织线程。

thread：一个CUDA的并行程序会被以许多个threads来执行

block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信

grid：多个blocks则会再构成grid

线程索引的计算公式

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。

CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化，这里给出Grid和Block不同划分方式下线程索引ID的计算公式。

1、 grid划分成1维，block划分为1维

int threadId = blockIdx.x *blockDim.x + threadIdx.x;

2、 grid划分成1维，block划分为2维

int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;

3、 grid划分成1维，block划分为3维

int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z

+ threadIdx.z * blockDim.y * blockDim.x

+ threadIdx.y * blockDim.x + threadIdx.x;

4、 grid划分成2维，block划分为1维

int blockId = blockIdx.y * gridDim.x + blockIdx.x;

int threadId = blockId * blockDim.x + threadIdx.x;

5、 grid划分成2维，block划分为2维

int blockId = blockIdx.x + blockIdx.y * gridDim.x;

int threadId = blockId * (blockDim.x * blockDim.y)

+ (threadIdx.y * blockDim.x) + threadIdx.x;

6、 grid划分成2维，block划分为3维

int blockId = blockIdx.x + blockIdx.y * gridDim.x;

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

+ (threadIdx.z * (blockDim.x * blockDim.y))

+ (threadIdx.y * blockDim.x) + threadIdx.x;

7、 grid划分成3维，block划分为1维

int blockId = blockIdx.x + blockIdx.y * gridDim.x

+ gridDim.x * gridDim.y * blockIdx.z;

int threadId = blockId * blockDim.x + threadIdx.x;

8、 grid划分成3维，block划分为2维

int blockId = blockIdx.x + blockIdx.y * gridDim.x

+ gridDim.x * gridDim.y * blockIdx.z;

int threadId = blockId * (blockDim.x * blockDim.y)

+ (threadIdx.y * blockDim.x) + threadIdx.x;

9、 grid划分成3维，block划分为3维

int blockId = blockIdx.x + blockIdx.y * gridDim.x

+ gridDim.x * gridDim.y * blockIdx.z;

int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)

+ (threadIdx.z * (blockDim.x * blockDim.y))

+ (threadIdx.y * blockDim.x) + threadIdx.x;

dim3数据结构

dim3是CUDA 中一个比较特殊的数据结构，我们可以用这个数据结构创建一个二维的线程块与线程网格。例如在长方形布局的方式中，每个线程块的X 轴方向上开启了32 个线程， Y轴方向上开启了4 个线程。在线程网格上， X 轴方向上有1 个线程块， Y 轴方向有4 个线程块。

dim3 threads_rect(32,4)
dim3 blocks_rect(1,4)

dim3 blocks(DIM/16,DIM/16);

dim3 threads(16,16);

kernel<<<blocks,threads>>>( d->dev_bitmap, ticks );

CUDA学习笔记