CUDA ---- Memory Access

時間  2019-11-16
標籤 cuda memory access 简体版
原文   原文鏈接

Memory Access Patterns

大部分device一開始從global Memory獲取數據，並且，大部分GPU應用表現會被帶寬限制。所以最大化應用對global Memory帶寬的使用時獲取高性能的第一步。也就是說，global Memory的使用就沒調節好，其它的優化方案也獲取不到什麼大效果,下面的內容會涉及到很多L1的知識，這部分了解下就好，L1在Maxwell以後就不用了，可是cache的知識點是不變的。程序員

Aligned and Coalesced Access

以下圖所示，global Memory的load/store要經由cache，全部的數據會初始化在DRAM，也就是物理的device Memory上，而kernel可以獲取的global Memory其實是一塊邏輯內存空間。Kernel對Memory的請求都是由DRAM和SM的片上內存以128-byte和32-byte傳輸解決的。算法

 

全部獲取global Memory都要通過L2 cache，也有許多還要通過L1 cache，主要由GPU的架構和獲取模式決定的。若是L1和L2都被使用，那麼Memory的獲取是以128-byte爲單位傳輸的，若是隻使用L2，則以32-byte爲單位傳輸，在容許使用L1的GPU中（Maxwell已經完全不使用L1，本來走L1都換成走texture cache），L1是能夠在編譯期被顯示使用或禁止的。編程

由上文可知，L1 cache中每一行是128bytes，這些數據映射到device Memory上的128位對齊的塊。若是warp中每一個thread請求一個4-byte的值，那麼每次請求會要求獲取128 bytes值，正好契合cache line大小和device Memory segment大小。數組

所以，咱們在設計代碼的時候，有兩個特徵須要注意：數據結構

  1. Aligned Memory access 對齊
  2. Coalesced Memory access 連續

當要獲取的Memory首地址是cache line的倍數時，就是Aligned Memory Access，若是是非對齊的，就會致使浪費帶寬。至於Coalesced Memory Access則是warp的32個thread請求的是連續的內存塊。架構

下圖就是很好的符合了連續和對齊原則，只有128-byte Memory傳輸的消耗：ide

 

下圖則沒有遵照連續和對齊原則，有三次傳輸消耗發生，一次是從偏移地址0開始，一次是從偏移地址256開始，還有一次是從偏移128開始，而此次包含了大部分須要的數據，另外兩次則有不少數據並非須要的，而致使帶寬浪費。函數

 

通常來說，咱們應該這樣優化傳輸效率：使用最少的傳輸次數來知足最大的獲取內存請求。固然，須要多少傳輸，多大的吞吐都是跟CC有關的。oop

Global Memory Reads

在SM中，數據運送是要通過下面三種cache/buffer的，主要依賴於要獲取的device Memory種類：性能

  1. L1/L2 cache
  2. Constant cache
  3. Read-only cache

L1/L2是默認路徑，另外兩條路須要應用顯示的說明，通常這樣作都是爲了提高性能（寫CUDA代碼的時候，能夠先都使用global Memory，而後根據須要慢慢調節，使用一些特殊的內存來提高性能）。Global Memory的load操做是否通過L1cache能夠有下面兩個因素決定：

  1. Device compute capability
  2. Compiler options

默認狀況下，L1是被開啓的，-Xptxas -dlcm=cg能夠用來禁用L1。L1被禁用後，全部去L1的都直接去L2了。當L2未命中時，就直接去DRAM。全部Memory transaction可能請求一個，兩個或者四個segment，每一個segment是32 bytes。固然L1也能夠被顯式的開啓-Xptxas -dlcm=ca，此時，全部Memory請求都先走L1，未命中則去L2。在Kepler K10，K20和K20x系列GPU，L1不在用來cache global Memory，L1的惟一用途就是來cache因爲register spill放到local Memory的那部分register。

Cache Loads

咱們以默認開啓L1爲例，說明下對齊和連續，下圖是理想的狀況，連續且對齊，warp中全部thread的Memory請求都落在同一塊cache line（128 bytes），只有一次傳輸消耗，沒有任何多餘的數據被傳輸，bus使用效率百分百。

 

