国产av日韩一区二区三区精品,成人性爱视频在线观看,国产,欧美,日韩,一区,www.成色av久久成人,2222eeee成人天堂

目錄
一、GEMM的基本特徵
1.1 GEMM計(jì)算過程及複雜度
2.2 解決Bank Conflict 問題
2.3 流水並行化:Double Buffering
三、cuBLAS 實(shí)作方式探究
首頁 科技週邊 人工智慧 CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

Mar 25, 2024 pm 12:30 PM
系統(tǒng) 計(jì)算 overflow 標(biāo)準(zhǔn)函式庫

通用矩陣乘法(General Matrix Multiplication,GEMM)是許多應(yīng)用程式和演算法中至關(guān)重要的一部分,也是評(píng)估電腦硬體效能的重要指標(biāo)之一。透過深入研究和優(yōu)化GEMM的實(shí)現(xiàn),可以幫助我們更好地理解高效能運(yùn)算以及軟硬體系統(tǒng)之間的關(guān)係。在電腦科學(xué)中,對(duì)GEMM進(jìn)行有效的最佳化可以提高運(yùn)算速度並節(jié)省資源,這對(duì)於提高電腦系統(tǒng)的整體效能至關(guān)重要。深入了解GEMM的工作原理和最佳化方法,有助於我們更好地利用現(xiàn)代計(jì)算硬體的潛力,並為各種複雜計(jì)算任務(wù)提供更有效率的解決方案。透過GEMM效能的最佳化與改進(jìn),可以加上

一、GEMM的基本特徵

1.1 GEMM計(jì)算過程及複雜度

#GEMM 的定義為:

CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

矩陣乘法的計(jì)算示意圖

1.2 簡(jiǎn)單實(shí)作及過程分析

下面是依照原始定義實(shí)現(xiàn)的CPU 上實(shí)現(xiàn)的程式碼,之後用以作為精度的對(duì)照
#define OFFSET(row, col, ld) ((row) * (ld) + (col))void cpuSgemm(float *a, float *b, float *c, const int M, const int N, const int K) {for (int m = 0; m <p></p>#下面使用CUDA實(shí)現(xiàn)最簡(jiǎn)單的矩陣乘法的Kernal,一共使用M * N 個(gè)執(zhí)行緒完成整個(gè)矩陣乘法。每個(gè)執(zhí)行緒負(fù)責(zé)矩陣C中一個(gè)元素的計(jì)算,需要完成K次乘累加。矩陣A,B,C皆存放與全域記憶體中(由修飾符<div>?</div>###__global__######?###確定),完整程式碼請(qǐng)見 sgemm_naive.cu 。 ######<pre class="brush:php;toolbar:false">__global__ void naiveSgemm(float * __restrict__ a, float * __restrict__ b, float * __restrict__ c,const int M, const int N, const int K) {int n = blockIdx.x * blockDim.x + threadIdx.x;int m = blockIdx.y * blockDim.y + threadIdx.y;if (m ######編譯完成,在Tesla V100-PCIE-32GB上執(zhí)行的結(jié)果如下,根據(jù)V100的白皮書,F(xiàn)P32 的峰值算力為15.7 TFLOPS,因此該方式算力利用率僅有11.5%。 ######<pre class="brush:php;toolbar:false">M N K =128128 1024, Time = 0.00010083 0.00010260 0.00010874 s, AVG Performance = 304.5951 GflopsM N K =192192 1024, Time = 0.00010173 0.00010198 0.00010253 s, AVG Performance = 689.4680 GflopsM N K =256256 1024, Time = 0.00010266 0.00010318 0.00010384 s, AVG Performance =1211.4281 GflopsM N K =384384 1024, Time = 0.00019475 0.00019535 0.00019594 s, AVG Performance =1439.7206 GflopsM N K =512512 1024, Time = 0.00037693 0.00037794 0.00037850 s, AVG Performance =1322.9753 GflopsM N K =768768 1024, Time = 0.00075238 0.00075558 0.00075776 s, AVG Performance =1488.9271 GflopsM N K = 1024 1024 1024, Time = 0.00121562 0.00121669 0.00121789 s, AVG Performance =1643.8068 GflopsM N K = 1536 1536 1024, Time = 0.00273072 0.00275611 0.00280208 s, AVG Performance =1632.7386 GflopsM N K = 2048 2048 1024, Time = 0.00487622 0.00488028 0.00488614 s, AVG Performance =1639.2518 GflopsM N K = 3072 3072 1024, Time = 0.01001603 0.01071136 0.01099990 s, AVG Performance =1680.4589 GflopsM N K = 4096 4096 1024, Time = 0.01771046 0.01792170 0.01803462 s, AVG Performance =1785.5450 GflopsM N K = 6144 6144 1024, Time = 0.03988969 0.03993405 0.04000595 s, AVG Performance =1802.9724 GflopsM N K = 8192 8192 1024, Time = 0.07119219 0.07139694 0.07160816 s, AVG Performance =1792.7940 GflopsM N K =1228812288 1024, Time = 0.15978026 0.15993242 0.16043369 s, AVG Performance =1800.7606 GflopsM N K =1638416384 1024, Time = 0.28559187 0.28567238 0.28573316 s, AVG Performance =1792.2629 Gflops
###

