1. CPU vs. GPU

1.1 四種計算機(jī)模型

GPU設(shè)計的初衷就是為了減輕CPU計算的負(fù)載，將一部分圖形計算的功能設(shè)計到一塊獨(dú)立的處理器中锭弊，將矩陣變換堪澎、頂點(diǎn)計算和光照計算等操作從 CPU 中轉(zhuǎn)移到 GPU中，從而一方面加速圖形處理味滞，另一方面減小了 CPU 的工作負(fù)載樱蛤，讓 CPU 有時間去處理其它的事情。

在GPU上的各個處理器采取異步并行的方式對數(shù)據(jù)流進(jìn)行處理剑鞍，根據(jù)費(fèi)林分類法（Flynn's Taxonomy）昨凡，可以將資訊流（information stream）分成指令（Instruction）和數(shù)據(jù)（Data）兩種，據(jù)此又可分成四種計算機(jī)類型：

• 單一指令流單一數(shù)據(jù)流計算機(jī)（SISD）：單核CPU
• 單一指令流多數(shù)據(jù)流計算機(jī)（SIMD）：GPU的計算模型
• 多指令流單一數(shù)據(jù)流計算機(jī)（MISD）：流水線模型
• 多指令流多數(shù)據(jù)流計算機(jī)（MIMD）：多核CPU

Flynn's Taxonomy

1.2 CPU 與 GPU 結(jié)構(gòu)差異

cpu vs. gpu

cpu vs. gpu

1.3 CPU設(shè)計理念：低延時

Latency oriented design

■ ALU：CPU有強(qiáng)大的ALU（算術(shù)運(yùn)算單元）,它可以在很少的時鐘周期內(nèi)完成算術(shù)計算蚁署。

• 當(dāng)今的CPU可以達(dá)到64bit 雙精度便脊。執(zhí)行雙精度浮點(diǎn)源算的加法和乘法只需要1～3個時鐘周期。

• CPU的時鐘周期的頻率是非常高的光戈，達(dá)到1.532～3gigahertz(千兆HZ, 10的9次方).

■ Cache：大的緩存也可以降低延時哪痰。保存很多的數(shù)據(jù)放在緩存里面遂赠，當(dāng)需要訪問的這些數(shù)據(jù)，只要在之前訪問過的晌杰，如今直接在緩存里面取即可跷睦。

■ Control：復(fù)雜的邏輯控制單元。

• 當(dāng)程序含有多個分支的時候肋演，它通過提供分支預(yù)測的能力來降低延時抑诸。

• 數(shù)據(jù)轉(zhuǎn)發(fā)。 當(dāng)一些指令依賴前面的指令結(jié)果時爹殊，數(shù)據(jù)轉(zhuǎn)發(fā)的邏輯控制單元決定這些指令在pipeline中的位置并且盡可能快的轉(zhuǎn)發(fā)一個指令的結(jié)果給后續(xù)的指令蜕乡。這些動作需要很多的對比電路單元和轉(zhuǎn)發(fā)電路單元。

1.4 GPU設(shè)計理念：大吞吐量

Throughput oriented design

■ ALU梗夸，Cache：GPU的特點(diǎn)是有很多的ALU和很少的cache. 緩存的目的不是保存后面需要訪問的數(shù)據(jù)的层玲，這點(diǎn)和CPU不同，而是為thread提高服務(wù)的绒瘦。如果有很多線程需要訪問同一個相同的數(shù)據(jù)称簿，緩存會合并這些訪問扣癣，然后再去訪問dram（因?yàn)樾枰L問的數(shù)據(jù)保存在dram中而不是cache里面）惰帽，獲取數(shù)據(jù)后cache會轉(zhuǎn)發(fā)這個數(shù)據(jù)給對應(yīng)的線程，這個時候是數(shù)據(jù)轉(zhuǎn)發(fā)的角色父虑。但是由于需要訪問dram该酗，自然會帶來延時的問題。

■ Control：控制單元（左邊黃色區(qū)域塊）可以把多個的訪問合并成少的訪問士嚎。

GPU的雖然有dram延時呜魄，卻有非常多的ALU和非常多的thread. 為了平衡內(nèi)存延時的問題，我們可以中充分利用多的ALU的特性達(dá)到一個非常大的吞吐量的效果莱衩。盡可能多的分配多的Threads.通常來看GPU ALU會有非常重的pipeline就是因?yàn)檫@樣爵嗅。

