CUDA恆定存儲器參考

我在恆定存儲器（它是一個全局變量）的陣列和獲得函數調用cudaGetSymbolAddress對它的引用。當我使用此引用來獲取常量數據而不是使用全局變量時，我的內核運行緩慢。這是什麼原因？CUDA恆定存儲器參考

__constant__ int g[2] = {1,2}; 
// __device__ int g[2] = {1,2}; 

// kernel: use by reference 
__global__ void add_1(int *a, int *b, int *c, int *f) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    c[tid] = f[0] * a[tid] + f[1] * b[tid]; 
} 

// kernel: use global variable 
__global__ void add_2(int *a, int *b, int *c, int *f) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    c[tid] = g[0] * a[tid] + f[1] * b[tid]; 
} 

int main() 
{ 
    ...... 
    // a,b,c are large arrays in device memory of size 40960. 

    int *f; 
    cudaGetSymbolAddress((void **)&f, (char *)&g); 

    add_1 <<< 160, 256 >>> (a, b, c, f); 

    ...... 
}

這是示例代碼和warp中的所有線程在同一時間加載相同的位置。註釋的代碼是通過直接訪問常量存儲器

解釋爲什麼恆定存儲器高速緩存未使用（由talonmies）

的原因是缺乏恆定緩存。僅當編譯器在標記爲處於恆定狀態空間的顯式變量上發出特定的PTX指令（ld.const）時，纔會發生緩存訪問。編譯器知道這樣做的方式是聲明一個變量__constant__ - 它是一個靜態編譯時間屬性，它影響代碼生成。運行時不會發生同樣的過程。

如果您通過在全局內存的指針，編譯器不能確定的是，在恆定的狀態空間的指針，它不會產生正確的PTX通過不斷的高速緩存來訪問內存。因此訪問速度會變慢。

懸而未決的問題

爲什麼即使當陣列g被聲明爲__device__變量，代碼時參照它使用慢。通過觀看PTX的代碼，對全局存儲器加載到寄存器：用於

2的ld.global.s32指令，它加載4個字節的寄存器。（在代碼中使用參考）
1的ld.global.v2.s32指令被使用，其加載8個字節到2個寄存器，（在代碼中使用全局變量）

有什麼不同和任何文件參考，將不勝感激？

來源

2012-10-17 Chaitanya Turubati

我想你並沒有利用這些常量功能，你如何訪問全局和常量內存中的數據？給你一個更好的答案 – pQB

pQB是對的，我剛剛看到了talonmies的答案，並且給出的信息是不可能確定哪一個適用於你的問題的。你的GPU？ – tera

@tera計算能力是1.3 –

與全局內存訪問常量內存將得到序列化（分成多次交易），如果他們不統一（一（半用於計算能力1.x的所有線程）經訪問相同的地址。

因此，如果訪問可能是統一的，那麼只能使用常量內存。

來源

2012-10-17 12:10:11 tera

我不認爲這是一個非一致的訪問，warp中的所有線程都請求加載指令的相同地址位置。 –

確實如此。我原本以爲[tid]索引的訪問將是不斷記憶的訪問。你能否提供來自'cuobjdump -sass -fun add'的反彙編內核代碼，以便我們可以更好地瞭解發生了什麼？ – tera

你有沒有嘗試過，如果沒有你的程序的其餘部分，如上所示的隔離代碼仍然顯示該問題？這個問題可能與我們在這裏看不到的其他東西有關...... – tera

CUDA恆定存儲器參考

回答

相關問題