下面以M=512,K=512,N=512,為例,詳細(xì)分析一下上述計(jì)算過程的workflow:

  1. 在Global Memory 中分別為矩陣A,B,C分配儲(chǔ)存空間.
  2. 由於矩陣C中每個(gè)元素的計(jì)算均相互獨(dú)立, 因此在並行度映射中讓每個(gè)thread 對(duì)應(yīng)矩陣C中1 個(gè)元素的計(jì)算.
  3. 執(zhí)行配置(execution configuration)中g(shù)ridSize 和blockSize 都有x(列向)、y(行向)兩個(gè)維度, 其中

CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

#nsys 記錄的naive 版本的profiling

CUDA之通用矩陣乘法:從入門到熟練!

二、GEMM的最佳化探究

# ###前文僅在功能上實(shí)現(xiàn)了GEMM,性能上還遠(yuǎn)遠(yuǎn)不及預(yù)期,本節(jié)將主要研究GEMM 性能上的優(yōu)化。 #########2.1 矩陣分塊利用Shared Memory#########上述的計(jì)算需要兩次Global Memory的load才能完成一次乘累加運(yùn)算,計(jì)算訪存比極低,沒有有效的數(shù)據(jù)復(fù)用。所以可以用 Shared Memory 來減少重複的記憶體讀取。 ######先將矩陣C等分成BMxBN大小的分塊,每個(gè)分塊由一個(gè) Block 計(jì)算,其中每個(gè)Thread負(fù)責(zé)計(jì)算矩陣C中的TMxTN個(gè)元素。之後計(jì)算所需的資料全部從 smem 讀取,就消除了一部分重複的A,B矩陣記憶體讀取??紤]到 Shared Memory 容量有限,可以在K維上每次讀取BK大小的分塊,這樣的循環(huán)一共需要K / BK次以完成整個(gè)矩陣乘法操作,即可得到 Block 的結(jié)果。其流程如下圖所示:###############利用Shared Memory 最佳化後,對(duì)每一個(gè)分塊,可得:############# ##由上式可知BM和BN越大,計(jì)算訪存比越高,效能就會(huì)越好。但由於 Shared Memory 容量的限制(V100 1個(gè)SM僅96KB),而一個(gè)Block需要佔(zhàn)用 BK * (BM BN) * 4 Bytes大小。 ###

TM和TN的取值也受到兩方面限制,一方面是線程數(shù)的限制,一個(gè)Block中有BM / TM * BN / TN個(gè)線程,這個(gè)數(shù)字不能超過1024,且不能太高防止影響SM內(nèi)Block間的平行;另一方面是暫存器數(shù)目的限制,一個(gè)執(zhí)行緒至少需要TM * TN個(gè)暫存器來存放矩陣C的部分和,再加上一些其它的暫存器,所有的暫存器數(shù)目不能超過256,且不能太高防止影響SM內(nèi)同時(shí)並行的執(zhí)行緒數(shù)目。

最終選取 BM = BN = 128,BK = 8,TM = TN = 8,則此時(shí)計(jì)算訪存比為32。根據(jù)V100的理論算力15.7TFLOPS,可得 15.7TFLOPS/32 = 490GB/s,根據(jù)實(shí)測(cè)的HBM頻寬為763GB/s,可知此時(shí)頻寬不再會(huì)限制運(yùn)算效能。

根據(jù)上述分析,kernel 函數(shù)實(shí)作過程如下,完整程式碼參考sgemm_v1.cu,主要步驟包括:

CUDA之通用矩陣乘法:從入門到熟練!CUDA之通用矩陣乘法:從入門到熟練!

CUDA之通用矩陣乘法:從入門到熟練!

A B 矩陣分塊的執(zhí)行緒索引關(guān)係