2. Nvidia GPU架構(gòu)

2.1 硬件架構(gòu)

■ SP：最基本的處理單元，streaming processor笨蚁，也稱為CUDA core睹晒。最后具體的指令和任務(wù)都是在SP上處理的。GPU進(jìn)行并行計算括细，也就是很多個SP同時做處理伪很。

■ SM：多個SP加上其他的一些資源組成一個streaming multiprocessor。也叫GPU大核奋单，其他資源如：warp scheduler锉试，register，shared memory等览濒。SM可以看做GPU的心臟（對比CPU核心）呆盖，register和shared memory是SM的稀缺資源拖云。CUDA將這些資源分配給所有駐留在SM中的threads。因此应又，這些有限的資源就使每個SM中active warps有非常嚴(yán)格的限制江兢，也就限制了并行能力。

gpu hardware structure

具有Tesla架構(gòu)的GPU是具有芯片共享存儲器的一組SIMT（單指令多線程）多處理器丁频。它以一個可伸縮的多線程流處理器（Streaming Multiprocessors杉允，SMs）陣列為中心實(shí)現(xiàn)了MIMD（多指令多數(shù)據(jù)）的異步并行機(jī)制，其中每個多處理器包含多個標(biāo)量處理器（Scalar Processor席里，SP）叔磷，為了管理運(yùn)行各種不同程序的數(shù)百個線程，SIMT架構(gòu)的多處理器會將各線程映射到一個標(biāo)量處理器核心奖磁，各標(biāo)量線程使用自己的指令地址和寄存器狀態(tài)獨(dú)立執(zhí)行改基。

GPU的共享存儲器的SIMT多處理器模型

如上圖所示，每個多處理器(Multiprocessor)都有一個屬于以下四種類型之一的芯片存儲器：

• 每個處理器上有一組本地 32 位寄存器（Registers）咖为；
• 并行數(shù)據(jù)緩存或共享存儲器(Shared Memory)秕狰，由所有標(biāo)量處理器核心共享，共享存儲器空間就位于此處躁染；
• 只讀固定緩存(Constant Cache)鸣哀，由所有標(biāo)量處理器核心共享，可加速從固定存儲器空間進(jìn)行的讀取操作（這是設(shè)備存儲器的一個只讀區(qū)域）吞彤；
• 一個只讀紋理緩存(Texture Cache)我衬，由所有標(biāo)量處理器核心共享，加速從紋理存儲器空間進(jìn)行的讀取操作（這是設(shè)備存儲器的一個只讀區(qū)域）饰恕，每個多處理器都會通過實(shí)現(xiàn)不同尋址模型和數(shù)據(jù)過濾的紋理單元訪問紋理緩存挠羔。

多處理器 SIMT 單元以32個并行線程為一組來創(chuàng)建、管理埋嵌、調(diào)度和執(zhí)行線程破加，這樣的線程組稱為 warp 塊(束)，即以線程束為調(diào)度單位雹嗦，但只有所有32個線程都在諸如內(nèi)存讀取這樣的操作時范舀，它們就會被掛起，如下所示的狀態(tài)變化俐银。當(dāng)主機(jī)CPU上的CUDA程序調(diào)用內(nèi)核網(wǎng)格時尿背，網(wǎng)格的塊將被枚舉并分發(fā)到具有可用執(zhí)行容量的多處理器；SIMT 單元會選擇一個已準(zhǔn)備好執(zhí)行的 warp 塊捶惜，并將下一條指令發(fā)送到該 warp 塊的活動線程田藐。一個線程塊的線程在一個多處理器上并發(fā)執(zhí)行，在線程塊終止時，將在空閑多處理器上啟動新塊汽久。

CPU五種狀態(tài)的轉(zhuǎn)換

線程束調(diào)度變化

2.2 軟件架構(gòu)

CUDA是一種新的操作GPU計算的硬件和軟件架構(gòu)鹤竭，它將GPU視作一個數(shù)據(jù)并行計算設(shè)備，而且無需把這些計算映射到圖形API景醇。操作系統(tǒng)的多任務(wù)機(jī)制可以同時管理CUDA訪問GPU和圖形程序的運(yùn)行庫臀稚，其計算特性支持利用CUDA直觀地編寫GPU核心程序。目前Tesla架構(gòu)具有在筆記本電腦三痰、臺式機(jī)吧寺、工作站和服務(wù)器上的廣泛可用性，配以C/C++語言的編程環(huán)境和CUDA軟件散劫，使這種架構(gòu)得以成為最優(yōu)秀的超級計算平臺稚机。

