tex1Dfetch意外返回0

Question

我不相信這是同一個問題，問題報告如下：

Bound CUDA texture reads zero

CUDA 1D texture fetch always return 0

在我的CUDA應用程序，我注意到，tex1Dfetch沒有返回預期值，超過緩衝區中的某個索引。在應用程序中的一個初步觀察是，可以正確讀取索引0處的值，但是在12705625處，讀取的值爲0.我製作了一個小型測試程序來調查此問題，如下所示。結果對我來說有點莫名其妙。我試圖探索哪些索引值不再被正確讀取。但是，隨着arraySize值的改變，「firstBadIndex」也會改變。即使arraySize = 2，第二個值也被錯誤地讀取！隨着arraySize變大，firstBadIndex變大。綁定到float，float2或float4數組時會發生這種情況。如果數據是從設備緩衝區讀取的（切換FetchTextureData中的註釋行），那麼一切都很好。這是使用特斯拉c2075上的CUDA 6.5。 感謝您提供的任何見解或建議。

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <stdio.h> 

#define FLOATTYPE float4 
texture<FLOATTYPE,cudaTextureType1D,cudaReadModeElementType> texture1D; 

const unsigned int arraySize = 1000; 
FLOATTYPE* host; 
FLOATTYPE* device; 
FLOATTYPE* dTemp; 
FLOATTYPE hTemp[1]; 

__global__ void FetchTextureData(FLOATTYPE* data,FLOATTYPE* arr,int idx) 
{ 
    data[0] = tex1Dfetch(texture1D, idx); 
    //data[0] = arr[idx]; 
} 

bool GetTextureValues(int idx){ 

    FetchTextureData<<<1,1>>>(dTemp,device,idx); 

    // copy to the host 
    cudaError_t err = cudaMemcpy(hTemp,dTemp,sizeof(FLOATTYPE),cudaMemcpyDeviceToHost); 
    if (err != cudaSuccess) { 
     throw "cudaMemcpy failed!"; 
    } 

    if (cudaDeviceSynchronize() != cudaSuccess) { 
     throw "cudaDeviceSynchronize failed!"; 
    } 

    return hTemp[0].x == 1.0f; 
} 

int main() 
{ 

    try{ 

     host = new FLOATTYPE[arraySize]; 
     cudaError_t err = cudaMalloc((void**)&device,sizeof(FLOATTYPE) * arraySize); 
     cudaError_t err1 = cudaMalloc((void**)&dTemp,sizeof(FLOATTYPE)); 
     if (err != cudaSuccess || err1 != cudaSuccess) { 
      throw "cudaMalloc failed!"; 
     } 

     // make some host data 
     for(unsigned int i=0; i<arraySize; i++){ 
      FLOATTYPE data = {1.0f, 0.0f, 0.0f, 0.0f}; 
      host[i] = data; 
     } 

     // and copy it to the device 
     err = cudaMemcpy(device,host,sizeof(FLOATTYPE) * arraySize,cudaMemcpyHostToDevice); 
     if (err != cudaSuccess){ 
      throw "cudaMemcpy failed!"; 
     } 

     // set up the textures 
     cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<FLOATTYPE>(); 
     texture1D.addressMode[0] = cudaAddressModeClamp; 
     texture1D.filterMode = cudaFilterModePoint; 
     texture1D.normalized = false; 
     cudaBindTexture(NULL, texture1D, device, channelDesc, arraySize); 

     // do a texture fetch and find where the fetches stop working 
     int lastGoodValue = -1, firstBadValue = -1; 
     float4 badValue = {-1.0f,0.0f,0.0f,0.0f}; 

     for(unsigned int i=0; i<arraySize; i++){ 

      if(i % 100000 == 0) printf("%d\n",i); 
      bool isGood = GetTextureValues(i); 

      if(firstBadValue == -1 && !isGood) 
       firstBadValue = i; 

      if(isGood) 
       lastGoodValue = i; 
      else 
       badValue = hTemp[0]; 
     } 

     printf("lastGoodValue %d, firstBadValue %d\n",lastGoodValue,firstBadValue); 
     printf("Bad value is (%.2f)\n",badValue.x); 

    }catch(const char* err){ 
     printf("\nCaught an error : %s\n",err); 
    } 
    return 0; 
} 

Answer 1

問題在於紋理的設置。此：

cudaBindTexture(NULL, texture1D, device, channelDesc, arraySize); 

應該是：

cudaBindTexture(NULL, texture1D, device, channelDesc, 
            arraySize * sizeof(FLOATTYPE)); 

作爲每documentation，尺寸參數是存儲器區域的以字節爲單位的大小，而不是元件的數量。我本來預料到，在鉗位尋址模式下，代碼仍能按預期工作。使用邊界模式時，您應該得到一個零值，這看起來會觸發不良值檢測。我實際上並沒有運行你的代碼，所以也許我在某處丟失了一些微妙的東西。對於這樣一個簡單的repro例子，你的代碼結構相當複雜，難以遵循（至少在我正在閱讀的手機屏幕上）。

編輯補充一點，我開始寫這和完成的時間之間，@njuffa指出，在意見相同的錯誤