CUDA 及其 golang 调用 - 从入门到放弃 - 3. 真·向量内积的尽头

经过上一回的努力，我们终于将 GPU 计算的时间缩减到同 CPU 一个数量级，但是发现内存和显存之间的 memcpy 成了最主要的性能损耗。

一、固定内存

这里有一个惊人的事实，从分页的虚拟内存中的地址复制到显存，需要经过两次复制，中间需要复制到非分页的固定内存地址。同时，CUDA 还提供了 cudaMallocHost 以分配固定内存，我们可以申请这样的内存并将其地址赋予 golang 中的 slice，在 golang 端直接写这个地址，将其拷贝到显存以达到缩减 memcpy 时间的目的。

const size_t NTB = 256;
const size_t EXT = 8;
#define divCeil(a, b) (((a) + (b) - 1) / (b))

struct Ctx {
    float *x, *y;
    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;
    cudaMallocHost(&(ctx->x), sz);
    cudaMallocHost(&(ctx->y), sz);
    cudaMalloc(&(ctx->xd), sz);
    cudaMalloc(&(ctx->yd), sz);
    cudaMallocManaged(&(ctx->rd), sizeof(float) * divCeil(n, NTB) / EXT);
    *p = ctx;
}

extern "C" __declspec(dllexport) void getInputs(Ctx *ctx, float **px, float **py) {
    *px = ctx->x;
    *py = ctx->y;
}

extern "C" __declspec(dllexport) void deinit(Ctx *ctx) {
    cudaFreeHost(ctx->x);
    cudaFreeHost(ctx->y);
    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 *r) {
    size_t sz = sizeof(float) * ctx->n;
    cudaMemcpy(ctx->xd, ctx->x, sz, cudaMemcpyHostToDevice);
    cudaMemcpy(ctx->yd, ctx->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;
}

package main

import (
    "math/rand"
    "reflect"
    "syscall"
    "time"
    "unsafe"
)

const N = 1 << 20

type Lib struct {
    dll        *syscall.DLL
    deinitProc *syscall.Proc
    dotProc    *syscall.Proc
    handler    uintptr
    X, Y       []float32
}

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))

    proc, err = l.dll.FindProc("getInputs")
    if nil != err {
        return nil, err
    }
    xh := (*reflect.SliceHeader)(unsafe.Pointer(&l.X))
    yh := (*reflect.SliceHeader)(unsafe.Pointer(&l.Y))
    xh.Len, xh.Cap, yh.Len, yh.Cap = N, N, N, N
    proc.Call(l.handler,
        uintptr(unsafe.Pointer(&xh.Data)), uintptr(unsafe.Pointer(&yh.Data)))

    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() float32 {
    var r float32
    l.dotProc.Call(l.handler, 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 := lib.X, lib.Y
    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()
    }
    println(time.Now().Sub(t).Microseconds())
    println(r)
}

CPU 版仍约 120ms，GPU 版进一步下降到约 174ms，其中 cudaMemcpy 下降到约 134ms

二、内存映射

CUDA 还提供了避免拷贝的内存映射，使用 cudaHostAlloc 分配固定内存并映射到显存，使核函数能够访问

c 端有改动的代码如下：

struct Ctx {
    float *x, *y, *r;
    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;
    cudaHostAlloc(&(ctx->x), sz, cudaHostAllocMapped);
    cudaHostAlloc(&(ctx->y), sz, cudaHostAllocMapped);
    cudaMallocManaged(&(ctx->r), sizeof(float) * divCeil(n, NTB) / EXT);
    *p = ctx;
}

extern "C" __declspec(dllexport) void deinit(Ctx *ctx) {
    cudaFreeHost(ctx->x);
    cudaFreeHost(ctx->y);
    cudaFree(ctx->r);
    free(ctx);
}

extern "C" __declspec(dllexport) void dot(Ctx *ctx, float *r) {
    size_t nb = divCeil(ctx->n, NTB) / EXT;
    float *rd = ctx->r;
    devDot<<<nb, NTB>>>(ctx->x, ctx->y, ctx->n, rd);
    cudaDeviceSynchronize();
    float s = 0.0;
    for (size_t i = 0; i < nb; i++) s += rd[i];
    *r = s;
}

CPU 版仍约 120ms，GPU 版进一步下降到约 156ms，其中 cudaMemcpy 已不再需要，而 devDot 却上升到约 125ms，因为和访问显存相比，访问内存产生严重的性能损耗

三、统一寻址

CUDA 6.0 开始提供统一寻址机制，使用 cudaMallocManaged 分配统一寻址的内存，透明地在内存和显存间进行映像，两边都能访问

c 端有改动的代码如下：

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;
    cudaMallocManaged(&(ctx->x), sz);
    cudaMallocManaged(&(ctx->y), sz);
    cudaMallocManaged(&(ctx->r), sizeof(float) * divCeil(n, NTB) / EXT);
    *p = ctx;
}

extern "C" __declspec(dllexport) void deinit(Ctx *ctx) {
    cudaFree(ctx->x);
    cudaFree(ctx->y);
    cudaFree(ctx->r);
    free(ctx);
}