CUDA軟件層次結(jié)構(gòu)

CUDA在軟件方面組成有：一個CUDA庫、一個應(yīng)用程序編程接口（API）及其運(yùn)行庫(Runtime)获搏、兩個較高級別的通用數(shù)學(xué)庫赖条，即CUFFT和CUBLAS。CUDA改進(jìn)了DRAM的讀寫靈活性常熙，使得GPU與CPU的機(jī)制相吻合纬乍。另一方面，CUDA 提供了片上（on-chip）共享內(nèi)存裸卫，使得線程之間可以共享數(shù)據(jù)仿贬。應(yīng)用程序可以利用共享內(nèi)存來減少DRAM的數(shù)據(jù)傳送，更少的依賴DRAM的內(nèi)存帶寬彼城。

■ thread：一個CUDA的并行程序會被以許多個threads來執(zhí)行诅蝶。

■ block：數(shù)個threads會被群組成一個block退个，同一個block中的threads可以同步募壕，也可以通過shared memory通信。

■ grid：多個blocks則會再構(gòu)成grid语盈。

■ warp：GPU執(zhí)行程序時的調(diào)度單位舱馅，目前cuda的warp的大小為32，同在一個warp的線程刀荒，以不同數(shù)據(jù)資源執(zhí)行相同的指令,這就是所謂 SIMT代嗤。

gpu software struture

2.3 軟硬件架構(gòu)對應(yīng)關(guān)系

1）SM像一個獨(dú)立的CPU core

從軟件上看，SM更像一個獨(dú)立的CPU core缠借。SM（Streaming Multiprocessors）是GPU架構(gòu)中非常重要的部分干毅，GPU硬件的并行性就是由SM決定的。以Fermi架構(gòu)為例泼返，其包含以下主要組成部分：

• CUDA cores
• Shared Memory/L1Cache
• Register File
• Load/Store Units
• Special Function Units
• Warp Scheduler

2）同一個block的threads在同一個SM并行執(zhí)行

GPU中每個sm都設(shè)計成支持?jǐn)?shù)以百計的線程并行執(zhí)行硝逢，并且每個GPU都包含了很多的SM，所以GPU支持成百上千的線程并行執(zhí)行。當(dāng)一個kernel啟動后渠鸽，thread會被分配到這些SM中執(zhí)行叫乌。大量的thread可能會被分配到不同的SM，同一個block中的threads必然在同一個SM中并行（SIMT）執(zhí)行徽缚。每個thread擁有它自己的程序計數(shù)器和狀態(tài)寄存器憨奸，并且用該線程自己的數(shù)據(jù)執(zhí)行指令，這就是所謂的Single Instruction Multiple Thread凿试。

3）warp是調(diào)度和運(yùn)行的基本單元排宰，一個warp占用一個SM運(yùn)行

一個SP可以執(zhí)行一個thread，但是實(shí)際上并不是所有的thread能夠在同一時刻執(zhí)行那婉。Nvidia把32個threads組成一個warp额各，warp是調(diào)度和運(yùn)行的基本單元。warp中所有threads并行的執(zhí)行相同的指令吧恃。一個warp需要占用一個SM運(yùn)行虾啦，多個warps需要輪流進(jìn)入SM。由SM的硬件warp scheduler負(fù)責(zé)調(diào)度痕寓。目前每個warp包含32個threads（Nvidia保留修改數(shù)量的權(quán)利）傲醉。所以，一個GPU上resident thread最多只有 SM*warp個呻率。

block是軟件概念硬毕，一個block只會由一個sm調(diào)度，程序員在開發(fā)時礼仗，通過設(shè)定block的屬性吐咳，告訴GPU硬件，我有多少個線程元践，線程怎么組織韭脊。而具體怎么調(diào)度由sm的warps scheduler負(fù)責(zé)，block一旦被分配好SM单旁，該block就會一直駐留在該SM中沪羔，直到執(zhí)行結(jié)束。一個SM可以同時擁有多個blocks象浑，但需要序列執(zhí)行蔫饰。下圖顯示了軟件硬件方面的術(shù)語對應(yīng)關(guān)系：

software v.s. hardware

