上一回真是一点牌面都没有，所以这里做出一些优化。
优化一：将 cudaMalloc 申请的显存地址保存在上下文里重复利用
优化二：启用多线程
const size_t NTB = 256;
const size_t EXT = 8;
#define divCeil(a, b) (((a) + (b) - 1) / (b))
struct Ctx {
float *xd, *yd, *rd;
size_t n;
};
extern "C" __declspec(dllexport) void init(Ctx **p, size_t n) {
Ctx *ctx = (Ctx *)malloc(sizeof(Ctx));
ctx->n = n;
size_t sz = sizeof(float) * n;
cudaMalloc(&(ctx->xd), sz);
cudaMalloc(&(ctx->yd), sz);
cudaMallocManaged(&(ctx->rd), sizeof(float) * divCeil(n, NTB) / EXT);
*p = ctx;
}
extern "C" __declspec(dllexport) void deinit(Ctx *ctx) {
cudaFree(ctx->xd);
cudaFree(ctx->yd);
cudaFree(ctx->rd);
free(ctx);
}
__global__ void devDot(float *x, float *y, size_t n, float *r) {
__shared__ float rb[NTB];
size_t itb = threadIdx.x;
size_t i = blockIdx.x * blockDim.x * EXT + itb;
float s = 0.0;
for (size_t j = 0; j < EXT && i < n; j++, i += blockDim.x) {
s += x[i] * y[i];
}
rb[itb] = s;
__syncthreads();
for (size_t i = NTB >> 1; i != 0; i >>= 1) {
if (itb < i) rb[itb] += rb[itb + i];
__syncthreads();
}
if (0 == itb) r[blockIdx.x] = rb[0];
}
extern "C" __declspec(dllexport) void dot(Ctx *ctx, float *x, float *y, float *r) {
size_t sz = sizeof(float) * ctx->n;
cudaMemcpy(ctx->xd, x, sz, cudaMemcpyHostToDevice);
cudaMemcpy(ctx->yd, y, sz, cudaMemcpyHostToDevice);
size_t nb = divCeil(ctx->n, NTB) / EXT;
float *rd = ctx->rd;
devDot<<<nb, NTB>>>(ctx->xd, ctx->yd, ctx->n, rd);
cudaDeviceSynchronize();
float s = 0.0;
for (size_t i = 0; i < nb; i++) s += rd[i];
*r = s;
}
GPU 的多线程和 CPU 的多线程是两回事。逻辑上分为 grid, block, thread 三层结构，在 GPU 函数调用处的 <<<m, n>>> 中指定整个 grid 包含 m 个 block，每个 block 包含的 n 个 thread。
在核函数中，gridDim.x 和 blockDim.x 为 grid 包含的 block 数和 block 包含的 thread 数，blockIdx.x 和 threadIdx.x 为 block 的序号和 thread 在 block 中的序号，__shared__ 指定局部变量在同一 block 中的线程间共享。这里，我们计算出每个线程对向量负责计算的范围，并行地求和（第一级）并放进共享的数组，再将共享数组中的值并行地累加（第二级），注意有两处需要调用 __syncthreads 进行 block 内所有线程的同步。最后在核函数中做第三级的累加。
因为是在并发的环境中，我们不能再用单个变量去承载整个累加的操作。在 GPU 中触犯并发的竞争问题，会让你死得比在 CPU 中惨烈得多，比如我的 GTX1050 是 640 核。
下面是 golang 的部分：
package main
import (
"math/rand"
"syscall"
"time"
"unsafe"
)
const N = 1 << 20
type Lib struct {
dll *syscall.DLL
deinitProc *syscall.Proc
dotProc *syscall.Proc
handler uintptr
}
func LoadLib() (*Lib, error) {
l := &Lib{}
var err error
defer func() {
if nil != err {
l.Release()
}
}()
if l.dll, err = syscall.LoadDLL("cuda.dll"); nil != err {
return nil, err
}
if l.deinitProc, err = l.dll.FindProc("deinit"); nil != err {
return nil, err
}
if l.dotProc, err = l.dll.FindProc("dot"); nil != err {
return nil, err
}
proc, err := l.dll.FindProc("init")
if nil != err {
return nil, err
}
proc.Call(uintptr(unsafe.Pointer(&l.handler)), uintptr(N))
return l, nil
}
func (l *Lib) Release() {
if nil != l.deinitProc && 0 != l.handler {
l.deinitProc.Call(l.handler)
}
if nil != l.dll {
l.dll.Release()
}
}
func (l *Lib) Dot(x, y []float32) float32 {
var r float32
l.dotProc.Call(
l.handler,
uintptr(unsafe.Pointer(&x[0])),
uintptr(unsafe.Pointer(&y[0])),
uintptr(unsafe.Pointer(&r)),
)
return r
}
func main() {
lib, err := LoadLib()
if nil != err {
println(err.Error())
return
}
defer lib.Release()
rand.Seed(time.Now().Unix())
x, y := make([]float32, N), make([]float32, N)
for i := 0; i < N; i++ {
x[i], y[i] = rand.Float32(), rand.Float32()
}
t := time.Now()
var r float32
for i := 0; i < 100; i++ {
r = 0
for i := 0; i < N; i++ {
r += x[i] * y[i]
}
}
println(time.Now().Sub(t).Microseconds())
println(r)
t = time.Now()
for i := 0; i < 100; i++ {
r = lib.Dot(x, y)
}
println(time.Now().Sub(t).Microseconds())
println(r)
}
仍然使用 nvprof 观察，在其中一次运行中，CPU 版计算 100 次耗时仍为约 120ms，而 GPU 版约 357ms，是上一次的十分之一不到啊！其中：
-
cudaMemcpy约 277ms，必然和上一次基本不变，而cudaMalloc被优化出了这 100 次循环中 -
cudaDeviceSynchronize约 47ms，其中：-
devDot约 8.661ms，性能提升三个数量级！与 CPU 版的比值约为 7.21%
-
可以看到，现在内存和显存之间的 memcpy 成了最主要的性能损耗！不知道后续有没有办法优化，这是否已来到在此环境下向量内积运算性能的尽头？
Licensed under CC BY-SA 4.0