確定好單一block的執(zhí)行過程,接下來需要確定多block處理的不同分塊在Global Memory中的對(duì)應(yīng)關(guān)係,仍以A為例進(jìn)行說明。由於分塊沿著行的方向移動(dòng),那麼首先需要確定行號(hào),根據(jù)Grid 的二維全域線性索引關(guān)係,by * BM?表示該分塊的起始行號(hào),同時(shí)我們已知load_a_smem_m?為分塊內(nèi)部的行號(hào),因此全域的行號(hào)為load_a_gmem_m = by * BM load_a_smem_m?##?##?##?##?##?##?##?##?##?##?##?。由於分塊沿著行的方向移動(dòng),因此列是變化的,需要在循環(huán)內(nèi)部計(jì)算,同樣也是先計(jì)算起始列號(hào)bk * BK?加速分塊內(nèi)部列號(hào)load_a_smem_k

?CUDA之通用矩陣乘法:從入門到熟練!得到CUDA之通用矩陣乘法:從入門到熟練!?

###load_a_gmem_k = bk * BK load_a_smem_k######?###,由此我們便可以確定了分塊在便原始資料中的位置OFFSET(###load_a_gmem_m, load_a_gmem_k###, K) 。同理可分析矩陣分塊 的情況,不再贅述。 ################

計(jì)算完後,還需要將其存入 Global Memory 中,這需要計(jì)算其在 Global Memory 中的對(duì)應(yīng)關(guān)係。由於存在較小的分塊,則行和列都由3部分構(gòu)成:全域行號(hào)store_c_gmem_m?等於大分塊的起始行號(hào)by * BM 小分塊的起始行號(hào)ty * TM 小分塊內(nèi)部的相對(duì)行號(hào)?i?。列同理。

__global__ void sgemm_V1(float * __restrict__ a, float * __restrict__ b, float * __restrict__ c,const int M, const int N, const int K) {const int BM = 128;const int BN = 128;const int BK = 8;const int TM = 8;const int TN = 8;const int bx = blockIdx.x;const int by = blockIdx.y;const int tx = threadIdx.x;const int ty = threadIdx.y;const int tid = ty * blockDim.x + tx;__shared__ float s_a[BM][BK];__shared__ float s_b[BK][BN];float r_c[TM][TN] = {0.0};int load_a_smem_m = tid >> 1;// tid/2, row of s_aint load_a_smem_k = (tid & 1) > 5; // tid/32, row of s_bint load_b_smem_n = (tid & 31) 

計(jì)算結(jié)果如下,效能達(dá)到了理論峰值效能的51.7%:

M N K =128128 1024, Time = 0.00031578 0.00031727 0.00032288 s, AVG Performance =98.4974 GflopsM N K =192192 1024, Time = 0.00031638 0.00031720 0.00031754 s, AVG Performance = 221.6661 GflopsM N K =256256 1024, Time = 0.00031488 0.00031532 0.00031606 s, AVG Performance = 396.4287 GflopsM N K =384384 1024, Time = 0.00031686 0.00031814 0.00032080 s, AVG Performance = 884.0425 GflopsM N K =512512 1024, Time = 0.00031814 0.00032007 0.00032493 s, AVG Performance =1562.1563 GflopsM N K =768768 1024, Time = 0.00032397 0.00034419 0.00034848 s, AVG Performance =3268.5245 GflopsM N K = 1024 1024 1024, Time = 0.00034570 0.00034792 0.00035331 s, AVG Performance =5748.3952 GflopsM N K = 1536 1536 1024, Time = 0.00068797 0.00068983 0.00069094 s, AVG Performance =6523.3424 GflopsM N K = 2048 2048 1024, Time = 0.00136173 0.00136552 0.00136899 s, AVG Performance =5858.5604 GflopsM N K = 3072 3072 1024, Time = 0.00271910 0.00273115 0.00274006 s, AVG Performance =6590.6331 GflopsM N K = 4096 4096 1024, Time = 0.00443805 0.00445964 0.00446883 s, AVG Performance =7175.4698 GflopsM N K = 6144 6144 1024, Time = 0.00917891 0.00950608 0.00996963 s, AVG Performance =7574.0999 GflopsM N K = 8192 8192 1024, Time = 0.01628838 0.01645271 0.01660790 s, AVG Performance =7779.8733 GflopsM N K =1228812288 1024, Time = 0.03592557 0.03597434 0.03614323 s, AVG Performance =8005.7066 GflopsM N K =1638416384 1024, Time = 0.06304122 0.06306373 0.06309302 s, AVG Performance =8118.7715 Gflops

下面仍以M=512,K=512,N=512為例,分析一下結(jié)果。首先透過 profiling 可以看到 Shared Memory 佔(zhàn)用為 8192 bytes,這與理論上(128 128)X8X4完全一致。

CUDA之通用矩陣乘法:從入門到熟練!nsys 記錄的V1 版本的profiling

profiling 顯示Occupancy 為12.5%,可以透過cuda-calculator 加以印證,該例中threads per block = 256, Registers per thread = 136, 由此可以計(jì)算得到每個(gè)SM中活躍的warp 為8,而對(duì)於V100,每個(gè)SM中的warp 總數(shù)為64,因此Occupancy 為8/64 = 12.5%。

CUDA之通用矩陣乘法:從入門到熟練!

2.2 解決Bank Conflict 問題

上節(jié)透過利用Shared Memory 大幅提高了訪存效率,進(jìn)而提高了效能,本節(jié)將進(jìn)一步優(yōu)化Shared Memory 的使用。

Shared Memory共分為32個(gè)Bank,每個(gè)Bank的寬度為4 Bytes,如果需要存取同一個(gè)Bank的多個(gè)數(shù)據(jù),就會(huì)發(fā)生Bank Conflict。例如一個(gè)Warp的32個(gè)線程,如果訪問的地址分別為0、4、8、...、124,就不會(huì)發(fā)生Bank Conflict,只佔(zhàn)用Shared Memory一拍的時(shí)間;如果訪問的地址為0、8 、16、...、248,這樣一來地址0和地址128對(duì)應(yīng)的資料位於同一Bank、地址4和地址132對(duì)應(yīng)的資料位於同一Bank,以此類推,那麼就需要佔(zhàn)用Shared Memory兩拍的時(shí)間才能讀出。

CUDA之通用矩陣乘法:從入門到熟練!

有Bank Conflict VS 無Bank Conflict

再看V1 版本計(jì)算部分的三層循環(huán),每次從Shared memory中取矩陣A的長(zhǎng)度為TM的向量和矩陣B的長(zhǎng)度為TN的向量,這兩個(gè)向量做外積並累加到部分和中,一次外積共TM * TN次乘累加,一共需要循環(huán)BK次取數(shù)和外積。

接下來分析從Shared Memory load的過程中存在的Bank Conflict:

i) 取矩陣A需要取一個(gè)列向量,而矩陣A在Shared Memory中是按行儲(chǔ)存的;

