CUDA 及其 golang 調(diào)用 - 從入門到放棄 - 3. 真·向量內(nèi)積的盡頭

經(jīng)過上一回的努力雳锋，我們終于將 GPU 計(jì)算的時(shí)間縮減到同 CPU 一個(gè)數(shù)量級黄绩，但是發(fā)現(xiàn)內(nèi)存和顯存之間的 memcpy 成了最主要的性能損耗。

一玷过、固定內(nèi)存

這里有一個(gè)驚人的事實(shí)爽丹，從分頁的虛擬內(nèi)存中的地址復(fù)制到顯存，需要經(jīng)過兩次復(fù)制习劫，中間需要復(fù)制到非分頁的固定內(nèi)存地址诽里。同時(shí)，CUDA 還提供了 cudaMallocHost 以分配固定內(nèi)存灸眼，我們可以申請這樣的內(nèi)存并將其地址賦予 golang 中的 slice焰宣，在 golang 端直接寫這個(gè)地址捕仔，將其拷貝到顯存以達(dá)到縮減 memcpy 時(shí)間的目的。

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 版進(jìn)一步下降到約 174ms帆调，其中 cudaMemcpy 下降到約 134ms

二番刊、內(nèi)存映射

CUDA 還提供了避免拷貝的內(nèi)存映射影锈，使用 cudaHostAlloc 分配固定內(nèi)存并映射到顯存，使核函數(shù)能夠訪問

c 端有改動(dòng)的代碼如下：

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 版進(jìn)一步下降到約 156ms，其中 cudaMemcpy 已不再需要佛吓，而 devDot 卻上升到約 125ms，因?yàn)楹驮L問顯存相比维雇，訪問內(nèi)存產(chǎn)生嚴(yán)重的性能損耗

三、統(tǒng)一尋址

CUDA 6.0 開始提供統(tǒng)一尋址機(jī)制逸贾，使用 cudaMallocManaged 分配統(tǒng)一尋址的內(nèi)存铝侵，透明地在內(nèi)存和顯存間進(jìn)行映像触徐，兩邊都能訪問

c 端有改動(dòng)的代碼如下：

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 版進(jìn)一步下降到約 86ms，其中 devDot 約 9.577ms温自，與 CPU 版的比值約為 7.98%皇钞，比上一篇稍慢一點(diǎn)。cudaMemcpy 也不存在了馆里，而主要的性能損耗變成了之前毫不顯眼的 cudaLaunchKernel 等 CUDA 運(yùn)行時(shí)鸠踪。

Licensed under CC BY-SA 4.0

