上一回真是一點(diǎn)牌面都沒有，所以這里做出一些優(yōu)化。

優(yōu)化一：將 cudaMalloc 申請(qǐng)的顯存地址保存在上下文里重復(fù)利用

優(yōu)化二：?jiǎn)⒂枚嗑€程

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 三層結(jié)構(gòu)，在 GPU 函數(shù)調(diào)用處的 <<<m, n>>> 中指定整個(gè) grid 包含 m 個(gè) block，每個(gè) block 包含的 n 個(gè) thread。

在核函數(shù)中，gridDim.x 和 blockDim.x 為 grid 包含的 block 數(shù)和 block 包含的 thread 數(shù)，blockIdx.x 和 threadIdx.x 為 block 的序號(hào)和 thread 在 block 中的序號(hào)，__shared__ 指定局部變量在同一 block 中的線程間共享。這里，我們計(jì)算出每個(gè)線程對(duì)向量負(fù)責(zé)計(jì)算的范圍，并行地求和（第一級(jí)）并放進(jìn)共享的數(shù)組，再將共享數(shù)組中的值并行地累加（第二級(jí)），注意有兩處需要調(diào)用 __syncthreads 進(jìn)行 block 內(nèi)所有線程的同步。最后在核函數(shù)中做第三級(jí)的累加。

因?yàn)槭窃诓l(fā)的環(huán)境中，我們不能再用單個(gè)變量去承載整個(gè)累加的操作。在 GPU 中觸犯并發(fā)的競(jìng)爭(zhēng)問題，會(huì)讓你死得比在 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 觀察，在其中一次運(yùn)行中，CPU 版計(jì)算 100 次耗時(shí)仍為約 120ms，而 GPU 版約 357ms，是上一次的十分之一不到??！其中：

• cudaMemcpy 約 277ms，必然和上一次基本不變，而 cudaMalloc 被優(yōu)化出了這 100 次循環(huán)中
• cudaDeviceSynchronize 約 47ms，其中：
  • devDot 約 8.661ms，性能提升三個(gè)數(shù)量級(jí)！與 CPU 版的比值約為 7.21%

可以看到，現(xiàn)在內(nèi)存和顯存之間的 memcpy 成了最主要的性能損耗！不知道后續(xù)有沒有辦法優(yōu)化，這是否已來到在此環(huán)境下向量?jī)?nèi)積運(yùn)算性能的盡頭？

Licensed under CC BY-SA 4.0