ii) 在TM = TN = 8的情況下,無論矩陣A或矩陣B,從Shared Memory中取數(shù)時(shí)需要取連續(xù)的8個(gè)數(shù),即便用LDS.128指令一條指令取四個(gè)數(shù),也需要兩個(gè)指令,由於一個(gè)執(zhí)行緒的兩個(gè)load指令的位址是連續(xù)的,那麼同一個(gè)Warp不同執(zhí)行緒的同一個(gè)load指令的訪存位址就是被間隔開的,便存在著Bank Conflict 。

為了解決上述的兩點(diǎn)Shared Memory的Bank Conflict,採(cǎi)用了一下兩點(diǎn)優(yōu)化:

i) 為矩陣A分配Shared Memory時(shí)形狀分配為[BK][BM],即讓矩陣A在Shared Memory中按列儲(chǔ)存

ii) 將原本每個(gè)執(zhí)行緒負(fù)責(zé)計(jì)算的TM * TN的矩陣C,分成下圖中這樣的兩塊TM/2 * TN的矩陣C,由於TM/2=4,一條指令即可完成A的一塊的load操作,兩個(gè)load可同時(shí)進(jìn)行。

CUDA之通用矩陣乘法:從入門到熟練!

kernel 函數(shù)的核心部分實(shí)作如下,完整程式碼見 sgemm_v2.cu 。

__shared__ float s_a[BK][BM];__shared__ float s_b[BK][BN];float r_load_a[4];float r_load_b[4];float r_comp_a[TM];float r_comp_b[TN];float r_c[TM][TN] = {0.0};int load_a_smem_m = tid >> 1;int load_a_smem_k = (tid & 1) > 5;int load_b_smem_n = (tid & 31) 

結(jié)果如下,相對(duì)未解 Bank Conflict 版(V1) 效能提高了 14.4%,達(dá)到了理論高峰的74.3%。