最后編輯于 ：
?著作權(quán)歸作者所有,轉(zhuǎn)載或內(nèi)容合作請聯(lián)系作者
  • 人面猴
    序言：七十年代末营密，一起剝皮案震驚了整個(gè)濱河市目锭，隨后出現(xiàn)的幾起案子纷捞，更是在濱河造成了極大的恐慌主儡，老刑警劉巖惨缆，帶你破解...
    沈念sama閱讀 206,968評論 6贊 482
  • 死咒
    序言：濱河連續(xù)發(fā)生了三起死亡事件踪央，死亡現(xiàn)場離奇詭異，居然都是意外死亡健无，警方通過查閱死者的電腦和手機(jī)液斜，發(fā)現(xiàn)死者居然都...
    沈念sama閱讀 88,601評論 2贊 382
  • 救了他兩次的神仙讓他今天三更去死
    文/潘曉璐 我一進(jìn)店門少漆，熙熙樓的掌柜王于貴愁眉苦臉地迎上來，“玉大人渗磅，你說我怎么就攤上這事始鱼。” “怎么了医清？”我有些...
    開封第一講書人閱讀 153,220評論 0贊 344
  • 道士緝兇錄：失蹤的賣姜人
    文/不壞的土叔 我叫張陵会烙，是天一觀的道長柏腻。 經(jīng)常有香客問我系吭，道長，這世上最難降的妖魔是什么贫导？ 我笑而不...
    開封第一講書人閱讀 55,416評論 1贊 279
  • ?港島之戀（遺憾婚禮）
    正文 為了忘掉前任孩灯，我火速辦了婚禮逾滥，結(jié)果婚禮上，老公的妹妹穿的比我還像新娘讥巡。我一直安慰自己舔哪，他們只是感情好，可當(dāng)我...
    茶點(diǎn)故事閱讀 64,425評論 5贊 374
  • 惡毒庶女頂嫁案：這布局不是一般人想出來的
    文/花漫 我一把揭開白布抬驴。 她就那樣靜靜地躺著布持，像睡著了一般陕悬。 火紅的嫁衣襯著肌膚如雪。 梳的紋絲不亂的頭發(fā)上捉超，一...
    開封第一講書人閱讀 49,144評論 1贊 285
  • 城市分裂傳說
    那天狂秦，我揣著相機(jī)與錄音，去河邊找鬼侧啼。 笑死堪簿，一個(gè)胖子當(dāng)著我的面吹牛，可吹牛的內(nèi)容都是我干的哪审。 我是一名探鬼主播虑瀑，決...
    沈念sama閱讀 38,432評論 3贊 401
  • 雙鴛鴦連環(huán)套：你想象不到人心有多黑
    文/蒼蘭香墨 我猛地睜開眼，長吁一口氣：“原來是場噩夢啊……” “哼叽奥！你這毒婦竟也來了？” 一聲冷哼從身側(cè)響起魔市，我...
    開封第一講書人閱讀 37,088評論 0贊 261
  • 萬榮殺人案實(shí)錄
    序言：老撾萬榮一對情侶失蹤待德，失蹤者是張志新（化名）和其女友劉穎枫夺，沒想到半個(gè)月后，有當(dāng)?shù)厝嗽跇淞掷锇l(fā)現(xiàn)了一具尸體涧偷，經(jīng)...
    沈念sama閱讀 43,586評論 1贊 300
  • ?護(hù)林員之死
    正文 獨(dú)居荒郊野嶺守林人離奇死亡燎潮，尸身上長有42處帶血的膿包…… 初始之章·張勛 以下內(nèi)容為張勛視角 年9月15日...
    茶點(diǎn)故事閱讀 36,028評論 2贊 325
  • ?白月光啟示錄
    正文 我和宋清朗相戀三年确封，在試婚紗的時(shí)候發(fā)現(xiàn)自己被綠了再菊。 大學(xué)時(shí)的朋友給我發(fā)了我未婚夫和他白月光在一起吃飯的照片。...
    茶點(diǎn)故事閱讀 38,137評論 1贊 334
  • 活死人
    序言：一個(gè)原本活蹦亂跳的男人離奇死亡秉剑，死狀恐怖稠诲，靈堂內(nèi)的尸體忽然破棺而出，到底是詐尸還是另有隱情略水，我是刑警寧澤劝萤，帶...
    沈念sama閱讀 33,783評論 4贊 324
  • ?日本核電站爆炸內(nèi)幕
    正文 年R本政府宣布，位于F島的核電站跨释，受9級特大地震影響，放射性物質(zhì)發(fā)生泄漏盖文。R本人自食惡果不足惜蚯姆，卻給世界環(huán)境...
    茶點(diǎn)故事閱讀 39,343評論 3贊 307
  • 男人毒藥：我在死后第九天來索命
    文/蒙蒙 一龄恋、第九天 我趴在偏房一處隱蔽的房頂上張望凶伙。 院中可真熱鬧，春花似錦显押、人聲如沸傻挂。這莊子的主人今日做“春日...
    開封第一講書人閱讀 30,333評論 0贊 19
  • 一樁弒父案兽肤，背后竟有這般陰謀
    文/蒼蘭香墨 我抬頭看了看天上的太陽绪抛。三九已至，卻和暖如春笤休，著一層夾襖步出監(jiān)牢的瞬間症副，已是汗流浹背。 一陣腳步聲響...
    開封第一講書人閱讀 31,559評論 1贊 262
  • 情欲美人皮
    我被黑心中介騙來泰國打工底洗， 沒想到剛下飛機(jī)就差點(diǎn)兒被人妖公主榨干…… 1. 我叫王不留咕娄，地道東北人。 一個(gè)月前我還...
    沈念sama閱讀 45,595評論 2贊 355
  • 代替公主和親
    正文 我出身青樓费变，卻偏偏與公主長得像，于是被迫代替她去往敵國和親扛稽。 傳聞我的和親對象是個(gè)殘疾皇子滑负，可洞房花燭夜當(dāng)晚...
    茶點(diǎn)故事閱讀 42,901評論 2贊 345