CUDA：__restrict__標籤使用

我不太瞭解CUDA的__restrict__標籤的概念。CUDA：__restrict__標籤使用

我讀過使用__restrict__可以避免指針別名，尤其是，如果指向的變量是隻讀的，則變量的讀取因爲被緩存而被優化。

這是一個簡化版本的代碼：

__constant__ float M[M_DIM1][M_DIM2]; 

__host__ void function(float N[][DIM2], float h_M[][M_DIM2], float P[][DIM2]); 

__global__ void kernel_function(const float* __restrict__ N, float *P); 

__host__ void function(float N[][DIM2], float h_M[][M_DIM2], float P[][DIM2]) { 

    int IOSize = DIM1 * DIM2 * sizeof(float); 
    int ConstSize = M_DIM1* M_DIM2* sizeof(float); 
    float* dN, *dP; 
    cudaMalloc((void**)&dN, IOSize); 
    cudaMemcpy(dN, N, IOSize, cudaMemcpyHostToDevice); 

    cudaMemcpyToSymbol(M, h_M, ConstSize); 

    cudaMalloc((void**)&dP, IOSize); 

    dim3 dimBlock(DIM1, DIM2); 
    dim3 dimGrid(1, 1); 

    kernel_function << <dimGrid, dimBlock >> >(dN, dP); 

    cudaMemcpy(P, dP, IOSize, cudaMemcpyDeviceToHost); 

    cudaFree(dN); 
    cudaFree(dP); 

}

我使用N個的__restrict__標籤，這是隻讀的，以正確的方式？此外，我讀過關於M的關鍵字__constant__，意思是說只讀和常量，那麼它們之間的區別是什麼，分配的類型是什麼？

來源

2017-04-05 Pleasant94

CUDA不是C.請勿垃圾標籤。 – Olaf

__restrict__如nvcc所使用的記錄爲here。（請注意，包括gnu編譯器在內的各種C++編譯器也支持這個確切的關鍵字，並且同樣使用它）。

它具有與C99 restrict關鍵字基本相同的語義，即an official part of that language standard。

簡而言之，__restrict__是一個合同，你作爲一個程序員做的編譯器，它說，大致，「我將只使用該指針指到底層數據」。從編譯器的角度來看，關鍵事件之一是指針別名，這可能會阻止編譯器進行各種優化。

如果您想對restrict或__restrict__的確切定義進行更長時間的正式論述，請參閱我已提供的某個鏈接或進行一些研究。

所以，__restrict__一般可到支持它，爲了優化的編譯器。

對於計算能力3.5或更高版本的設備，這些設備有一個單獨的高速緩存稱爲read only cache，其獨立普通的L1類型的高速緩存。

如果同時使用__restrict__和const裝飾傳遞給內核的全球指針，那麼這也是一個強烈的暗示，編譯器，用於cc3.5及更高版本的設備生成代碼時，爲了使這些全球內存負載流通過只讀緩存。這可以提供應用程序性能優勢，通常只需很少的其他代碼重構。這不保證只讀緩存的使用情況，即使您不使用這些裝飾器，編譯器也會經常試圖積極使用只讀緩存，只要它能滿足必要的條件即可。

__constant__是指不同hardware resource on the GPU。有許多不同之處：

__constant__是適用於所有的GPU，僅在cc3.5只讀緩存和更高的使用__constant__標籤（包含上線指定分配
內存分配的內存）限制爲最大64KB。只讀緩存沒有這種限制。我們不把__restrict__放在分配內存的行上;它用於裝飾指針。
高速緩存在只讀高速緩存中的數據具有典型的全局內存訪問注意事項 - 通常我們希望通過只讀高速緩存來實現全局內存讀取的最佳合併的相鄰和連續訪問。該機制，OTOH，預計所謂的統一的獲得最快的性能。統一訪問實質上意味着每個線程都在請求來自的相同位置/地址/索引。

兩個__constant__內存，並標有const裝飾傳遞給內核代碼指針全局存儲器，只讀從內核代碼的角度。

我沒有看到任何明顯的問題，你已經顯示的代碼，無論是使用__restrict__或其他。我唯一的意見就是爲了獲得最大利益，您可能希望在您的內核聲明/原型中使用__restrict__來裝飾N和P指針，以獲得最大利益，如果這是您的意圖。（顯然，你不會裝飾P與const。）

來源

2017-04-05 20:11:15

CUDA：restrict標籤使用

回答

CUDA：__restrict__標籤使用

回答

相關問題

CUDA：restrict標籤使用