M N K =128128 1024, Time = 0.00029699 0.00029918 0.00030989 s, AVG Performance = 104.4530 GflopsM N K =192192 1024, Time = 0.00029776 0.00029828 0.00029882 s, AVG Performance = 235.7252 GflopsM N K =256256 1024, Time = 0.00029485 0.00029530 0.00029619 s, AVG Performance = 423.2949 GflopsM N K =384384 1024, Time = 0.00029734 0.00029848 0.00030090 s, AVG Performance = 942.2843 GflopsM N K =512512 1024, Time = 0.00029853 0.00029945 0.00030070 s, AVG Performance =1669.7479 GflopsM N K =768768 1024, Time = 0.00030458 0.00032467 0.00032790 s, AVG Performance =3465.1038 GflopsM N K = 1024 1024 1024, Time = 0.00032406 0.00032494 0.00032621 s, AVG Performance =6155.0281 GflopsM N K = 1536 1536 1024, Time = 0.00047990 0.00048224 0.00048461 s, AVG Performance =9331.3912 GflopsM N K = 2048 2048 1024, Time = 0.00094426 0.00094636 0.00094992 s, AVG Performance =8453.4569 GflopsM N K = 3072 3072 1024, Time = 0.00187866 0.00188096 0.00188538 s, AVG Performance =9569.5816 GflopsM N K = 4096 4096 1024, Time = 0.00312589 0.00319050 0.00328147 s, AVG Performance = 10029.7885 GflopsM N K = 6144 6144 1024, Time = 0.00641280 0.00658940 0.00703498 s, AVG Performance = 10926.6372 GflopsM N K = 8192 8192 1024, Time = 0.01101130 0.01116194 0.01122950 s, AVG Performance = 11467.5446 GflopsM N K =1228812288 1024, Time = 0.02464854 0.02466705 0.02469344 s, AVG Performance = 11675.4946 GflopsM N K =1638416384 1024, Time = 0.04385955 0.04387468 0.04388355 s, AVG Performance = 11669.5995 Gflops

分析一下profiling 可以看到Static Shared Memory 仍然是使用了8192 Bytes,奇怪的是,Shared Memory executed 卻翻倍變成了16384 Bytes(知友如果知道原因可以告訴我一下)。

CUDA之通用矩陣乘法:從入門到熟練!

2.3 流水並行化:Double Buffering

Double Buffering,即雙重緩衝,即透過增加buffer的方式,使得?存取-計(jì)算?的序列模式管線化,以減少等待時(shí)間,提高運(yùn)算效率,其原理如下圖所示:

CUDA之通用矩陣乘法:從入門到熟練!

Single Buffering VS Double Buffering

具體到GEMM 任務(wù)中來,就是需要兩倍的Shared Memory,之前只需要BK * (BM BN) * 4 Bytes的Shared Memory ,採(cǎi)用Double Buffering之後需要2BK * (BM BN) * 4 Bytes的Shared Memory,然後使其pipeline 流動(dòng)。

程式碼核心部分如下所示,完整程式碼請(qǐng)參考 sgemm_v3.cu 。有以下幾點(diǎn)要注意:

1)主循環(huán)從bk = 1?開始,第一次資料載入在主循環(huán)之前,最後一次計(jì)算在主循環(huán)之後,這是pipeline 的特點(diǎn)決定的;

2)由於計(jì)算和下一次訪問使用的Shared Memory不同,因此主循環(huán)中每次循環(huán)只需要一次__syncthreads()即可

3)由於GPU無法向CPU那樣支援亂序執(zhí)行,主循環(huán)中需要先將下一次迴圈運(yùn)算所需的Gloabal Memory中的資料load 到暫存器,然後進(jìn)行本次計(jì)算,之後再將load到暫存器中的資料寫到Shared Memory,這樣在LDG指令向Global Memory做load時(shí),不會(huì)影響後續(xù)FFMA及其它運(yùn)算指令的launch 執(zhí)行,也就達(dá)到了Double Buffering的目的。

__shared__ float s_a[2][BK][BM];__shared__ float s_b[2][BK][BN];float r_load_a[4];float r_load_b[4];float r_comp_a[TM];float r_comp_b[TN];float r_c[TM][TN] = {0.0};int load_a_smem_m = tid >> 1;int load_a_smem_k = (tid & 1) > 5;int load_b_smem_n = (tid & 31) 

效能如下所示,達(dá)到了理論高峰的 80.6%。

