紋理內存與讀取和寫入

Question

我正在開發一個CUDA應用程序，內核必須多次訪問全局內存。這個內存可以被所有的CTA隨機訪問（沒有局部性，所以不能使用共享內存）。我需要優化它。我聽說紋理內存可以緩解這個問題，但是內核可以讀寫紋理內存嗎？ 1D紋理內存？ 2D紋理內存？那麼CUDA數組呢？

Answer 1

CUDA紋理是隻讀的。紋理讀取被緩存。所以性能增益是概率性的。

CUDA Toolkit 3.1向前也具有稱爲Surfaces的可寫紋理，但它們僅適用於Compute Capability> = 2.0的設備。表面就像紋理，但優點是它們也可以由內核編寫。

表面只能綁定到cudaArray用標記cudaArraySurfaceLoadStore創建。

Answer 2

我會建議你宣佈你的記憶爲傾斜的線性記憶並將其與紋理結合。我還沒有嘗試新的無紋理紋理。任何人都試過嗎？

所提到的紋理mem是隻讀通過緩存。將它視爲只讀內存。 因此，請注意，在內核本身中，不要寫入綁定到紋理的內存，因爲它可能不會更新到紋理緩存。

Answer 3

我碰到過這個問題，並用一點搜索我發現this問題和this答案對它有用。 基本上紋理內存是全局內存。紋理內存是指可以與全局內存關聯的特殊緩存機制，內容爲。所以內核可以操縱全局內存綁定到紋理。但正如它在provided link中所示，沒有指令，如tex1D(ref, x) = 12.0。

Answer 4

這是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; 
} 

Answer 5

這是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; 
} 

請注意以下幾點：

1. 我紋理不具約束力的cudaArray，因爲cudaArray s不能從內核中進行修改;
2. 我不結合紋理到cudaMalloc版陣列，由於紋理結合到cudaMalloc版陣列可以僅通過tex1Dfetch被擷取並tex1Dfetch不所述cudaAddressModeWrap尋址模式保證其邊界以外的信號的週期延拓;
3. 我將紋理綁定到cudaMallocPitch ed數組，因爲這使得可以通過tex1D獲取紋理，這允許cudaAddressModeWrap尋址模式;
4. 我使用標準座標，因爲只有那些啓用cudaAddressModeWrap尋址模式。

我需要點#2，#3和#4，因爲我從我在寫一個代碼中提取這個例子。