核函数
- 核函数是cuda编程关键
- 通过创建.cu创建cudac程序文件,然后交给nvcc编译
- 加上__global__前缀的函数由host端调用,__device修饰的函数为device函数,由设备调用
- __host修饰的函数为host函数,由设备调用
- host调用核函数的方式是function<<<gridDim,blockDim,sharedMemorySize,stream>>>(args,…)
Cuda线程&线程id
CUDA里面用了grid和block作为线程的组织单位,一个grid可以包含多个block,一个block包含多个thread,grid和block都是三维向量,下段代码展示了如何定义一个grid和block:
dim3 grid(x.y,z);
dim3 block(x,y,z);
在计算线程索引的时候,经常会看到gridDim、blockDim、blockIdx、threadIdx。关于这些名词的解释如下:
- gridDim:block在gird里面的数量维度,
- blockDIm: threads在blcok中的维度
- blockIdx:block在grid中的索引
- threadIdx: thread在block中的索引
计算线程坐标
怎么从那么多的线程中找到目标线程,这就需要计算线程坐标,可以拆成以下三步去计算:
- 获取thread在block中的位置
- 获取block在grid中的位置
- 计算block有多少线程,求解位置索引
计算thread在block中的位置
t
h
r
e
a
d
I
n
B
l
o
c
k
=
t
h
r
e
a
I
d
x
.
x
+
t
h
r
e
a
d
I
d
x
.
y
∗
b
l
o
c
k
D
i
m
.
x
+
t
h
r
e
a
d
I
d
x
.
z
∗
b
l
o
c
k
D
i
m
.
x
∗
b
l
o
c
k
D
i
m
.
y
threadInBlock=threaIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y
threadInBlock=threaIdx.x+threadIdx.y∗blockDim.x+threadIdx.z∗blockDim.x∗blockDim.y
一开始我并不理解这个计算公式,后来为了加强记忆,理解成下面这样(不对的话请指证):
在三维block中计算线程的一维索引时,可以想像成将三维数据展平成一维,按照“优先顺序”的话,x维度是变化最快的,其次是y维度,最后是z维度。因此计算方式是:
- threadIdx.x 直接加到索引上,因为它是最内层的循环(X 维度)。
- threadIdx.y * blockDim.x 是因为每增加一个 Y 索引,你跳过了整个 X 维度的线程数。
- threadIdx.z * blockDim.x * blockDim.y 是因为每增加一个 Z 索引,你跳过了整个 X-Y 平面的线程数
计算block在grid中的位置
b l o c k I n G r i d = b l o c k I d x . x + b l o c k I d x . y ∗ g i r d D i m . x + b l o c k I d x . z ∗ g i r d D i m . x ∗ g r i d D i m . y blockInGrid=blockIdx.x+blockIdx.y*girdDim.x+blockIdx.z*girdDim.x*gridDim.y blockInGrid=blockIdx.x+blockIdx.y∗girdDim.x+blockIdx.z∗girdDim.x∗gridDim.y
求解全局idx:
一个block中的thread总数为:
p
e
r
B
l
o
c
k
S
I
z
e
=
b
l
o
c
k
D
i
m
.
x
∗
b
l
o
c
k
D
i
m
.
y
perBlockSIze=blockDim.x*blockDim.y
perBlockSIze=blockDim.x∗blockDim.y
因此
i
d
x
=
t
h
r
e
a
d
I
n
B
l
o
c
k
+
b
l
o
c
k
I
n
g
r
i
d
∗
p
e
r
B
l
o
c
k
S
i
z
e
idx=threadInBlock+blockIngrid*perBlockSize
idx=threadInBlock+blockIngrid∗perBlockSize
上式可以理解为:每个block有perBlockSize个thread,对应线程所在的block起始位置就是blockInGrid*perperBlockSize,然后加上偏移就是threadInblock