本文主要介紹基於CUDA的向量逆向，並結合三種CUDA異構記憶體進行相應優化。

1. 向量逆向

基於 Global memory 的核心演算法

Global memory空間大，但是速度是所有Device記憶體中最慢的

__global__ void array_reverse(int *array_a_dev, int *array_a_rev_dev, int len)
{
    int tid = threadIdx.x;
    array_a_rev_dev[len - tid - 1] = array_a_dev[tid];
}

基於 Shared memory 的核心演算法

當存在多個執行緒共同訪問同一個資料的時候，可以將該資料放到shared memory中，以提高資料讀取效率

__global__ void array_reverse_shared(int *array_a_dev, int *array_a_rev_dev, int len)
{
    int tid = threadIdx.x;
    __shared__ int array_shared[9];
    array_shared[tid] = array_a_dev[tid];
    __syncthreads();
    array_a_rev_dev[len - tid -
 
 1] = array_shared[tid];
}

基於Dynamic shared memory 核心演算法

Dynamic shared memory 通過採用extern的宣告而不定義的方式，提高了程式碼的靈活性，比方說，dynamic shared memory可以在同一個程式的不同kernel中被定義，也可以動態地根據需求進行shared memory記憶體空間的分配。如下，

__global__ void array_reverse_dynamic_shared(int *array_a_dev, int *array_a_rev_dev, int len)
{
    int
 
 tid = threadIdx.x;
    extern __shared__ int array_shared[];
    // __shared__ int array_shared[9];
    array_shared[tid] = array_a_dev[tid];
    __syncthreads();
    array_a_rev_dev[len - tid - 1] = array_shared[tid];
}

Dynamic shared memory的kernel的呼叫有些不同，

array_reverse_dynamic_shared<<<dimGrid, dimBlock, len*sizeof(int)>>>(array_a_dev, array_a_rev_dev, len);

注：動態變化 (或者說，非預編譯的變數) 的記憶體空間大小 (比如，這裡的len*sizeof(int)) 需要在kernel函式呼叫時定義。

2. 編譯除錯

基於 Global memory 的程式碼

原始碼：array_rev.cu

編譯

nvcc array_rev.cu -o main

執行

./main

基於 Shared memory 的程式碼

原始碼：array_rev_sm.cu

編譯

nvcc array_rev_sm.cu -o main

執行

./main

基於 Dynamic shared memory 的程式碼

原始碼：array_rev_dynamic_sm.cu

編譯

nvcc array_rev_dynamic_sm.cu -o main

執行

./main