GPU 內(nèi)存的分級（gpu memory hierarchy）

小普 中科院化學(xué)所在讀博士研究生

研究課題，計算機模擬并行軟件的開發(fā)與應(yīng)用

Email:? yaopu2019@126.com （歡迎和我討論問題）

CSDN：

博客園：

摘要(Abstact)

GPU 的存儲是多樣化的, 其速度和數(shù)量并不相同，了解GPU存儲對于程序的性能調(diào)優(yōu)有著重要的意義。本文介紹如下幾個問題:

1.內(nèi)存類型有什么？2)查詢自己設(shè)備的內(nèi)存大小 3）內(nèi)存訪問速度4）不同級別的存儲關(guān)系5）使用注意事項。各種存儲結(jié)構(gòu)的優(yōu)缺點。

?

正文

GPU結(jié)構(gòu)圖

?

?

①寄存器內(nèi)存（Register memory）

優(yōu)點：訪問速度的冠軍！

缺點：數(shù)量有限

使用：在__global__函數(shù) ，或者_(dá)__device__ 函數(shù)內(nèi)，定義的普通變量，就是寄存器變量。

例子：

//kernel.cu__global__ void register_test(){int a = 1.0;double b = 2.0;}//main.cuint main(){int nBlock = 100;register_test <<<nBlock,128>>>();return 0;}

?

②共享內(nèi)存（Shared memory）

優(yōu)點：

1緩存速度快 比全局內(nèi)存 快2兩個數(shù)量級

2 線程塊內(nèi)，所有線程可以讀寫。

?3 生命周期與線程塊同步

缺點：大小有限制

使用：關(guān)鍵詞 __shared__ ?如 __shared__ double A[128];

適用條件：

使用場合，如規(guī)約求和 : a = sum A[i]

如果不是頻繁修改的變量，比如矢量加法。

是編程優(yōu)化中的重要手段！

C[i] = A[i] + B[i] 則沒有必要將A,B進行緩存到shared memory 中。

?

//kernel.cu__global__ void shared_test(){__shared__ double A[128];int a = 1.0;double b = 2.0;int tid = threadIdx.x;A[tid] = a;}

?

③全局內(nèi)存 (Global Memory)

優(yōu)點：

1空間最大（GB級別）

2.可以通過cudaMemcpy 等與Host端，進行交互。

3.生命周期比Kernel函數(shù)長

4.所有線程都能訪問

缺點：訪存最慢

//kernel.cu__global__ void shared_test(int *B){double b = 2.0;int tid = threadIdx.x;int id = blockDim.x*128 + threadIdx.x;int a = B[id] ;}

?

④紋理內(nèi)存

優(yōu)點，比普通的global memory 快

缺點：使用起來，需要四個步驟，麻煩一點

適用場景：比較大的只需要讀取array，采用紋理方式訪問，會實現(xiàn)加速

使用的四個步驟（這里以1維float數(shù)組為例子），初學(xué)者，自己手敲一遍代碼！！！

第一步，聲明紋理空間,全局變量：

texture<float, 1, cudaReadModeElementType> tex1D_load;

第二步,綁定紋理

?

?

聲明語句：

#include <iostream>#include <time.h>#include <assert.h>#include <cuda_runtime.h>#include "helper_cuda.h"#include <iostream>#include <ctime>#include <stdio.h>using namespace std;texture<float, 1, cudaReadModeElementType> tex1D_load;//第一步，聲明紋理空間,全局變量__global__ void kernel(float *d_out, int size){//tex1D_load 為全局變量，不在參數(shù)表中int index;index = blockIdx.x * blockDim.x + threadIdx.x;if (index < size){d_out[index] = tex1Dfetch(tex1D_load, index); //第三步，抓取紋理內(nèi)存的值//從紋理中抓取值printf("%f\n", d_out[index]);}}int main(){int size = 120;size_t Size = size * sizeof(float);float *harray;float *d_in;float *d_out;harray = new float[size];checkCudaErrors(cudaMalloc((void **)&d_out, Size));checkCudaErrors(cudaMalloc((void **)&d_in, Size));//initial host memoryfor (int m = 0; m < 4; m++){printf("m = %d\n", m);for (int loop = 0; loop < size; loop++){harray[loop] = loop + m * 1000;}//拷貝到d_in中checkCudaErrors(cudaMemcpy(d_in, harray, Size, cudaMemcpyHostToDevice));//第二步,綁定紋理checkCudaErrors(cudaBindTexture(0, tex1D_load, d_in, Size));//0表示沒有偏移int nBlocks = (Size - 1) / 128 + 1;kernel<<<nBlocks, 128>>>(d_out, size); //第三步cudaUnbindTexture(tex1D_load);???????? //第四，解紋理getLastCudaError("Kernel execution failed");checkCudaErrors(cudaDeviceSynchronize());}delete[] harray;cudaUnbindTexture(&tex1D_load);checkCudaErrors(cudaFree(d_in));checkCudaErrors(cudaFree(d_out));return 0;}

?

總結(jié)如下表

?

結(jié)束語

小普 中科院化學(xué)所在讀博士研究生

研究課題，計算機模擬并行軟件的開發(fā)與應(yīng)用

Email:? yaopu2019@126.com （歡迎和我討論問題，私信和郵件都OK!）

讓程序使得更多人受益！

參考文獻(xiàn)

CUDA專家手冊 GPU編程權(quán)威指南 [M] 2014

CUDA Toolkit Documentation v10.1.168 https://docs.nvidia.com/cuda/

總結(jié)

以上是生活随笔為你收集整理的GPU 内存的分级综述（gpu memory hierarchy）的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。