喜讯！喜讯！GPU 终于赢了！CPU 版仍约 120ms，GPU 版进一步下降到约 86ms，其中 devDot 约 9.577ms，与 CPU 版的比值约为 7.98%，比上一篇稍慢一点。cudaMemcpy 也不存在了，而主要的性能损耗变成了之前毫不显眼的 cudaLaunchKernel 等 CUDA 运行时。

Licensed under CC BY-SA 4.0

最后编辑于 ：
©著作权归作者所有,转载或内容合作请联系作者
  • 人面猴
    序言：七十年代末，一起剥皮案震惊了整个滨河市，随后出现的几起案子，更是在滨河造成了极大的恐慌，老刑警刘岩，带你破解...
    沈念sama阅读 194,491评论 5赞 459
  • 死咒
    序言：滨河连续发生了三起死亡事件，死亡现场离奇诡异，居然都是意外死亡，警方通过查阅死者的电脑和手机，发现死者居然都...
    沈念sama阅读 81,856评论 2赞 371
  • 救了他两次的神仙让他今天三更去死
    文/潘晓璐 我一进店门，熙熙楼的掌柜王于贵愁眉苦脸地迎上来，“玉大人，你说我怎么就摊上这事。” “怎么了？”我有些...
    开封第一讲书人阅读 141,745评论 0赞 319
  • 道士缉凶录：失踪的卖姜人
    文/不坏的土叔 我叫张陵，是天一观的道长。 经常有香客问我，道长，这世上最难降的妖魔是什么？ 我笑而不...
    开封第一讲书人阅读 52,196评论 1赞 263
  • 港岛之恋（遗憾婚礼）
    正文 为了忘掉前任，我火速办了婚礼，结果婚礼上，老公的妹妹穿的比我还像新娘。我一直安慰自己，他们只是感情好，可当我...
    茶点故事阅读 61,073评论 4赞 355
  • 恶毒庶女顶嫁案：这布局不是一般人想出来的
    文/花漫 我一把揭开白布。 她就那样静静地躺着，像睡着了一般。 火红的嫁衣衬着肌肤如雪。 梳的纹丝不乱的头发上，一...
    开封第一讲书人阅读 46,112评论 1赞 272
  • 城市分裂传说
    那天，我揣着相机与录音，去河边找鬼。 笑死，一个胖子当着我的面吹牛，可吹牛的内容都是我干的。 我是一名探鬼主播，决...
    沈念sama阅读 36,531评论 3赞 381
  • 双鸳鸯连环套：你想象不到人心有多黑
    文/苍兰香墨 我猛地睁开眼，长吁一口气：“原来是场噩梦啊……” “哼！你这毒妇竟也来了？” 一声冷哼从身侧响起，我...
    开封第一讲书人阅读 35,215评论 0赞 253
  • 万荣杀人案实录
    序言：老挝万荣一对情侣失踪，失踪者是张志新（化名）和其女友刘颖，没想到半个月后，有当地人在树林里发现了一具尸体，经...
    沈念sama阅读 39,485评论 1赞 290
  • 护林员之死
    正文 独居荒郊野岭守林人离奇死亡，尸身上长有42处带血的脓包…… 初始之章·张勋 以下内容为张勋视角 年9月15日...
    茶点故事阅读 34,578评论 2赞 309
  • 白月光启示录
    正文 我和宋清朗相恋三年，在试婚纱的时候发现自己被绿了。 大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
    茶点故事阅读 36,356评论 1赞 326
  • 活死人
    序言：一个原本活蹦乱跳的男人离奇死亡，死状恐怖，灵堂内的尸体忽然破棺而出，到底是诈尸还是另有隐情，我是刑警宁泽，带...
    沈念sama阅读 32,215评论 3赞 312
  • 日本核电站爆炸内幕
    正文 年R本政府宣布，位于F岛的核电站，受9级特大地震影响，放射性物质发生泄漏。R本人自食恶果不足惜，却给世界环境...
    茶点故事阅读 37,583评论 3赞 299
  • 男人毒药：我在死后第九天来索命
    文/蒙蒙 一、第九天 我趴在偏房一处隐蔽的房顶上张望。 院中可真热闹，春花似锦、人声如沸。这庄子的主人今日做“春日...
    开封第一讲书人阅读 28,898评论 0赞 17
  • 一桩弑父案，背后竟有这般阴谋
    文/苍兰香墨 我抬头看了看天上的太阳。三九已至，却和暖如春，着一层夹袄步出监牢的瞬间，已是汗流浃背。 一阵脚步声响...
    开封第一讲书人阅读 30,174评论 1赞 250
  • 情欲美人皮
    我被黑心中介骗来泰国打工， 没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留，地道东北人。 一个月前我还...
    沈念sama阅读 41,497评论 2赞 341
  • 代替公主和亲
    正文 我出身青楼，却偏偏与公主长得像，于是被迫代替她去往敌国和亲。 传闻我的和亲对象是个残疾皇子，可洞房花烛夜当晚...
    茶点故事阅读 40,697评论 2赞 335