需要注意的是，大部分threads只是邏輯上并行愉豺，并不是所有的thread可以在物理上同時執(zhí)行篓吁。例如，遇到分支語句（if else蚪拦，while杖剪，for等）時节腐，各個thread的執(zhí)行條件不一樣必然產(chǎn)生分支執(zhí)行，這就導(dǎo)致同一個block中的線程可能會有不同步調(diào)摘盆。另外翼雀，并行thread之間的共享數(shù)據(jù)會導(dǎo)致競態(tài)：多個線程請求同一個數(shù)據(jù)會導(dǎo)致未定義行為。CUDA提供了cudaThreadSynchronize()來同步同一個block的thread以保證在進(jìn)行下一步處理之前孩擂，所有thread都到達(dá)某個時間點(diǎn)狼渊。

同一個warp中的thread可以以任意順序執(zhí)行，active warps被sm資源限制类垦。當(dāng)一個warp空閑時狈邑，SM就可以調(diào)度駐留在該SM中另一個可用warp。在并發(fā)的warp之間切換是沒什么消耗的蚤认，因?yàn)橛布Y源早就被分配到所有thread和block米苹，所以該新調(diào)度的warp的狀態(tài)已經(jīng)存儲在SM中了。不同于CPU砰琢，CPU切換線程需要保存/讀取線程上下文（register內(nèi)容）蘸嘶，這是非常耗時的棒搜，而GPU為每個threads提供物理register指煎，無需保存/讀取上下文。

4）SIMT和SIMD

CUDA是一種典型的SIMT架構(gòu)（單指令多線程架構(gòu)）凿滤，SIMT和SIMD（Single Instruction, Multiple Data）類似挚冤，SIMT應(yīng)該算是SIMD的升級版况增，更靈活，但效率略低训挡，SIMT是NVIDIA提出的GPU新概念澳骤。二者都通過將同樣的指令廣播給多個執(zhí)行官單元來實(shí)現(xiàn)并行。一個主要的不同就是澜薄，SIMD要求所有的vector element在一個統(tǒng)一的同步組里同步的執(zhí)行为肮，而SIMT允許線程們在一個warp中獨(dú)立的執(zhí)行。SIMT有三個SIMD沒有的主要特征：

• 每個thread擁有自己的instruction address counter
• 每個thread擁有自己的狀態(tài)寄存器
• 每個thread可以有自己獨(dú)立的執(zhí)行路徑

更細(xì)節(jié)的差異可以看這里表悬。

3. CUDA C編程入門

Nvidia官方教程看這里

3.1 編程模型

CUDA程序構(gòu)架分為兩部分：Host和Device弥锄。一般而言，Host指的是CPU蟆沫，Device指的是GPU。在CUDA程序構(gòu)架中温治，主程序還是由 CPU 來執(zhí)行饭庞，而當(dāng)遇到數(shù)據(jù)并行處理的部分，CUDA 就會將程序編譯成 GPU 能執(zhí)行的程序熬荆，并傳送到GPU舟山。而這個程序在CUDA里稱做核（kernel）。CUDA允許程序員定義稱為核的C語言函數(shù)，從而擴(kuò)展了 C 語言累盗，在調(diào)用此類函數(shù)時寒矿，它將由N個不同的CUDA線程并行執(zhí)行N次，這與普通的C語言函數(shù)只執(zhí)行一次的方式不同若债。執(zhí)行核的每個線程都會被分配一個獨(dú)特的線程ID符相，可通過內(nèi)置的threadIdx變量在內(nèi)核中訪問此ID。

在 CUDA 程序中蠢琳，主程序在調(diào)用任何 GPU 內(nèi)核之前啊终，必須對核進(jìn)行執(zhí)行配置，即確定線程塊數(shù)和每個線程塊中的線程數(shù)以及共享內(nèi)存大小傲须。

1）線程層次結(jié)構(gòu)

在GPU中要執(zhí)行的線程蓝牲，根據(jù)最有效的數(shù)據(jù)共享來創(chuàng)建塊(Block)，其類型有一維泰讽、二維或三維例衍。在同一個塊內(nèi)的線程可彼此協(xié)作，通過一些共享存儲器來共享數(shù)據(jù)已卸，并同步其執(zhí)行來協(xié)調(diào)存儲器訪問肄渗。一個塊中的所有線程都必須位于同一個處理器核心中。因而咬最，一個處理器核心的有限存儲器資源制約了每個塊的線程數(shù)量翎嫡。在早起的 NVIDIA 架構(gòu)中，一個線程塊最多可以包含 512 個線程永乌，而在后期出現(xiàn)的一些設(shè)備中則最多可支持1024個線程惑申。一般 GPGPU 程序線程數(shù)目是很多的，所以不能把所有的線程都塞到同一個塊里翅雏。但一個內(nèi)核可由多個大小相同的線程塊同時執(zhí)行圈驼，因而線程總數(shù)應(yīng)等于每個塊的線程數(shù)乘以塊的數(shù)量。這些同樣維度和大小的塊將組織為一個一維或二維線程塊網(wǎng)格(Grid)望几。具體框架如下圖所示绩脆。