M N K =128128 1024, Time = 0.00024000 0.00024240 0.00025792 s, AVG Performance = 128.9191 GflopsM N K =192192 1024, Time = 0.00024000 0.00024048 0.00024125 s, AVG Performance = 292.3840 GflopsM N K =256256 1024, Time = 0.00024029 0.00024114 0.00024272 s, AVG Performance = 518.3728 GflopsM N K =384384 1024, Time = 0.00024070 0.00024145 0.00024198 s, AVG Performance =1164.8394 GflopsM N K =512512 1024, Time = 0.00024173 0.00024237 0.00024477 s, AVG Performance =2062.9786 GflopsM N K =768768 1024, Time = 0.00024291 0.00024540 0.00026010 s, AVG Performance =4584.3820 GflopsM N K = 1024 1024 1024, Time = 0.00024534 0.00024631 0.00024941 s, AVG Performance =8119.7302 GflopsM N K = 1536 1536 1024, Time = 0.00045712 0.00045780 0.00045872 s, AVG Performance =9829.5167 GflopsM N K = 2048 2048 1024, Time = 0.00089632 0.00089970 0.00090656 s, AVG Performance =8891.8924 GflopsM N K = 3072 3072 1024, Time = 0.00177891 0.00178289 0.00178592 s, AVG Performance = 10095.9883 GflopsM N K = 4096 4096 1024, Time = 0.00309763 0.00310057 0.00310451 s, AVG Performance = 10320.6843 GflopsM N K = 6144 6144 1024, Time = 0.00604826 0.00619887 0.00663078 s, AVG Performance = 11615.0253 GflopsM N K = 8192 8192 1024, Time = 0.01031738 0.01045051 0.01048861 s, AVG Performance = 12248.2036 GflopsM N K =1228812288 1024, Time = 0.02283978 0.02285837 0.02298272 s, AVG Performance = 12599.3212 GflopsM N K =1638416384 1024, Time = 0.04043287 0.04044823 0.04046151 s, AVG Performance = 12658.1556 Gflops

從profiling 可以看到雙倍的Shared Memory 的佔(zhàn)用

CUDA之通用矩陣乘法:從入門到熟練!

三、cuBLAS 實(shí)作方式探究

本節(jié)我們將認(rèn)識(shí)CUDA的標(biāo)準(zhǔn)函式庫-cuBLAS,即NVIDIA版本的基本線性代數(shù)子程式(Basic Linear Algebra Subprograms, BLAS) 規(guī)範(fàn)實(shí)作程式碼。它支援 Level 1 (向量與向量運(yùn)算) ,Level 2 (向量與矩陣運(yùn)算) ,Level 3 (矩陣與矩陣運(yùn)算) 層級(jí)的標(biāo)準(zhǔn)矩陣運(yùn)算。

CUDA之通用矩陣乘法:從入門到熟練!

cuBLAS/CUTLASS GEMM的基本過程

#如上圖所示,計(jì)算過程分解成線程區(qū)塊片(thread block tile )、線程束片(warp tile)和線程片(thread tile)的層次結(jié)構(gòu)並將AMP的策略應(yīng)用於此層次結(jié)構(gòu)來高效率的完成基於GPU的拆分成tile的GEMM。這個(gè)層次結(jié)構(gòu)緊密地反映了NVIDIA CUDA程式設(shè)計(jì)模型。可以看到從global memory到shared memory的資料移動(dòng)(矩陣到thread block tile);從shared memory到暫存器的資料移動(dòng)(thread block tile到warp tile);從暫存器到CUDA core的計(jì)算(warp tile到thread tile )。

cuBLAS 實(shí)作了單精確度矩陣乘的函數(shù)cublasSgemm,其主要參數(shù)如下:

