我正在開發一個CUDA應用程序,內核必須多次訪問全局內存。這個內存可以被所有的CTA隨機訪問(沒有局部性,所以不能使用共享內存)。我需要優化它。我聽說紋理內存可以緩解這個問題,但是內核可以讀寫紋理內存嗎? 1D紋理內存? 2D紋理內存?那麼CUDA數組呢?紋理內存與讀取和寫入
回答
CUDA紋理是隻讀的。紋理讀取被緩存。所以性能增益是概率性的。
CUDA Toolkit 3.1向前也具有稱爲Surfaces的可寫紋理,但它們僅適用於Compute Capability> = 2.0的設備。表面就像紋理,但優點是它們也可以由內核編寫。
表面只能綁定到cudaArray
用標記cudaArraySurfaceLoadStore
創建。
我可以在表面存儲器上執行原子操作嗎? –
我會建議你宣佈你的記憶爲傾斜的線性記憶並將其與紋理結合。我還沒有嘗試新的無紋理紋理。任何人都試過嗎?
所提到的紋理mem是隻讀通過緩存。將它視爲只讀內存。 因此,請注意,在內核本身中,不要寫入綁定到紋理的內存,因爲它可能不會更新到紋理緩存。
我碰到過這個問題,並用一點搜索我發現this問題和this答案對它有用。 基本上紋理內存是全局內存。紋理內存是指可以與全局內存關聯的特殊緩存機制,內容爲。所以內核可以操縱全局內存綁定到紋理。但正如它在provided link中所示,沒有指令,如tex1D(ref, x) = 12.0
。
這是sgarizvi的答案的後續行動。
如今,計算能力爲>=2.0
的卡比2012
要普遍得多,也就是在問這個問題的時候。
下面,對如何使用CUDA表面存儲器寫入到紋理最小例子。
#include <stdio.h>
#include "TimingGPU.cuh"
#include "Utilities.cuh"
surface<void, cudaSurfaceType1D> surfD;
/*******************/
/* KERNEL FUNCTION */
/*******************/
__global__ void SurfaceMemoryWrite(const int N) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
surf1Dwrite((float)tid, surfD, tid * sizeof(float), cudaBoundaryModeTrap);
}
/********/
/* MAIN */
/********/
int main() {
const int N = 10;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
//Alternatively
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaArray *d_arr; gpuErrchk(cudaMallocArray(&d_arr, &channelDesc, N, 1, cudaArraySurfaceLoadStore));
gpuErrchk(cudaBindSurfaceToArray(surfD, d_arr));
SurfaceMemoryWrite<<<1, N>>>(N);
float *h_arr = new float[N];
gpuErrchk(cudaMemcpyFromArray(h_arr, d_arr, 0, 0, N * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) printf("h_arr[%i] = %f\n", i, h_arr[i]);
return 0;
}
這是Farzad的回答的後續行動。
法扎德的觀點是強調了CUDA C編程指南:
的質地和表面內存緩存(請參閱設備內存訪問) 和同一內核中,高速緩存沒有保持連貫與 尊重全局內存寫入和表面內存寫入,因此任何通過在同一內核調用中通過全局寫入或表面寫入寫入到 的地址的任何 紋理讀取或表面讀取都會返回 未定義的數據。換句話說,一個線程可以安全地只讀如果此內存位置已被 由以前的內核調用或存儲器複製更新一些紋理 或表面內存的位置,但如果它已 以前由同一個線程或其他線程更新從內核調用 。
這意味着可以修改紋理綁定的全局內存位置,但是這不能發生在紋理提取操作的同一個內核中。在另一邊,「寫紋理」在上述意義上,可以跨自從紋理緩存在內核啓動清仁,見cuda kernel for add(a,b,c) using texture objects for a & b - works correctly for 'increment operation' add(a,b,a)?。
下面,我提供其中全局存儲器位置的紋理被綁定到被修改的例子。在這個例子中,我呼籲通過以下方式
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
...
square<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_vec, pitch, N);
...
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
在median_filter_periodic_boundary
內核,紋理拾取進行操作的CUDA內核,而在square
內核,紋理綁定的全局內存位置被修改。
下面是代碼:
#include <stdio.h>
#include "TimingGPU.cuh"
#include "Utilities.cuh"
texture<float, 1, cudaReadModeElementType> signal_texture;
#define BLOCKSIZE 32
/*************************************************/
/* KERNEL FUNCTION FOR MEDIAN FILTER CALCULATION */
/*************************************************/
__global__ void median_filter_periodic_boundary(float * __restrict__ d_out, const unsigned int N){
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) {
float signal_center = tex1D(signal_texture, (float)(tid + 0.5 - 0)/(float)N);
float signal_before = tex1D(signal_texture, (float)(tid + 0.5 - 1)/(float)N);
float signal_after = tex1D(signal_texture, (float)(tid + 0.5 + 1)/(float)N);
d_out[tid] = (signal_center + signal_before + signal_after)/3.f;
}
}
/*************************************************/
/* KERNEL FUNCTION FOR MEDIAN FILTER CALCULATION */
/*************************************************/
__global__ void square(float * __restrict__ d_vec, const size_t pitch, const unsigned int N){
unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N) d_vec[tid] = 2.f * tid;
}
/********/
/* MAIN */
/********/
int main() {
const int N = 10;
// --- Input/output host array declaration and initialization
float *h_vec = (float *)malloc(N * sizeof(float));
for (int i = 0; i < N; i++) h_vec[i] = (float)i;
// --- Input/output host and device array vectors
size_t pitch;
float *d_vec; gpuErrchk(cudaMallocPitch(&d_vec, &pitch, N * sizeof(float), 1));
printf("pitch = %i\n", pitch);
float *d_out; gpuErrchk(cudaMalloc(&d_out, N * sizeof(float)));
gpuErrchk(cudaMemcpy(d_vec, h_vec, N * sizeof(float), cudaMemcpyHostToDevice));
// --- CUDA texture memory binding and properties definition
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
//Alternatively
//cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
size_t texture_offset = 0;
gpuErrchk(cudaBindTexture2D(&texture_offset, signal_texture, d_vec, channelDesc, N, 1, pitch));
signal_texture.normalized = true;
signal_texture.addressMode[0] = cudaAddressModeWrap;
// --- Median filter kernel execution
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(h_vec, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
printf("\n\nFirst filtering\n");
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);
// --- Square kernel execution
square<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_vec, pitch, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaMemcpy(h_vec, d_vec, N * sizeof(float), cudaMemcpyDeviceToHost));
printf("\n\nSquaring\n");
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);
// --- Median filter kernel execution
median_filter_periodic_boundary<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_out, N);
gpuErrchk(cudaPeekAtLastError());
gpuErrchk(cudaDeviceSynchronize());
printf("\n\nSecond filtering\n");
gpuErrchk(cudaMemcpy(h_vec, d_out, N * sizeof(float), cudaMemcpyDeviceToHost));
for (int i=0; i<N; i++) printf("h_vec[%i] = %f\n", i, h_vec[i]);
printf("Test finished\n");
return 0;
}
請注意以下幾點:
- 我紋理不具約束力的
cudaArray
,因爲cudaArray
s不能從內核中進行修改; - 我不結合紋理到
cudaMalloc
版陣列,由於紋理結合到cudaMalloc
版陣列可以僅通過tex1Dfetch
被擷取並tex1Dfetch
不所述cudaAddressModeWrap
尋址模式保證其邊界以外的信號的週期延拓; - 我將紋理綁定到
cudaMallocPitch
ed數組,因爲這使得可以通過tex1D
獲取紋理,這允許cudaAddressModeWrap
尋址模式; - 我使用標準座標,因爲只有那些啓用
cudaAddressModeWrap
尋址模式。
我需要點#2
,#3
和#4
,因爲我從我在寫一個代碼中提取這個例子。
- 1. 讀取和寫入一個紋理(OpenGL)
- 2. WebGL紋理同時讀取/寫入
- 3. Nodejs:讀取和寫入內存地址
- 4. 讀取和寫入內存位置
- 5. 用於讀取和寫入DDS紋理的Java庫
- 6. OpenGl讀取和寫入相同的紋理
- 7. 在內存中讀取/寫入位
- 8. 一次寫入多次讀取內存
- 9. msn加腳本讀取/寫入內存
- 10. 內存塊C寫入/讀取檢測
- 11. sqlalchemy,從mysql讀取並寫入內存
- 12. 無法讀出紋理內存的值
- 13. 寫入紋理GLSL
- 14. 寫入GL_R8紋理
- 15. 寫入紋理陣列的一層並從另一層讀取
- 16. Linux內核寫入()和讀取()函數
- 17. 文件讀取,寫入和保存
- 18. OpenCL紋理內存
- 19. 使用OpenCL寫入Nvidia設備上的紋理內存
- 20. 避免依賴紋理與「大」內核讀取
- 21. SFML紋理內存管理
- 22. Android將文件讀取和寫入內部存儲
- 23. 在C內存中讀取和寫入BMP文件錯誤
- 24. i2c從機寫入和從內存位置讀取vhdl
- 25. 多個動畫和紋理與Cocos2d - 如何從內存中刪除紋理?
- 26. OpenGL:紋理大小和視頻內存
- 27. PIC16F84 - eeprom讀取和寫入
- 28. Plist讀取和寫入iPhone
- 29. C#app.config讀取和寫入
- 30. 讀取和寫入文件
閱讀CUDA編程指南的第3部分中有關Surface存儲器的文檔。 – talonmies
如果你的內存訪問在大塊內存中是真正的隨機存儲而沒有局部性,那麼沒有一種緩存會給你帶來顯着的改進。您或者需要找到一種方法來改善內存訪問模式,或者使用當前的低性能。 – tera