線程塊網(wǎng)格

核函數(shù)只能在主機(jī)端調(diào)用，其調(diào)用形式為：Kernel<<<Dg,Db, Ns, S>>>(param list)

• Dg：用于定義整個grid的維度和尺寸橄抹，即一個grid有多少個block靴迫。為dim3類型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x個block楼誓，每列有Dg.y個block玉锌，第三維恒為1(目前一個核函數(shù)只有一個grid)。整個grid中共有Dg.x*Dg.y個block疟羹，其中Dg.x和Dg.y最大值為65535主守。
• Db：用于定義一個block的維度和尺寸禀倔，即一個block有多少個thread。為dim3類型参淫。Dim3 Db(Db.x, Db.y, Db.z)表示整個block中每行有Db.x個thread救湖，每列有Db.y個thread，高度為Db.z涎才。Db.x和Db.y最大值為512鞋既，Db.z最大值為62。一個block中共有Db.xDb.yDb.z個thread憔维。計算能力為1.0,1.1的硬件該乘積的最大值為768涛救，計算能力為1.2,1.3的硬件支持的最大值為1024。
• Ns：是一個可選參數(shù)业扒，用于設(shè)置每個block除了靜態(tài)分配的shared Memory以外检吆，最多能動態(tài)分配的shared memory大小，單位為byte程储。不需要動態(tài)分配時該值為0或省略不寫蹭沛。
• S：是一個cudaStream_t類型的可選參數(shù)，初始值為零章鲤，表示該核函數(shù)處在哪個流之中摊灭。

如下是一個CUDA簡單的求和程序：

CUDA求和程序

2）存儲器層次結(jié)構(gòu)

CUDA 設(shè)備擁有多個獨(dú)立的存儲空間，其中包括：全局存儲器败徊、本地存儲器帚呼、共享存儲器、常量存儲器皱蹦、紋理存儲器和寄存器煤杀，如下圖所示。

CUDA設(shè)備上的存儲器

CUDA線程可在執(zhí)行過程中訪問多個存儲器空間的數(shù)據(jù)沪哺，如下圖所示其中：

• 每個線程都有一個私有的本地存儲器沈自。
• 每個線程塊都有一個共享存儲器，該存儲器對于塊內(nèi)的所有線程都是可見的辜妓，并且與塊具有相同的生命周期枯途。
• 所有線程都可訪問相同的全局存儲器。
• 此外還有兩個只讀的存儲器空間籍滴，可由所有線程訪問酪夷，這兩個空間是常量存儲器空間和紋理存儲器空間。全局异逐、固定和紋理存儲器空間經(jīng)過優(yōu)化捶索，適于不同的存儲器用途。紋理存儲器也為某些特殊的數(shù)據(jù)格式提供了不同的尋址模式以及數(shù)據(jù)過濾灰瞻，方便 Host對流數(shù)據(jù)的快速存取腥例。

存儲器的應(yīng)用層次

3）主機(jī)（Host）和設(shè)備（Device）

如下圖所示，CUDA 假設(shè)線程可在物理上獨(dú)立的設(shè)備上執(zhí)行酝润，此類設(shè)備作為運(yùn)行C語言程序的主機(jī)的協(xié)處理器操作燎竖。內(nèi)核在GPU上執(zhí)行，而C語言程序的其他部分在CPU上執(zhí)行（即串行代碼在主機(jī)上執(zhí)行要销，而并行代碼在設(shè)備上執(zhí)行）构回。此外，CUDA還假設(shè)主機(jī)和設(shè)備均維護(hù)自己的DRAM疏咐，分別稱為主機(jī)存儲器和設(shè)備存儲器纤掸。因而，一個程序通過調(diào)用CUDA運(yùn)行庫來管理對內(nèi)核可見的全局浑塞、固定和紋理存儲器空間借跪。這種管理包括設(shè)備存儲器的分配和取消分配，還包括主機(jī)和設(shè)備存儲器之間的數(shù)據(jù)傳輸酌壕。

