本文參加2022CUDA on Platform線上訓(xùn)練營學(xué)習(xí)筆記
歡迎各位
大犇提意見
一、矩陣轉(zhuǎn)置(Matrix Transpose)基礎(chǔ)
上圖中將m * n的矩陣A通過矩陣轉(zhuǎn)置變成了n * m的 AT,簡單來講矩陣轉(zhuǎn)置即為將原始矩陣的第一行轉(zhuǎn)置為目標(biāo)矩陣的第一列，以此類推,相信基礎(chǔ)扎實的你簡單地看看CPU端的代碼就能理解
二、矩陣轉(zhuǎn)置的CPU端實現(xiàn)
__host__ void cpu_transpose(int *matrix,int *tr_matrix,int m,int n) {
for (int i = 0; i < n; i++) {
for (int j = 0; j < m; j++) {
tr_matrix[i * m + j] = matrix[j * n + i];
}
}
return;
}
定義一個名為cpu_transpose的函數(shù)，將矩陣matrix轉(zhuǎn)置為矩陣tr_matrix，通過觀察代碼不難發(fā)現(xiàn)tr_matrix[i][j]=matrix[j][i],這里需要注意到的是坐標(biāo)的轉(zhuǎn)換,轉(zhuǎn)置后的矩陣行數(shù)和列數(shù)發(fā)生變換，留意m和n不要乘錯了。
原始矩陣:
CPU端執(zhí)行結(jié)果:
三、矩陣轉(zhuǎn)置的GPU端實現(xiàn)(share Memory)
1、核函數(shù)的編寫
GPU端的實現(xiàn)與CPU端類似，首先根據(jù)各個線程的index(索引)計算出當(dāng)前線程在原始矩陣中的位置row和col,在原始矩陣中的row行，col列
int row = blockDim.y * blockIdx.y + threadIdx.y;
int col = blockDim.x * blockIdx.x + threadIdx.x;
下邊我們申請同一個block中的線程可以訪問的shared Memory
__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
在GPU中申請了一塊名為smem_matrix大小為sizeof(int)*BLOCK_SIZE^2的共享內(nèi)存，在執(zhí)行賦值操作之前將當(dāng)前block中的線程需要訪問到的數(shù)據(jù)從Global_Memory中復(fù)制到share_Memory
smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
賦值時需要注意的是:由于我們?yōu)閮?nèi)核函數(shù)設(shè)置執(zhí)行配置的時候通常會向上取整，會申請多于實際需求的線程數(shù)，所以在我們賦值之前需要判斷當(dāng)前線程的坐標(biāo)是否是需求坐標(biāo)，以此來防止訪問matrix時row*n+col成為野指針，對我們的數(shù)據(jù)造成重大的危害
有了同一個block中的線程申請一個share Memory的概念后，需要做的是同步同一個BLock中的線程
__syncthreads();
通過上邊一系列的操作，我們就可以開始真正的轉(zhuǎn)置操作了,需要注意的是，我們已經(jīng)把線程所需的數(shù)據(jù)賦值到share Memory當(dāng)中,所以我們在賦值時只需調(diào)用smem_matrix,同樣,賦值操作之前，我們需要判斷當(dāng)前的坐標(biāo)是否實際有效
if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
上述分析使我們獲得了完整的GPU代碼
__global__ void cuda_transpose(int *matrix,int *tr_matrix,int m,int n) {
int row = blockDim.y * blockIdx.y + threadIdx.y;
int col = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
__syncthreads();
if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
return;
}
2、核函數(shù)的啟動
在設(shè)備端申請兩個指針并為其分配內(nèi)存
int* d_matrix, *dtr_matrix;
cudaMalloc((void**)&d_matrix, sizeof(int) * m * n);
cudaMalloc((void**)&dtr_matrix, sizeof(int) * m * n);
手動將matrix中的數(shù)據(jù)通過Pcie復(fù)制到設(shè)備端的Global Memory當(dāng)中
cudaMemcpy(d_matrix, matrix, sizeof(int) * m * n, cudaMemcpyHostToDevice);
核函數(shù)執(zhí)行設(shè)置的設(shè)定,一個warp通常為32個線程所以我們一個Block中的線程數(shù)最好設(shè)置為32的整數(shù)倍，從此提高使用率，有效防止inactive code的出現(xiàn)
dim3 block = { BLOCK_SIZE,BLOCK_SIZE,1 }; //BLOCK_SIZE = 16
gridDim的設(shè)置最需關(guān)注的就是申請的線程能夠有效的覆蓋真?zhèn)€矩陣，寧可多申請，通過核函數(shù)中的if屏蔽，也不少申請，導(dǎo)致計算的缺失，所以我們在計算中采用向上取整的方法
需要注意的使 dim3 類型中的三個成員都是要求unsigned int 類型的所以我們在前面添加(unsigned int)來強制將我們的計算結(jié)果轉(zhuǎn)換為無符號
dim3 gird = { (unsigned int)(n - 1 + BLOCK_SIZE) / BLOCK_SIZE, (unsigned int)(m - 1 + BLOCK_SIZE) / BLOCK_SIZE,1 };
核函數(shù)啟動！
cuda_transpose << < gird , block >> > (d_matrix, dtr_matrix, m, n);
3、核函數(shù)性能計數(shù)
在CUDA中有一種特殊的類型cudaEvent_t，可以幫助我們記錄核函數(shù)的執(zhí)行信息
cudaEvent_t kernel_start;
cudaEvent_t kernel_end;
cudaEventCreate(&kernel_start);
cudaEventCreate(&kernel_end);
kernel_start用于記錄核函數(shù)開始執(zhí)行時的信息，kernel_end用來記錄核函數(shù)運行結(jié)束時的信息，這里使用到了兩個函數(shù)cudaEventQuery(kernel_start);,cudaEventSynchronize(kernel_end);,前者是非阻塞的，只要執(zhí)行到就直接記錄，后者是阻塞式的，需要前面的執(zhí)行完畢才能運行,具體的性能計數(shù)函數(shù)如下
通過簡單的邏輯組合，就可以得到核函數(shù)的實際運行時間，具體代碼如下
cudaEventCreate(&kernel_start);
cudaEventCreate(&kernel_end);
cudaEventRecord(kernel_start);
cudaEventQuery(kernel_start);
cuda_transpose << < gird , block >> > (d_matrix, dtr_matrix, m, n);
cudaEventRecord(kernel_end);
cudaEventSynchronize(kernel_end);
float ms;
cudaEventElapsedTime(&ms, kernel_start, kernel_end);
四、代碼參考
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <stdio.h>
#include <stdlib.h>
#include <string>
#include <iostream>
#define BLOCK_SIZE 32
using namespace std;
__global__ void cuda_transpose(int *matrix,int *tr_matrix,int m,int n) {
int row = blockDim.y * blockIdx.y + threadIdx.y;
int col = blockDim.x * blockIdx.x + threadIdx.x;
__shared__ int smem_matrix[BLOCK_SIZE][BLOCK_SIZE];
smem_matrix[threadIdx.y][threadIdx.x] = row < m&& col < n ? matrix[row*n+col] : 0;
__syncthreads();
if(blockIdx.x * blockDim.x + threadIdx.y < n && threadIdx.x + blockIdx.y * blockDim.x < m)
tr_matrix[threadIdx.x+blockIdx.y*blockDim.x+m*(blockIdx.x*blockDim.x+threadIdx.y)] = smem_matrix[threadIdx.x][threadIdx.y];
return;
}
__host__ void cpu_transpose(int *matrix,int *tr_matrix,int m,int n) {
for (int i = 0; i < n; i++) {
for (int j = 0; j < m; j++) {
tr_matrix[i * m + j] = matrix[j * n + i];
}
}
return;
}
__host__ void init_matrix(int* matrix,int m,int n) {
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
matrix[i*n+j] = rand();
}
}
}
void print(int*, string,int,int);
bool check(int*, int*, int, int);
int main() {
int m = 1111;
int n = 113;
int *matrix;
cudaMallocHost((void**)&matrix, sizeof(int) * m * n);
init_matrix(matrix,m,n);
//print(matrix, "init matrix", m, n);
int* htr_matrix;
cudaMallocHost((void**)&htr_matrix, sizeof(int) * m * n);
cpu_transpose(matrix, htr_matrix, m, n);
//print(htr_matrix, "CPU", n, m);
//將CPU端執(zhí)行的結(jié)果存放在htr_matrix中
int* d_matrix, *dtr_matrix;
cudaMalloc((void**)&d_matrix, sizeof(int) * m * n);
cudaMalloc((void**)&dtr_matrix, sizeof(int) * m * n);
cudaMemcpy(d_matrix, matrix, sizeof(int) * m * n, cudaMemcpyHostToDevice);
dim3 gird = { (unsigned int)(n - 1 + BLOCK_SIZE) / BLOCK_SIZE, (unsigned int)(m - 1 + BLOCK_SIZE) / BLOCK_SIZE,1 };
dim3 block = { BLOCK_SIZE,BLOCK_SIZE,1 };
cudaEvent_t kernel_start;
cudaEvent_t kernel_end;
cudaEventCreate(&kernel_start);
cudaEventCreate(&kernel_end);
cudaEventRecord(kernel_start);
cudaEventQuery(kernel_start);
cuda_transpose << < gird , block >> > (d_matrix, dtr_matrix, m, n);
cudaEventRecord(kernel_end);
cudaEventSynchronize(kernel_end);
float ms;
cudaEventElapsedTime(&ms, kernel_start, kernel_end);
int* hdtr_matrix;
cudaMallocHost((void**)&hdtr_matrix, sizeof(int) * m * n);
cudaMemcpy(hdtr_matrix, dtr_matrix, sizeof(int) * m * n, cudaMemcpyDeviceToDevice);
//print(hdtr_matrix, "GPU", n, m);
if (check(hdtr_matrix, htr_matrix, n, m)) {
cout << "pass\n";
}
else {
cout << "error\n";
}
printf("GPU time is : %f \n", ms);
cudaFree(hdtr_matrix);
cudaFree(dtr_matrix);
cudaFree(matrix);
cudaFree(htr_matrix);
cudaFree(d_matrix);
return 0;
}
void print(int* a, string name,int m,int n) {
cout << "NAME : " << name << endl;
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
printf("%6d ", a[i * n + j]);
}
printf("\n");
}
}
bool check(int* a, int* b, int m, int n) {
bool check_flag = true;
for (int i = 0; i < m; i++) {
for (int j = 0; j < n; j++) {
if (a[i * n + j] != b[i * n + j]) {
return false;
}
}
}
return check_flag;
}
執(zhí)行結(jié)果如圖
五、實踐心得
本次實踐通過GPU端中的share Memory對核函數(shù)運行時的讀寫問題做了優(yōu)化，當(dāng)線程與線程之間為連續(xù)讀寫時，global Memory的效率是比較高的，不使用share Memory時，使用GPU進(jìn)行矩陣轉(zhuǎn)置會出現(xiàn)兩難問題（1.讀row-major 寫col-major，2寫col-major 讀row-major），而在share Memory中row-major和col-major的效率幾乎相同，很好地解決了global memory上的問題，在編寫過程中，需要注意的是，要順著global memory寫，首先保證global memory讀寫時是row-major，以達(dá)到最高的優(yōu)化效率。
遇到的最大問題是，邊界的判斷問題，GPU轉(zhuǎn)置過程中，由于要保證global memory 是 row-major，所以坐標(biāo)不像是CPU端中的簡單調(diào)換，具體表現(xiàn)為（在對share 數(shù)字賦值時該線程無意義，而在寫global操作中該線程有意義），所以在__syncthreads();后需要判斷當(dāng)前線程是否有意義文章來源：http://www.zghlxwxcb.cn/news/detail-532243.html
鄙人第一次寫實操博客，有建議必洗耳恭聽
再次感謝偉大的NV 開發(fā)者社區(qū)文章來源地址http://www.zghlxwxcb.cn/news/detail-532243.html
到了這里，關(guān)于CUDA:矩陣轉(zhuǎn)置的GPU實現(xiàn)(Share Memory)的文章就介紹完了。如果您還想了解更多內(nèi)容，請在右上角搜索TOY模板網(wǎng)以前的文章或繼續(xù)瀏覽下面的相關(guān)文章，希望大家以后多多支持TOY模板網(wǎng)！