下圖是對齊但線程ID和地址不是連續一一對應的狀況，不過因爲全部數據仍然在一個連續對齊的塊中，全部依然沒有額外的傳輸消耗，咱們仍然只須要一次128 bytes的傳輸就能完成。

 

下圖則是非連續未對齊的狀況，數據落在了兩個128-byte的塊中，因此就有兩個128-byte的傳輸消耗，而其中有一半是無效數據，bus使用是百分之五十。

 

下圖是最壞的狀況，一樣是請求32個4 bytes數據，可是每一個地址分佈的至關不規律，咱們只想要須要的那128 bytes數據，可是，實際上下圖這樣的分佈，卻須要N∈(0,32)個cache line，也就是N次數據傳輸消耗。

 

CPU的L1 cache是根據時間和空間局部性作出的優化，可是GPU的L1僅僅被設計成針對空間局部性而不包括時間局部性。頻繁的獲取L1不會致使某些數據駐留在cache中，只要下次用不到，直接刪。

Uncached Loads

這裏就是指不走L1可是仍是要走L2，也就是cache line從128-byte變爲32-byte了。依然以上文warp 32個thread每一個4 bytes請求，總計128 bytes爲例，下圖是理想的對齊且連續情形，全部的128 bytes都落在四塊32 bytes的塊中。

 

下圖請求沒有對齊，請求落在了160-byte範圍內，bus有效使用率是百分之八十，相對使用L1，性能要好很多。

 

下圖是全部thread都請求同一塊數據的情形，bus有效使用率爲4bytes/32bytes=12.5%，依然要比L1表現好。

 

下圖是狀況最糟糕的，數據很是分散，可是因爲所請求的128 bytes落在了多個以32 bytes爲單位的segment中，所以無效的數據傳輸要少的多。

 

Example of Misaligned Reads

內存獲取模式通常都是有應用的實現和算法來決定的，一些狀況下，要知足連續內存是很是難的。可是對於對齊來講，是有一些方法來幫助應用實現的。

下面以代碼來檢驗上述知識，kernel中多了一個k索引，是用來配置偏移地址的，經過他就能夠配置對齊狀況，只有在load兩個數組A和B時纔會使用k。對C的寫操做則繼續使用原來的代碼，從而保證寫操做 保持很好的對齊。

__global__ void readOffset(float *A, float *B, float *C, const int n,int offset) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < n) C[i] = A[k] + B[k];
}

下面是main代碼，offset默認是零：

int main(int argc, char **argv) {
// set up device
int dev = 0;
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
printf("%s starting reduction at ", argv[0]);
printf("device %d: %s ", dev, deviceProp.name);
cudaSetDevice(dev);
// set up array size
int nElem = 1<<20; // total number of elements to reduce
printf(" with array size %d\n", nElem);
size_t nBytes = nElem * sizeof(float);
// set up offset for summary
int blocksize = 512;
int offset = 0;
if (argc>1) offset = atoi(argv[1]);
if (argc>2) blocksize = atoi(argv[2]);
// execution configuration
dim3 block (blocksize,1);
dim3 grid ((nElem+block.x-1)/block.x,1);
// allocate host memory
float *h_A = (float *)malloc(nBytes);
float *h_B = (float *)malloc(nBytes);
float *hostRef = (float *)malloc(nBytes);
float *gpuRef = (float *)malloc(nBytes);
// initialize host array
initialData(h_A, nElem);
memcpy(h_B,h_A,nBytes);
// summary at host side
sumArraysOnHost(h_A, h_B, hostRef,nElem,offset);
// allocate device memory
float *d_A,*d_B,*d_C;
cudaMalloc((float**)&d_A, nBytes);
cudaMalloc((float**)&d_B, nBytes);
cudaMalloc((float**)&d_C, nBytes);
// copy data from host to device
cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_A, nBytes, cudaMemcpyHostToDevice);
// kernel 1:
double iStart = seconds();
warmup <<< grid, block >>> (d_A, d_B, d_C, nElem, offset);
cudaDeviceSynchronize();
double iElaps = seconds() - iStart;
printf("warmup <<< %4d, %4d >>> offset %4d elapsed %f sec\n",
grid.x, block.x,
offset, iElaps);
iStart = seconds();
readOffset <<< grid, block >>> (d_A, d_B, d_C, nElem, offset);
cudaDeviceSynchronize();
iElaps = seconds() - iStart;
printf("readOffset <<< %4d, %4d >>> offset %4d elapsed %f sec\n",
grid.x, block.x,
offset, iElaps);
// copy kernel result back to host side and check device results
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
checkResult(hostRef, gpuRef, nElem-offset);
// copy kernel result back to host side and check device results
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
checkResult(hostRef, gpuRef, nElem-offset);
// copy kernel result back to host side and check device results
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
checkResult(hostRef, gpuRef, nElem-offset);
// free host and device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
// reset device
cudaDeviceReset();
return EXIT_SUCCESS;
}
View Code

