上周幫同事排查一個(gè)CUDA kernel時(shí)靈時(shí)不靈的問(wèn)題。同一個(gè)block里寫(xiě)全局內(nèi)存另一個(gè)block輪詢flag按理說(shuō)是很常見(jiàn)的生產(chǎn)者-消費(fèi)者模式結(jié)果在A卡上能跑在RTX 4090上偶發(fā)卡死。折騰了兩天最后問(wèn)題落到了內(nèi)存柵欄函數(shù)上——準(zhǔn)確地說(shuō)是缺少__threadfence()這顆“定心丸”。這篇對(duì)應(yīng)CUDA C編程指南語(yǔ)言擴(kuò)展部分的內(nèi)存柵欄函數(shù)7.5節(jié)與同步函數(shù)第6章兩兄弟我合到一期講。很多寫(xiě)CUDA的C程序員調(diào)度kernel、優(yōu)化shared memory都挺順手但對(duì)__syncthreads和__threadfence的理解停留在“加上就對(duì)了”的層面。這篇不打算逐條翻譯手冊(cè)而是把這兩類函數(shù)的語(yǔ)義、適用邊界、死鎖陷阱以及編程指南第6版以后新增的同步原語(yǔ)一次性講清楚。適合誰(shuí)看?已經(jīng)寫(xiě)過(guò)基本kernel、想搞明白跨線程通信為什么需要柵欄、以及被__syncthreads死鎖坑過(guò)的人。1. GPU弱內(nèi)存模型為什么跨線程通信必須手動(dòng)加?xùn)艡?.1 先接受一個(gè)事實(shí)GPU不會(huì)主動(dòng)幫你排序大多數(shù)CPU程序員轉(zhuǎn)CUDA時(shí)腦子里帶著一套x86的經(jīng)驗(yàn)寫(xiě)一個(gè)變量再寫(xiě)一個(gè)flag另一個(gè)線程看到flag為1就一定能看到前面的變量值。這個(gè)經(jīng)驗(yàn)在x86上大體成立因?yàn)閤86用的是TSOTotal Store Order內(nèi)存模型store是按順序排著隊(duì)對(duì)外可見(jiàn)的緩存一致性也做得極強(qiáng)。GPU不一樣。CUDA官方文檔里明確說(shuō)GPU是弱內(nèi)存模型weak memory model。弱體現(xiàn)在三個(gè)層面第一編譯器重排。C編譯器默認(rèn)不知道你的代碼里有跨線程通信它認(rèn)為普通變量的讀寫(xiě)只作用于當(dāng)前線程。于是它可以自由地把*data 42和*flag 1換順序甚至把輪詢循環(huán)里的*flag讀到寄存器里緩存起來(lái)。第二硬件執(zhí)行亂序。一個(gè)SM內(nèi)部warp調(diào)度器會(huì)交錯(cuò)的發(fā)射多條指令同一線程的不同指令之間只要沒(méi)有數(shù)據(jù)依賴就可能亂序執(zhí)行。即使按順序發(fā)射了store的結(jié)果也會(huì)先停留在流水線里不會(huì)立刻變成別的線程能看到的狀態(tài)。第三緩存可見(jiàn)性延遲。每個(gè)SM有自己的L1 cache全局內(nèi)存的寫(xiě)入要先寫(xiě)到L1/L2路徑上最終到L2才可能被其他SM看到。寫(xiě)入什么時(shí)候“抵達(dá)”L2沒(méi)有統(tǒng)一的時(shí)間點(diǎn)保證。所以“寫(xiě)一個(gè)flag別人就能看到”這個(gè)樸素想法在GPU上大概率翻車。你得自己顯式地告訴硬件這組內(nèi)存操作需要按什么順序、在多大范圍內(nèi)對(duì)外可見(jiàn)。這就是內(nèi)存柵欄函數(shù)存在的意義。1.2 必掛的實(shí)驗(yàn)無(wú)柵欄的塊間flag通信我直接給一個(gè)最簡(jiǎn)復(fù)現(xiàn)__global__ void flag_test(int* data, int* flag) { if (blockIdx.x 0) { if (threadIdx.x 0) { *data 42; // 普通store *flag 1; // 普通store } } else if (blockIdx.x 1) { if (threadIdx.x 0) { while (*flag ! 1); // 輪詢等待 printf(data%d\n, *data); } } }邏輯上應(yīng)該輸出data42實(shí)際跑起來(lái)你可能會(huì)遇到三種情況卡死。block 1在while里永遠(yuǎn)出不來(lái)。原因可能是*flag被編譯器讀到寄存器里緩存了也可能是flag的寫(xiě)一直沒(méi)刷新到block 1所在SM可見(jiàn)的位置。輸出data0或未初始化值。說(shuō)明flag先于data被看到兩個(gè)store的順序被重排了。一切正常。恭喜你但那是運(yùn)氣換個(gè)架構(gòu)或調(diào)個(gè)編譯器優(yōu)化級(jí)別就翻車。這個(gè)demo看起來(lái)平平無(wú)奇但它精準(zhǔn)踩中了兩個(gè)坑編譯器重排硬件弱可見(jiàn)性。后面我們會(huì)用柵欄和原子操作把這兩個(gè)坑都堵上。1.3 柵欄與原子各管一段可見(jiàn)性的三塊基石想徹底搞懂fence得先明白“讓一次跨線程通信成立”需要哪三塊基石寫(xiě)入方要把自己的寫(xiě)操作按順序提交出去。寫(xiě)操作是否已經(jīng)離開(kāi)執(zhí)行管線。柵欄函數(shù)管這一塊。寫(xiě)入要到達(dá)另一個(gè)線程能看見(jiàn)的緩存層級(jí)。全局內(nèi)存至少要刷到L2原子操作和部分fence會(huì)push系統(tǒng)做到這一點(diǎn)。讀取方不能把讀操作緩存到寄存器。需要用volatile或原子讀防止編譯器把循環(huán)讀優(yōu)化成只讀一次。三者缺一不可。很多人以為_(kāi)_threadfence()是萬(wàn)能同步其實(shí)它只是第一塊基石。原子操作是第二塊volatile或atomic load是第三塊。理解這個(gè)分工是后面所有內(nèi)容的地基。2. 三種柵欄的語(yǔ)義與選型__threadfence_block / __threadfence / __threadfence_system2.1 柵欄是“提交點(diǎn)”不是“等待點(diǎn)”先修正一個(gè)常見(jiàn)誤解柵欄不是讓線程站在那等別的線程它不影響執(zhí)行流。柵欄的作用是給當(dāng)前線程的內(nèi)存操作劃定一個(gè)“提交點(diǎn)”——保證柵欄之前的所有內(nèi)存寫(xiě)入在柵欄之后的內(nèi)存操作開(kāi)始之前對(duì)指定范圍內(nèi)的觀察者可見(jiàn)。類比一下快遞普通store是你把包裹交給了快遞員快遞員什么時(shí)候送、送到哪一站你管不著。__threadfence()是在跟快遞員說(shuō)把之前交給你所有的包裹全部送到目的地物流站再回來(lái)。你不需要等快遞員回來(lái)但后續(xù)再交出去的包裹一定排在前面那批之后。注意柵欄只約束“當(dāng)前線程自身的訪問(wèn)順序”它不去等其他線程讀到什么也不保證其他線程馬上來(lái)讀。它和barrier徹底是兩回事混用就會(huì)出問(wèn)題這點(diǎn)第5節(jié)再展開(kāi)。2.2 作用范圍對(duì)比從線程塊到整個(gè)系統(tǒng)CUDA提供三個(gè)柵欄函數(shù)作用域從窄到寬函數(shù)作用范圍覆蓋的內(nèi)存典型場(chǎng)景__threadfence_block()當(dāng)前線程塊內(nèi)所有線程共享內(nèi)存 全局內(nèi)存同一block內(nèi)寫(xiě)共享內(nèi)存后再讀__threadfence()當(dāng)前設(shè)備上所有線程全局內(nèi)存block間通過(guò)全局內(nèi)存通信_(tái)_threadfence_system()設(shè)備 主機(jī)所有線程全局內(nèi)存 鎖頁(yè)主機(jī)內(nèi)存與鎖頁(yè)內(nèi)存交互、跨設(shè)備邊界選型原則很簡(jiǎn)單能用窄范圍就不用寬范圍。__threadfence_block()通常被編譯成輕量的內(nèi)存屏障指令基本不觸碰L2__threadfence()會(huì)強(qiáng)制L2層面的可見(jiàn)性代價(jià)高一個(gè)數(shù)量級(jí)__threadfence_system()最貴它要求設(shè)備內(nèi)存和主機(jī)鎖頁(yè)內(nèi)存在整個(gè)系統(tǒng)范圍內(nèi)可見(jiàn)通常意味著跨PCIe/驅(qū)動(dòng)層的同步開(kāi)銷。實(shí)戰(zhàn)里我見(jiàn)過(guò)不少人不管三七二十一所有通信一律__threadfence()。如果通信雙方本來(lái)就在同一個(gè)block內(nèi)每用一次全設(shè)備fence都是白給性能。2.3 經(jīng)典三段式寫(xiě)法store fence atomicExch正確的塊間flag通信業(yè)內(nèi)已經(jīng)形成了一套標(biāo)準(zhǔn)三段式。先上代碼__global__ void producer_consumer(int* data, int* flag) { if (blockIdx.x 0 threadIdx.x 0) { *data 42; __threadfence(); // 確保data寫(xiě)入對(duì)device所有線程可見(jiàn) atomicExch(flag, 1); // 原子寫(xiě)放行消費(fèi)者 } if (blockIdx.x 1 threadIdx.x 0) { while (atomicAdd(flag, 0) ! 1); // 原子讀避免寄存器緩存 printf(data%d\n, *data); } }這段代碼的每一步都有講究*data 42是普通store放在fence前面。它不需要原子因?yàn)樗灰蟆霸趂lag1之前data的寫(xiě)已經(jīng)被提交”。fence保證這一點(diǎn)。__threadfence()放在flag寫(xiě)入前。它把前面所有普通store強(qiáng)制提交到device作用域可見(jiàn)的位置。atomicExch(flag, 1)放在fence后。它本身是原子操作又帶有副作用編譯器不會(huì)把它和前面的普通store交換順序同時(shí)它作為“釋放鎖”的動(dòng)作標(biāo)志著producer完成了所有數(shù)據(jù)準(zhǔn)備。消費(fèi)者用atomicAdd(flag, 0)做原子讀。為什么不直接讀*flag因?yàn)槠胀ㄗx可能被優(yōu)化成寄存器緩存死循環(huán)。原子讀天然有副作用且int對(duì)齊的原子訪問(wèn)是硬件保證的。這套模式就是CUDA手寫(xiě)版的release/acquire協(xié)議。后面第4節(jié)我們會(huì)看到編程指南第6版引入了正式的內(nèi)存模型可以用cuda::atomic_ref把這三段式折疊成兩行語(yǔ)義更清晰。3. __syncthreads塊內(nèi)同步的邊界與死鎖陷阱3.1 同步了執(zhí)行順帶做了內(nèi)存柵欄__syncthreads()是block內(nèi)最常用的同步函數(shù)。它做兩件事第一執(zhí)行屏障block內(nèi)所有線程必須都到達(dá)這個(gè)調(diào)用點(diǎn)任何一個(gè)線程沒(méi)到其他線程就得等。第二內(nèi)存柵欄所有線程在__syncthreads()之前對(duì)共享內(nèi)存和全局內(nèi)存的寫(xiě)入在屏障之后對(duì)block內(nèi)所有線程可見(jiàn)。正因?yàn)檫@兩件事綁在一起很多人誤以為_(kāi)_syncthreads就是“線程安全的萬(wàn)能鑰匙”。其實(shí)它很重重在所有線程都必須到齊。如果代碼路徑上有一個(gè)線程繞過(guò)去了整個(gè)block就死鎖。經(jīng)典用法是這樣的__shared__ int tmp[32]; tmp[threadIdx.x] threadIdx.x; __syncthreads(); // 確保所有線程寫(xiě)完tmp int v tmp[(threadIdx.x 1) % 32];去掉__syncthreads()tmp[(threadIdx.x 1) % 32]很可能讀到鄰居線程還沒(méi)寫(xiě)入的舊值。這不是“偶爾出錯(cuò)”在弱內(nèi)存模型下就是未定義行為。3.2 統(tǒng)一到達(dá)原則條件分支和變長(zhǎng)循環(huán)里的死鎖我見(jiàn)過(guò)最多的大坑是有人在條件分支里放_(tái)_syncthreads()if (threadIdx.x 10) { __syncthreads(); // 只有10個(gè)線程會(huì)執(zhí)行 }結(jié)果必然是死鎖。原因很簡(jiǎn)單hardware barrier的計(jì)數(shù)器需要block內(nèi)所有線程都arrive現(xiàn)在只有前10個(gè)線程在傻等剩下22個(gè)線程根本不會(huì)來(lái)湊數(shù)。更隱蔽的是變長(zhǎng)循環(huán)for (int i 0; i threadIdx.x; i) { __syncthreads(); // 每個(gè)線程循環(huán)次數(shù)不同遲早死鎖 }線程0執(zhí)行0次直接跳過(guò)了線程31要執(zhí)行31次兩邊永遠(yuǎn)等不到彼此。還有一種看似安全實(shí)則危險(xiǎn)的寫(xiě)法是在if-else兩個(gè)分支里各放一個(gè)__syncthreads()if (cond) { __syncthreads(); } else { __syncthreads(); }這里所有線程最終都會(huì)執(zhí)行某個(gè)__syncthreads但如果按Volta之后獨(dú)立線程調(diào)度的視角看一部分線程先到達(dá)if分支的barrier另一部分后到達(dá)else分支的barrier——它們等在不同的PC地址上依舊死鎖或產(chǎn)生未定義行為。CUDA要求的是所有線程在源代碼層面到達(dá)同一個(gè)__syncthreads()調(diào)用點(diǎn)。所以社區(qū)有個(gè)不成文的規(guī)矩__syncthreads()永遠(yuǎn)放在所有線程必然執(zhí)行的、無(wú)分支的代碼路徑上。如果你確實(shí)需要分支內(nèi)同步應(yīng)該改用cuda::barrier這類“可分離到達(dá)與等待”的原語(yǔ)而不是往__syncthreads上硬湊。3.3 輕量替代__syncwarp與cooperative_groups很多時(shí)候你并不需要整個(gè)block都同步。比如warp內(nèi)reduce、warp內(nèi)shuffle只需要這一個(gè)warp的線程步調(diào)一致。這時(shí)候用__syncthreads()就太虧了正確的選擇是__syncwarp()。__syncwarp(); // 當(dāng)前warp內(nèi)所有線程到達(dá)后才繼續(xù)默認(rèn)掩碼是全warp還可以按位指定只同步一部分laneunsigned mask __activemask(); // 當(dāng)前活躍的lane集合 __syncwarp(mask);注意__syncwarp在Volta架構(gòu)引入了獨(dú)立線程調(diào)度后行為敏感如果mask與實(shí)際活躍的線程不一致結(jié)果是未定義的。所以最好用__activemask()動(dòng)態(tài)獲取或者直接調(diào)用無(wú)參版本。再進(jìn)一步cooperative_groups庫(kù)把同步表達(dá)得更清晰#include cooperative_groups.h namespace cg cooperative_groups; cg::this_thread_block().sync(); // 等價(jià)于 __syncthreads() auto tiled cg::tiled_partition16(cg::this_thread_block()); tiled.sync(); // 只同步一個(gè)tile內(nèi)的16個(gè)線程cooperative_groups的好處是語(yǔ)義自文檔化讀者一眼看出你同步的范圍是block還是tile。代碼里那種“滿屏__syncthreads靠注釋解釋”的寫(xiě)法用CG之后會(huì)清爽很多。4. 編程指南第6版以來(lái)的現(xiàn)代原語(yǔ)內(nèi)存序、atomic_ref與barrier4.1 從“經(jīng)驗(yàn)性fence”到正式內(nèi)存模型早期寫(xiě)CUDA內(nèi)存同步基本靠一套口口相傳的“經(jīng)驗(yàn)口訣”數(shù)據(jù)寫(xiě)完加fenceflag用atomic輪詢用volatile?？谠E能解決90%的問(wèn)題但剩下10%會(huì)讓人崩潰——因?yàn)闆](méi)人能說(shuō)清fence和atomic到底保證了什么、不保證什么。編程指南第6版引入的正式內(nèi)存模型本質(zhì)上是把C11的內(nèi)存模型搬到了CUDA里給開(kāi)發(fā)者提供了四檔內(nèi)存序memory_order_relaxed只要求原子性不限制順序memory_order_acquire該讀之后的普通讀/寫(xiě)不能重排到它之前memory_order_release該寫(xiě)之前的普通讀/寫(xiě)不能重排到它之后memory_order_seq_cst全序最強(qiáng)的排序約束對(duì)應(yīng)的作用域也有三檔thread_scope_block、thread_scope_device、thread_scope_system正好映射到前文三種fence的范圍。有了正式模型編譯器終于能根據(jù)語(yǔ)義做優(yōu)化而不是靠程序員手動(dòng)插入全局fence“一刀切”保證順序。4.2 cuda::atomic_refrelease/acquire替代裸fencecuda::atomic_ref是libcu提供的原子引用封裝它引用一塊既有內(nèi)存你可以把它當(dāng)原子變量用。先看改寫(xiě)過(guò)后的生產(chǎn)者-消費(fèi)者#include cuda/atomic using cuda::atomic_ref; using cuda::thread_scope_device; using cuda::memory_order_release; using cuda::memory_order_acquire; __global__ void producer_consumer(int* data, int* flag) { atomic_refint, thread_scope_device flag_ref(*flag); if (blockIdx.x 0 threadIdx.x 0) { *data 42; flag_ref.store(1, memory_order_release); } if (blockIdx.x 1 threadIdx.x 0) { while (flag_ref.load(memory_order_acquire) ! 1); printf(data%d\n, *data); } }這段代碼和手寫(xiě)三段式是等價(jià)語(yǔ)義但明顯更精確release store保證*data 42這個(gè)普通store一定在flag寫(xiě)生效前對(duì)消費(fèi)者可見(jiàn)。它不需要額外的fence指令因?yàn)閞elease語(yǔ)義已經(jīng)把這個(gè)約束寫(xiě)進(jìn)編譯器和硬件要遵守的規(guī)則里了。acquire load保證一旦讀到flag1后面讀取*data時(shí)一定能看到release之前所有寫(xiě)的內(nèi)容。作用域被限定在device不會(huì)有多余的系統(tǒng)級(jí)開(kāi)銷。我自己的體會(huì)是用cuda::atomic_ref之后代碼的可讀性和性能都上了一個(gè)臺(tái)階。它把“為什么這里要fence”變成了“這里是一次release/acquire配對(duì)”其他人review代碼時(shí)基本不需要猜。4.3 cuda::barrier把到達(dá)與等待解耦__syncthreads()的問(wèn)題是到達(dá)和等待必須發(fā)生在同一個(gè)調(diào)用點(diǎn)所有線程要么一起到要么一起死。而cuda::barrier提供了分離的arrive和wait讓生產(chǎn)者先標(biāo)記“我到了”然后去干別的活消費(fèi)者等所有生產(chǎn)者都arrive后再繼續(xù)?；灸Ｊ绞沁@樣的#include cuda/barrier using barrier_t cuda::barriercuda::thread_scope_block; __global__ void barrier_demo(int* out, int n) { __shared__ barrier_t bar; if (threadIdx.x 0) { // 初始化參與的線程數(shù)不同CUDA版本初始化API略有差異 // placement new 或 init() 均可詳見(jiàn)官方libcu文檔 init(bar, blockDim.x); } __syncthreads(); // 每個(gè)線程做自己的階段一 int v out[threadIdx.x] * 2; auto token bar.arrive(); // “我階段一干完了” bar.wait(cuda::std::move(token)); // 等其他人都干完 // 階段二此時(shí)所有線程都可以安全讀取階段一的數(shù)據(jù) out[threadIdx.x] v out[(threadIdx.x 1) % blockDim.x]; }arrive返回一個(gè)tokenwait消費(fèi)這個(gè)token。這一步把“完成信號(hào)”和“等待條件”解耦了正好能套進(jìn)流水線算法一批線程arrive后馬上開(kāi)始計(jì)算下一塊數(shù)據(jù)而不是傻傻等著別人全部就位才開(kāi)始。cuda::barrier與__syncthreads的另一個(gè)區(qū)別是barrier可以跨block配置作用域比如cuda::thread_scope_device的barrier能讓不同block的線程互相等待——這正是__syncthreads做不到的。當(dāng)然跨block barrier需要確保所有block同時(shí)駐留這通常配合cooperative launch使用。4.4 異步拷貝與barrier的流水線配合更進(jìn)一步cuda::barrier還經(jīng)常和cuda::memcpy_async配合做異步共享內(nèi)存拷貝的完成同步。這個(gè)套路在科學(xué)計(jì)算里幾乎是標(biāo)配cuda::memcpy_async(shared_buf[0], global_data[offset], shared_size, cuda::pipeline::memcpy_async_thread_scope_block); auto token bar.arrive(); bar.wait(cuda::std::move(token)); // 此時(shí)shared_buf才安全可讀memcpy_async發(fā)起的是異步拷貝數(shù)據(jù)真正到位的時(shí)間點(diǎn)是不確定的。用barrier的arrive/wait去承接“拷貝完成”事件比靠固定延遲的__syncthreads空等要精準(zhǔn)得多也能讓計(jì)算和顯存搬運(yùn)重疊起來(lái)。如果在Ampere及更新的架構(gòu)上底層還有硬件級(jí)mbarrier可以直接操作吞吐更高但API也更深。建議一般項(xiàng)目從cuda::barrier入手理解清楚再往下鉆。5. 真實(shí)項(xiàng)目里的避坑記錄柵欄與同步的取舍5.1 最常被混淆的一對(duì)fence與sync過(guò)去一年我評(píng)審過(guò)的CUDA代碼里出現(xiàn)率最高的錯(cuò)誤是把__threadfence()和__syncthreads()當(dāng)成同一個(gè)東西。有人寫(xiě)floating-point累加想讓一個(gè)block先寫(xiě)完另一個(gè)block來(lái)讀于是在producer加了__syncthreads()期望“同步之后別人就能看到了”。結(jié)果__syncthreads只同步本block內(nèi)線程對(duì)block 2毫無(wú)約束力對(duì)方照樣讀到舊值。反過(guò)來(lái)有人處理共享內(nèi)存復(fù)用在消費(fèi)者側(cè)死等__threadfence()從來(lái)不調(diào)用__syncthreads()結(jié)果共享內(nèi)存數(shù)據(jù)還沒(méi)寫(xiě)完就開(kāi)始讀。fence不等待其他線程它只是單方面承諾“我的寫(xiě)已經(jīng)提交了”但沒(méi)人保證對(duì)方已經(jīng)執(zhí)行到該讀的位置。一句話總結(jié)需要“所有線程到達(dá)同一個(gè)位置” → 用同步__syncthreads、barrier需要“我的寫(xiě)入對(duì)外可見(jiàn)” → 用柵欄fence、release/acquire兩樣都要 → 用barrier或組合原語(yǔ)判斷表記牢能省掉一半調(diào)試時(shí)間。5.2 作用域?yàn)E用全量fence拖慢熱循環(huán)__threadfence_system()看著很穩(wěn)但它是三個(gè)fence里最貴的一個(gè)。它的語(yǔ)義覆蓋到主機(jī)端系統(tǒng)內(nèi)存往往需要刷新比L2更遠(yuǎn)的路徑。設(shè)備端通信本來(lái)不需要觸碰主機(jī)內(nèi)存用system fence就是純浪費(fèi)。我實(shí)際項(xiàng)目里測(cè)過(guò)一組數(shù)據(jù)一個(gè)塊間通信的熱循環(huán)每秒做約10萬(wàn)次flag交換。三個(gè)版本耗時(shí)對(duì)比大致如下實(shí)現(xiàn)相對(duì)耗時(shí)__threadfence_system() atomicExch1.4x__threadfence() atomicExch1.0xcuda::atomic_refrelease/acquire0.8x不同架構(gòu)比例會(huì)有浮動(dòng)但趨勢(shì)穩(wěn)定作用域越寬越慢語(yǔ)義越精確越快。所以寫(xiě)代碼前先問(wèn)自己一句通信雙方到底在什么范圍只在block內(nèi)就用thread_scope_block只在設(shè)備內(nèi)就用thread_scope_device別一上來(lái)就system。5.3 編譯器的二次重排volatile與原子讀的必要性還有一個(gè)坑和編譯器有關(guān)。即使你寫(xiě)了fence如果輪詢進(jìn)程里用的是普通int*指針編譯器完全可能在-O3下把整個(gè)循環(huán)優(yōu)化成int tmp *flag; while (tmp ! 1) {}然后你的fence再正確也沒(méi)用——線程壓根沒(méi)在讀內(nèi)存。我遇到過(guò)一次在啟用了--use_fast_math和激進(jìn)優(yōu)化后flag輪詢直接“熔斷”程序掛死。解決方案有兩條路把flag聲明成volatile int*強(qiáng)制每次讀都走內(nèi)存用原子讀比如atomicAdd(flag, 0)或cuda::atomic_ref的load。我個(gè)人推薦后者因?yàn)関olatile只保證“讀內(nèi)存”不保證“原子性”和“內(nèi)存序語(yǔ)義”。用原子讀配合acquire語(yǔ)義完整得多。當(dāng)你需要編譯器別亂動(dòng)又需要精確排序時(shí)原子操作內(nèi)存序是正解volatile只是應(yīng)急手段。5.4 可見(jiàn)性問(wèn)題的定位手段與工具鏈最后給一套我平時(shí)排查內(nèi)存可見(jiàn)性問(wèn)題的三板斧。第一板斧最小復(fù)現(xiàn)。把通信邏輯拆出來(lái)做成一個(gè)只有2個(gè)block、每block只有1個(gè)活躍線程的kernel。如果這個(gè)最簡(jiǎn)模型還出錯(cuò)那問(wèn)題100%出在同步原語(yǔ)本身如果最簡(jiǎn)模型好了說(shuō)明是周圍代碼的重排邏輯在搗亂。第二板斧工具掃描。用compute-sanitizer的racecheck工具compute-sanitizer --tool racecheck ./my_app它對(duì)共享內(nèi)存的race檢測(cè)很成熟對(duì)全局內(nèi)存的可見(jiàn)性問(wèn)題會(huì)有一定誤報(bào)但能幫你快速縮小范圍。注意racecheck報(bào)出的每一處都要人工確認(rèn)不能盲目照單全收。第三板斧看SASS。到Nsight Compute里看一眼生成的指令序列找membar.gl、fence.acq_rel.gpu這類指令。如果以及寫(xiě)了__threadfence()卻沒(méi)看到任何fence指令說(shuō)明編譯器認(rèn)為這個(gè)fence是多余的——這本身就是一個(gè)信號(hào)告訴你內(nèi)存訪問(wèn)在編譯層已經(jīng)被重排了可能需要改用原子操作讓編譯器保留語(yǔ)義。三板斧走下來(lái)絕大多數(shù)可見(jiàn)性問(wèn)題都能定位。剩下的那部分通常不是fence放少了而是作用域選錯(cuò)了——回頭看看5.2的判斷表。我在實(shí)際項(xiàng)目里最深的體會(huì)是柵欄和同步函數(shù)不是“性能優(yōu)化技巧”而是CUDA正確性的基礎(chǔ)設(shè)施。早期寫(xiě)kernel時(shí)覺(jué)得它們礙事能省則省后來(lái)被時(shí)好時(shí)壞的bug折磨過(guò)幾輪才明白該用的地方一個(gè)都不能省。如果你剛接觸CUDA建議把這三類API當(dāng)核心語(yǔ)法對(duì)待而不是等出問(wèn)題了再回頭補(bǔ)課。