cublasStatus_t cublasSgemm( cublasHandle_t handle, // 調(diào)用 cuBLAS 庫時(shí)的句柄 cublasOperation_t transa, // A 矩陣是否需要轉(zhuǎn)置 cublasOperation_t transb, // B 矩陣是否需要轉(zhuǎn)置 int m, // A 的行數(shù) int n, // B 的列數(shù) int k, // A 的列數(shù) const float *alpha, // 系數(shù) α, host or device pointer const float *A, // 矩陣 A 的指針,device pointer int lda, // 矩陣 A 的主維,if A 轉(zhuǎn)置, lda = max(1, k), else max(1, m) const float *B, // 矩陣 B 的指針, device pointer int ldb, // 矩陣 B 的主維,if B 轉(zhuǎn)置, ldb = max(1, n), else max(1, k) const float *beta, // 系數(shù) β, host or device pointer float *C, // 矩陣 C 的指針,device pointer int ldc // 矩陣 C 的主維,ldc >= max(1, m) );

呼叫方式如下:

cublasHandle_t cublas_handle;cublasCreate(&cublas_handle);float cublas_alpha = 1.0;float cublas_beta = 0;cublasSgemm(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, &cublas_alpha, d_b, N, d_a, K, &cublas_beta, d_c, N);

性能如下所示,達(dá)到了理論峰值的82.4%。

M N K =128128 1024, Time = 0.00002704 0.00003634 0.00010822 s, AVG Performance = 860.0286 GflopsM N K =192192 1024, Time = 0.00003155 0.00003773 0.00007267 s, AVG Performance =1863.6689 GflopsM N K =256256 1024, Time = 0.00003917 0.00004524 0.00007747 s, AVG Performance =2762.9438 GflopsM N K =384384 1024, Time = 0.00005318 0.00005978 0.00009120 s, AVG Performance =4705.0655 GflopsM N K =512512 1024, Time = 0.00008326 0.00010280 0.00013840 s, AVG Performance =4863.9646 GflopsM N K =768768 1024, Time = 0.00014278 0.00014867 0.00018816 s, AVG Performance =7567.1560 GflopsM N K = 1024 1024 1024, Time = 0.00023485 0.00024460 0.00028150 s, AVG Performance =8176.5614 GflopsM N K = 1536 1536 1024, Time = 0.00046474 0.00047607 0.00051181 s, AVG Performance =9452.3201 GflopsM N K = 2048 2048 1024, Time = 0.00077930 0.00087862 0.00092307 s, AVG Performance =9105.2126 GflopsM N K = 3072 3072 1024, Time = 0.00167904 0.00168434 0.00171114 s, AVG Performance = 10686.6837 GflopsM N K = 4096 4096 1024, Time = 0.00289619 0.00291068 0.00295904 s, AVG Performance = 10994.0128 GflopsM N K = 6144 6144 1024, Time = 0.00591766 0.00594586 0.00596915 s, AVG Performance = 12109.2611 GflopsM N K = 8192 8192 1024, Time = 0.01002384 0.01017465 0.01028435 s, AVG Performance = 12580.2896 GflopsM N K =1228812288 1024, Time = 0.02231159 0.02233805 0.02245619 s, AVG Performance = 12892.7969 GflopsM N K =1638416384 1024, Time = 0.03954650 0.03959291 0.03967242 s, AVG Performance = 12931.6086 Gflops

由此可以對(duì)比以上各種方法的性能情況,可見手動(dòng)實(shí)現(xiàn)的性能已接近官方的性能,如下:

CUDA之通用矩陣乘法:從入門到熟練!


#

以上是CUDA之通用矩陣乘法:從入門到熟練!的詳細(xì)內(nèi)容。更多資訊請(qǐng)關(guān)注PHP中文網(wǎng)其他相關(guān)文章!

本網(wǎng)站聲明
本文內(nèi)容由網(wǎng)友自願(yuàn)投稿,版權(quán)歸原作者所有。本站不承擔(dān)相應(yīng)的法律責(zé)任。如發(fā)現(xiàn)涉嫌抄襲或侵權(quán)的內(nèi)容,請(qǐng)聯(lián)絡(luò)admin@php.cn

熱AI工具

Undress AI Tool

Undress AI Tool

免費(fèi)脫衣圖片

Undresser.AI Undress

Undresser.AI Undress

人工智慧驅(qū)動(dòng)的應(yīng)用程序,用於創(chuàng)建逼真的裸體照片

AI Clothes Remover

AI Clothes Remover

用於從照片中去除衣服的線上人工智慧工具。

Clothoff.io

Clothoff.io

AI脫衣器

Video Face Swap

Video Face Swap

使用我們完全免費(fèi)的人工智慧換臉工具,輕鬆在任何影片中換臉!

熱工具

記事本++7.3.1

記事本++7.3.1

好用且免費(fèi)的程式碼編輯器

SublimeText3漢化版

SublimeText3漢化版

中文版,非常好用

禪工作室 13.0.1

禪工作室 13.0.1

強(qiáng)大的PHP整合開發(fā)環(huán)境

Dreamweaver CS6

Dreamweaver CS6

視覺化網(wǎng)頁開發(fā)工具

SublimeText3 Mac版

SublimeText3 Mac版

神級(jí)程式碼編輯軟體(SublimeText3)

熱門話題

Laravel 教程
1601
29
PHP教程
1502
276
.NET Core快速入門教程 1、開篇:說說.NET Core的那些事兒 .NET Core快速入門教程 1、開篇:說說.NET Core的那些事兒 May 07, 2025 pm 04:54 PM

一、.NETCore的起源談到.NETCore,就不能不提它的前身.NET。當(dāng)年Java風(fēng)頭正盛,微軟也對(duì)Java青睞有加,Windows平臺(tái)上的Java虛擬機(jī)就是微軟依據(jù)JVM標(biāo)準(zhǔn)開發(fā)的,據(jù)稱是當(dāng)時(shí)性能最佳的Java虛擬機(jī)。然而,微軟有自己的小算盤,試圖將Java與Windows平臺(tái)捆綁,增加一些Windows特有的功能。 Sun公司對(duì)此不滿,導(dǎo)致雙方關(guān)係破裂,微軟隨後推出了.NET。 .NET從誕生之初就借鑒了Java的許多特性,並在語言特性和窗體開發(fā)等方面逐漸超越了Java。 Java在1.6版

如何理解C  中的ABI兼容性? 如何理解C 中的ABI兼容性? Apr 28, 2025 pm 10:12 PM

C 中的ABI兼容性是指不同編譯器或版本生成的二進(jìn)制代碼能否在不重新編譯的情況下兼容。 1.函數(shù)調(diào)用約定,2.名稱修飾,3.虛函數(shù)表佈局,4.結(jié)構(gòu)體和類的佈局是主要涉及的方面。

H5頁面製作是前端開發(fā)嗎 H5頁面製作是前端開發(fā)嗎 Apr 05, 2025 pm 11:42 PM

是的,H5頁面製作是前端開發(fā)的重要實(shí)現(xiàn)方式,涉及HTML、CSS和JavaScript等核心技術(shù)。開發(fā)者通過巧妙結(jié)合這些技術(shù),例如使用&lt;canvas&gt;標(biāo)籤繪製圖形或使用JavaScript控制交互行為,構(gòu)建出動(dòng)態(tài)且功能強(qiáng)大的H5頁面。

C  中的chrono庫如何使用? C 中的chrono庫如何使用? Apr 28, 2025 pm 10:18 PM

使用C 中的chrono庫可以讓你更加精確地控制時(shí)間和時(shí)間間隔,讓我們來探討一下這個(gè)庫的魅力所在吧。 C 的chrono庫是標(biāo)準(zhǔn)庫的一部分,它提供了一種現(xiàn)代化的方式來處理時(shí)間和時(shí)間間隔。對(duì)於那些曾經(jīng)飽受time.h和ctime折磨的程序員來說,chrono無疑是一個(gè)福音。它不僅提高了代碼的可讀性和可維護(hù)性,還提供了更高的精度和靈活性。讓我們從基礎(chǔ)開始,chrono庫主要包括以下幾個(gè)關(guān)鍵組件:std::chrono::system_clock:表示系統(tǒng)時(shí)鐘,用於獲取當(dāng)前時(shí)間。 std::chron

如何通過JavaScript或CSS控制瀏覽器打印設(shè)置中的頁首和頁尾? 如何通過JavaScript或CSS控制瀏覽器打印設(shè)置中的頁首和頁尾? Apr 05, 2025 pm 10:39 PM

如何使用JavaScript或CSS控制瀏覽器打印設(shè)置中的頁首和頁尾在瀏覽器的打印設(shè)置中,有一個(gè)選項(xiàng)可以控制是否顯?...

在移動(dòng)端如何兼容多行溢出省略? 在移動(dòng)端如何兼容多行溢出省略? Apr 05, 2025 pm 10:36 PM

移動(dòng)端多行溢出省略在不同設(shè)備上的兼容問題在使用Vue2.0開發(fā)移動(dòng)端應(yīng)用時(shí),常常會(huì)遇到需要對(duì)文本進(jìn)行多行溢...

為什麼inline-block元素會(huì)出現(xiàn)錯(cuò)位現(xiàn)象?如何解決這個(gè)問題? 為什麼inline-block元素會(huì)出現(xiàn)錯(cuò)位現(xiàn)象?如何解決這個(gè)問題? Apr 04, 2025 pm 10:39 PM

關(guān)於inline-block元素錯(cuò)位顯示的原因及解決方案在編寫網(wǎng)頁佈局時(shí),我們常常會(huì)遇到一些看似奇怪的顯示問題。比...

Bootstrap列表如何改變大??? Bootstrap列表如何改變大小? Apr 07, 2025 am 10:45 AM

Bootstrap 列表的大小取決於包含列表的容器的大小,而不是列表本身。使用 Bootstrap 的網(wǎng)格系統(tǒng)或 Flexbox 可以控制容器的大小,從而間接調(diào)整列表項(xiàng)的大小。

See all articles