編譯運行：

$ nvcc -O3 -arch=sm_20 readSegment.cu -o readSegment
$ ./readSegment 0
readOffset <<< 32768, 512 >>> offset 0 elapsed 0.001820 sec
$ ./readSegment 11
readOffset <<< 32768, 512 >>> offset 11 elapsed 0.001949 sec
$ ./readSegment 128
readOffset <<< 32768, 512 >>> offset 128 elapsed 0.001821 sec

當offset=11時，會致使從A和B load數據時不對齊。其運行時間消耗也是最大的，咱們可使用nvcc的gld_efficiency來檢驗一下：

$ nvprof --devices 0 --metrics gld_efficiency ./readSegment 0
$ nvprof --devices 0 --metrics gld_efficiency ./readSegment 11
$ nvprof --devices 0 --metrics gld_efficiency ./readSegment 128

輸出：

Offset 0: gld_efficiency 100.00%
Offset 11: gld_efficiency 49.81%
Offset 128: gld_efficiency 100.00%

能夠看到offset=11時，效率減半，能夠預見其吞吐必然很高，也可使用gld_transactions來檢驗：

$ nvprof --devices 0 --metrics gld_transactions ./readSegment $OFFSET

輸出爲：

Offset 0: gld_transactions 65184
Offset 11: gld_transactions 131039
Offset 128: gld_transactions 65744

而後咱們使用-Xptxas -dlcm=cg來禁用L1，看一下直接使用L2的表現：

$ ./readSegment 0
readOffset <<< 32768, 512 >>> offset 0 elapsed 0.001825 sec
$ ./readSegment 11
readOffset <<< 32768, 512 >>> offset 11 elapsed 0.002309 sec
$ ./readSegment 128
readOffset <<< 32768, 512 >>> offset 128 elapsed 0.001823 sec

從該結果看出，未對齊的狀況更糟糕了，而後看下gld_efficiency：

Offset 0: gld_efficiency 100.00%
Offset 11: gld_efficiency 80.00%
Offset 128: gld_efficiency 100.00%

由於L1被禁用後，每次load操做都是以32-byte爲單位而不是128，因此無用數據會減小很是多。

這裏未對齊反而狀況變糟是一種特例，高Occupancy狀況下，uncached會幫助提高bus有效使用率，而對於未對齊的狀況，無用數據的傳輸將明顯減小。

Read-Only Cache

最開始，read-only cache是用來爲texture Memory load服務的，對於CC3.5以上，該cache能夠替換L1（Maxwell以後，L1的功能徹底就被這個cache取代了）。Read-only cache的單位是32 bytes，通常來說是比L1要好用得多。

有兩種方式來使用read-only cache：

  1. Using the function __ldg
  2. Using a declaration qualifier on the pointer being dereferenced

例如：

__global__ void copyKernel(int *out, int *in) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = in[idx];
}

改寫後：

__global__ void copyKernel(int *out, int *in) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = __ldg(&in[idx]);
}

或者使用 const __restrict__來修飾指針。該修飾符幫助nvcc編譯器識別non-aliased指針，nvcc會自動使用該non-alias 指針從read-cache讀出數據。

__global__ void copyKernel(int * __restrict__ out,const int * __restrict__ in) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = in[idx];
}

Global Memory Writes