CUDA異構(gòu)編程模型

3.2 編程入門

3.2.1 CUDA C基礎(chǔ)

CUDA C是對C/C++語言進(jìn)行拓展后形成的變種掏愁，兼容C/C++語法，文件類型為".cu"文件卵牍，編譯器為"nvcc"果港，相比傳統(tǒng)的C/C++，主要添加了以下幾個方面：

• 函數(shù)類型限定符
• 執(zhí)行配置運(yùn)算符
• 五個內(nèi)置變量
• 變量類型限定符
• 其他的還有數(shù)學(xué)函數(shù)糊昙、原子函數(shù)辛掠、紋理讀取、綁定函數(shù)等

1）函數(shù)類型限定符
用來確定某個函數(shù)是在CPU還是GPU上運(yùn)行释牺，以及這個函數(shù)是從CPU調(diào)用還是從GPU調(diào)用

• device表示從GPU調(diào)用萝衩，在GPU上執(zhí)行
• global表示從CPU調(diào)用，在GPU上執(zhí)行船侧，也稱之為kernel函數(shù)
• host表示在CPU上調(diào)用欠气，在CPU上執(zhí)行

在計算能力3.0及以后的設(shè)備中，global類型的函數(shù)也可以調(diào)用__global類型的函數(shù)镜撩。

#include <stdio.h>

__device__ void device_func(void) {
}

__global__ void global_func(void) {
    device_func();
}

int main() {
    printf("%s\n", __FILE__);
    global_func<<<1,1>>>();
    return 0;
}

2）執(zhí)行配置運(yùn)算符
執(zhí)行配置運(yùn)算符<<<>>>预柒，用來傳遞內(nèi)核函數(shù)的執(zhí)行參數(shù)。格式如下：
kernel<<<gridDim, blockDim, memSize, stream>>>(para1, para2, ...);

• gridDim表示網(wǎng)格的大小袁梗，可以是1,2,3維
• blockDim表示塊的·大小宜鸯，可以是1,2,3維
• memSize表示動態(tài)分配的共享存儲器大小，默認(rèn)為0
• stream表示執(zhí)行的流遮怜，默認(rèn)為0
• para1, para2等為核函數(shù)參數(shù)

#include <stdio.h>

__global__ void func(int a, int b) {
}

int main() {
    int a = 0, b = 0;
    func<<<128, 128>>>(a, b);
    func<<<dim3(128, 128), dim3(16, 16)>>>(a, b);
    func<<<dim3(128, 128, 128), dim3(16, 16, 2)>>>(a, b);
    return 0;
}

3）五個內(nèi)置變量
這些內(nèi)置變量用來在運(yùn)行時獲得Grid和Block的尺寸及線程索引等信息

• gridDim: 包含三個元素x, y, z的結(jié)構(gòu)體淋袖，表示Grid在三個方向上的尺寸，對應(yīng)于執(zhí)行配置中的第一個參數(shù)
• blockDim: 包含上元素x, y, z的結(jié)構(gòu)體锯梁，表示Block在三個方向上的尺寸即碗，對應(yīng)于執(zhí)行配置中的第二個參數(shù)
• blockIdx: 包含三個元素x, y, z的結(jié)構(gòu)體焰情，分別表示當(dāng)前線程所在塊在網(wǎng)格中x, y, z方向上的索引
• threadIdx: 包含三個元素x, y, z的結(jié)構(gòu)體，分別表示當(dāng)前線程在其所在塊中x, y, z方向上的索引
• warpSize: 表明warp的尺寸

線程tid的計算

線程tid的計算

4）變量類型限定符
用來確定某個變量在設(shè)備上的內(nèi)存位置

• _device_表示位于全局內(nèi)存空間剥懒，默認(rèn)類型
• _share_表示位于共享內(nèi)存空間
• _constant_表示位于常量內(nèi)存空間
• texture表示其綁定的變量可以被紋理緩存加速訪問

3.2.2 示例

向量的點(diǎn)積——假設(shè)向量大小為N内舟，按照上下文的方法，將申請N的空間初橘，用來存放向量元素互乘的結(jié)果验游，然后在CPU上對N個乘積進(jìn)行累加。
類似于Map-Reduce：

• Map——GPU上對N個變量分別相乘
• Reduce——CPU上對N個乘積進(jìn)行累加

點(diǎn)積

點(diǎn)積