寫操做相對要簡單的多，L1壓根就不使用了。數據只會cache在L2中，因此寫操做也是以32bytes爲單位的。Memory transaction一次能夠是一個、兩個或四個segment。例如，若是兩個地址落在了同一個128-byte的區域內，可是在不一樣的兩個64-byte對齊的區域，一個四個segment的transaction就會被執行（也就是說，一個單獨的4-segment的傳輸要比兩次1-segment的傳輸性能好）。

下圖是一個理想的狀況，連續且對齊，只須要一次4 segment的傳輸：

 

下圖是離散的狀況，會由三次1-segment傳輸完成。

 

下圖是對齊且地址在一個連續的64-byte範圍內的狀況，由一次2-segment傳輸完成：

 

Example of Misaligned Writes

再次修改代碼，load變回使用i，而對C的寫則使用k：

__global__ void writeOffset(float *A, float *B, float *C,const int n, int offset) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int k = i + offset;
    if (k < n) C[k] = A[i] + B[i];
}

修改host的計算函數;

void sumArraysOnHost(float *A, float *B, float *C, const int n,int offset) {
    for (int idx = offset, k = 0; idx < n; idx++, k++) {
        C[idx] = A[k] + B[k];
    }
}

編譯運行：

$ nvcc -O3 -arch=sm_20 writeSegment.cu -o writeSegment
$ ./writeSegment 0
writeOffset <<< 2048, 512 >>> offset 0 elapsed 0.000134 sec
$ ./writeSegment 11
writeOffset <<< 2048, 512 >>> offset 11 elapsed 0.000184 sec
$ ./writeSegment 128
writeOffset <<< 2048, 512 >>> offset 128 elapsed 0.000134 sec

顯而易見，Misaligned表現最差，而後查看gld_efficiency：

$ nvprof --devices 0 --metrics gld_efficiency --metrics gst_efficiency ./writeSegment $OFFSET
writeOffset Offset 0: gld_efficiency 100.00%
writeOffset Offset 0: gst_efficiency 100.00%
writeOffset Offset 11: gld_efficiency 100.00%
writeOffset Offset 11: gst_efficiency 80.00%
writeOffset Offset 128: gld_efficiency 100.00%
writeOffset Offset 128: gst_efficiency 100.00%

除了offset=11的store外，全部load和store都是百分百。當offset=11時，128-bytes的寫請求會被一個4-segment和一個1-segment的傳輸服務，所以，咱們雖然須要寫128bytes可是卻有160bytes數據被load，從而致使百分之八十的效率。

Array of Structure versus Structure of Arrays

做爲C程序員，咱們應該熟悉兩種組織數據的方式：array of structures（AoS）和structure of arrays（SoA）。兩者的使用是一個有趣的話題，主要是數據排列組織。

觀察下面代碼，首先考慮該數據結構集合在使用AoS組織時，是怎樣存儲的：

struct innerStruct {
    float x;
    float y;
};
struct innerStruct myAoS[N]; //每一對x和y的存儲，空間上是連續的

而後是SoA：

struct innerArray {
    float x[N];
    float y[N];
};
struct innerArray moa; //x和y是分別存儲的，全部x和y是分別存儲在兩段不一樣的連續地址裏。

下圖顯示了AoS和SoA在內存中的存儲格式，當對x進行操做時，會致使通常的帶寬浪費，由於在操做x時，y也會隱式的被load，而SoA的表現就要好得多，由於全部x都是相鄰的。

 

許多並行編程規範裏，特別是SIMD-style風格的規範，都更傾向於使用SoA，在CUDA C裏，SoA也是很是建議使用的，由於數據已經預先排序連續了。

Example：Simple Math with the AoS Data Layout

__global__ void testInnerStruct(innerStruct *data,innerStruct *result, const int n) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        innerStruct tmp = data[i];
        tmp.x += 10.f;
        tmp.y += 20.f;
        result[i] = tmp;
    }
}

輸入長度是1M，#define LEN 1<<20。

初始化數據：

void initialInnerStruct(innerStruct *ip, int size) {
    for (int i = 0; i < size; i++) {
        ip[i].x = (float)(rand() & 0xFF) / 100.0f;
        ip[i].y = (float)(rand() & 0xFF) / 100.0f;
    }
    return;
}

Main代碼：

int main(int argc, char **argv) {
// set up device
int dev = 0;
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
printf("%s test struct of array at ", argv[0]);
printf("device %d: %s \n", dev, deviceProp.name);
cudaSetDevice(dev);
// allocate host memory
int nElem = LEN;
size_t nBytes = nElem * sizeof(innerStruct);
innerStruct *h_A = (innerStruct *)malloc(nBytes);
innerStruct *hostRef = (innerStruct *)malloc(nBytes);
innerStruct *gpuRef = (innerStruct *)malloc(nBytes);
// initialize host array
initialInnerStruct(h_A, nElem);
testInnerStructHost(h_A, hostRef,nElem);
// allocate device memory
innerStruct *d_A,*d_C;
cudaMalloc((innerStruct**)&d_A, nBytes);
cudaMalloc((innerStruct**)&d_C, nBytes);
// copy data from host to device
cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
// set up offset for summary
int blocksize = 128;
if (argc>1) blocksize = atoi(argv[1]);
// execution configuration
dim3 block (blocksize,1);
dim3 grid ((nElem+block.x-1)/block.x,1);
// kernel 1: warmup
double iStart = seconds();
warmup <<< grid, block >>> (d_A, d_C, nElem);
cudaDeviceSynchronize();
double iElaps = seconds() - iStart;
printf("warmup <<< %3d, %3d >>> elapsed %f sec\n",grid.x,
block.x,iElaps);
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
checkInnerStruct(hostRef, gpuRef, nElem);
// kernel 2: testInnerStruct
iStart = seconds();
testInnerStruct <<< grid, block >>> (d_A, d_C, nElem);
cudaDeviceSynchronize();
iElaps = seconds() - iStart;
printf("innerstruct <<< %3d, %3d >>> elapsed %f sec\n",grid.x,
block.x,iElaps);
cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
checkInnerStruct(hostRef, gpuRef, nElem);
// free memories both host and device
cudaFree(d_A);
cudaFree(d_C);
free(h_A);
free(hostRef);
free(gpuRef);
// reset device
cudaDeviceReset();
return EXIT_SUCCESS;
}
View Code

編譯運行(Fermi M2070)：

$ nvcc -O3 -arch=sm_20 simpleMathAoS.cu -o simpleMathAoS
$ ./simpleMathAoS
innerStruct <<< 8192, 128 >>> elapsed 0.000286 sec

查看load和store性能：

$ nvprof --devices 0 --metrics gld_efficiency,gst_efficiency ./simpleMathAoS
gld_efficiency 50.00%
gst_efficiency 50.00%

正如預期那樣，都只達到了通常，由於額外那部分消耗都用來load/store 另外一個元素了，而這部分不是咱們須要的。

Example：Simple Math with the SoA Data Layout

__global__ void testInnerArray(InnerArray *data,InnerArray *result, const int n) {
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i<n) {
        float tmpx = data->x[i];
        float tmpy = data->y[i];
        tmpx += 10.f;
        tmpy += 20.f;
        result->x[i] = tmpx;
        result->y[i] = tmpy;
    }
}

分配global Memory：

int nElem = LEN;
size_t nBytes = sizeof(InnerArray);
InnerArray *d_A,*d_C;
cudaMalloc((InnerArray **)&d_A, nBytes);
cudaMalloc((InnerArray **)&d_C, nBytes);

編譯運行：

$ nvcc -O3 -arch=sm_20 simpleMathSoA.cu -o simpleSoA
$ ./simpleSoA
innerArray <<< 8192, 128 >>> elapsed 0.000200 sec

查看load/store性能：

$ nvprof --devices 0 --metrics gld_efficiency,gst_efficiency ./simpleMathSoA
gld_efficiency 100.00%
gst_efficiency 100.00%

Performance Tuning

調節device Memory帶寬利用性能時，主要是力求達到下面兩個目標：

  1. Aligned and Coalesced Memory accesses that reduce wasted bandwidth
  2. Sufficient concurrent Memory operations to hide Memory latency

Unrolling Techniques

展開循環能夠增長更多的獨立的Memory操做，咱們在以前博文有詳細介紹如何展開loop，考慮以前的redSegment的例子，咱們修改下readOffset來使每一個thread執行四個獨立Memory操做，就像下面那樣：

__global__ void readOffsetUnroll4(float *A, float *B, float *C,const int n, int offset) {
    unsigned int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    unsigned int k = i + offset;
    if (k + 3 * blockDim.x < n) {
        C[i] = A[k]
        C[i + blockDim.x] = A[k + blockDim.x] + B[k + blockDim.x];
        C[i + 2 * blockDim.x] = A[k + 2 * blockDim.x] + B[k + 2 * blockDim.x];
        C[i + 3 * blockDim.x] = A[k + 3 * blockDim.x] + B[k + 3 * blockDim.x];
    }
}

編譯運行（可能須要使用-Xptxas -dlcm=ca來啓用L1）：

$ ./readSegmentUnroll 0
warmup <<< 32768, 512 >>> offset 0 elapsed 0.001990 sec
unroll4 <<< 8192, 512 >>> offset 0 elapsed 0.000599 sec
$ ./readSegmentUnroll 11
warmup <<< 32768, 512 >>> offset 11 elapsed 0.002114 sec
unroll4 <<< 8192, 512 >>> offset 11 elapsed 0.000615 sec
$ ./readSegmentUnroll 128
warmup <<< 32768, 512 >>> offset 128 elapsed 0.001989 sec
unroll4 <<< 8192, 512 >>> offset 128 elapsed 0.000598 sec

咱們看到，unrolling技術會對性能有巨大影響，比地址對齊影響還大。對於這類I/O-bound的kernel，提升內存獲取的並行性對性能提高的影響，有更高的優先級。不過，咱們應該看到，對齊的test比未對齊的test表現依然要好。

Unrolling並不能影響內存操做的總數目（只是影響並行的操做數目），咱們能夠查看下相關屬性：

$ nvprof --devices 0 --metrics gld_efficiency,gst_efficiency ./readSegmentUnroll 11
readOffset gld_efficiency 49.69%
readOffset gst_efficiency 100.00%
readOffsetUnroll4 gld_efficiency 50.79%
readOffsetUnroll4 gst_efficiency 100.00%
$ nvprof --devices 0 --metrics gld_transactions,gst_transactions
./readSegmentUnroll 11
readOffset gld_transactions 132384
readOffset gst_transactions 32928
readOffsetUnroll4 gld_transactions 33152
readOffsetUnroll4 gst_transactions 8064

Exposing More Parallelism

這方面就是調整grid和block的配置，下面是加上unrolling後的結果：

$ ./readSegmentUnroll 0 1024 22
unroll4 <<< 1024, 1024 >>> offset 0 elapsed 0.000169 sec
$ ./readSegmentUnroll 0 512 22
unroll4 <<< 2048, 512 >>> offset 0 elapsed 0.000159 sec
$ ./readSegmentUnroll 0 256 22
unroll4 <<< 4096, 256 >>> offset 0 elapsed 0.000157 sec
$ ./readSegmentUnroll 0 128 22
unroll4 <<< 8192, 128 >>> offset 0 elapsed 0.000158 sec

表現最好的是block配置256 thread的kernel，雖然128thread會增長並行性，可是依然比256少那麼一點點性能，這個主要是CC版本對應的資源限制決定的，以本代碼爲例，Fermi每一個SM最多有8個block，每一個SM可以並行的的warp是48個，當使用128個thread（per block）時，每一個block中有4個warp，由於每一個SM最多8個block可以同時運行，所以該kernel每一個SM最多隻能有32個warp，還有16個warp的計算性能沒用上，因此性能差了就，可使用Occupancy來驗證下。

 

 

參考書：《professional cuda c programming》

相關文章
相關標籤/搜索
每日一句
    每一个你不满意的现在，都有一个你没有努力的曾经。
本站公眾號
   歡迎關注本站公眾號,獲取更多信息
聯繫我們 最近搜索 最新文章 沪ICP备13005482号-10 MyBatis教程 SQL 教程 MySQL教程 Java 教程 Thymeleaf 教程 Hibernate教程 Spring教程 Redis教程