CUDA读取纹理内存出错

2013-10-18 11:51:45 535 阅读 AI 写同款 · GEO 优化AI 选题数据监测一键分发 ›

CUDA Kepler（一） Read-Only Cache特性解析及其性能对比测试 Kepler 计算设备 Kepler是NVIDIA发布的第三代CUDA计算设备的代号，这一系列产品主要有两大类：GK104和GK110，没有记错的话，GK104的运算能力是3.0，而GK110则具备了完整的Kepler架构组件，计算能力为3.5。两种计算能力的硬件在体系结构上面有一些差异，这些差异可能会在这一系列文章的后续版本中进行详细地讲解。今天我们讨论的是计算能力3.5的GK110架构新引入阅读详情

如题，这个问题很让人困惑，我使用CUDA数组和cudaMallocPitch分配的一维数组都不好用，都出现乱码
我的kernel函数为：





texture<float, cudaTextureType2D, cudaReadModeElementType> texSrc;

__inline__ __global__ void convolutionRows_kernel( float *d_Dst, int imageH, int imageW)

{ 

        int ix = __mul24(blockIdx.x, blockDim.x) + threadIdx.x;

        int iy = __mul24(blockIdx.y, blockDim.y) + threadIdx.y; 

        while(ix < imageW && iy < imageH) 

       {

                 float fx = (float) ix + 0.5f; 

                 float fy = (float) iy + 0.5f; 

                 d_Dst[ix + __mul24(iy, imageW)] = tex2D(texSrc, fx, fy);

                 ix += blockDim.x; 

                 iy += blockDim.y; 

         } 

}

#1-- CUDA数组的使用是：
#1.1--下面的代码是绑定CUDA数组到texture及函数调用



static void convolutionRowsGPU(float* d_DstData, cudaArray* d_srcDataArray, cudaChannelFormatDesc floatTex, unsigned int imageHeight, unsigned int imageWidth)

{

          dim3 myBlock(16, 12);

          dim3 myGrid((imageWidth+(blockDim.x -1))/blockDim.x, (imageHeight +(blockDim.y -1))/ blockDim.y ); 

          //#1-- bind array to texture 

          CUDA_SAFE_CALL( cudaBindTextureToArray(texSrc, d_srcDataArray, floatTex)); 

          //#2 -- call kernel function for row convolution 

          convolutionRows_kernel <<< myGrid ,  myBlock >>>(d_DstData, imageHeight, imageWidth);

         cudaGetLastErrorWrapper("convolutionRows_kernel"); 

        //#3--unbind texture CUDA_SAFE_CALL(cudaUnbindTexture(texSrc)); 

}

#1.2--下面是CUDA数组相关的纹理使用



unsigned int imageHeight = 640; 

 unsigned int imageWidth = 480; 

float* h_sorce2DImage_f = new float[ imageHeight  * imageWidth ];

//分配CUDA数组并赋值

cudaArray* d_fDataArray; 

cudaChannelFormatDesc floatTex = cudaCreateChannelDesc<float>();

cudaMallocArrayWrapper(&d_fDataArray, & floatTex, imageWidth, imageHeight); 

//d_sorce2DImage_f 是定义的device端变量，大小为imageHeight * imageWidth *sizeof(float) 

cudaMemcpyToArrayWrapper(d_fDataArray, 0, 0, d_sorce2DImage_f, imageSize * sizeof(float), cudaMemcpyDeviceToDevice); 



//测试1，测试CUDA数组内容

cudaError_t myError = cudaMemcpy2DFromArray(h_sorce2DImage_f, imageWidth, d_fDataArray, 0, 0, imageWidth, imageHeight, cudaMemcpyDeviceToHost);

WiteDataToFile("filted1.txt", h_sorce2DImage_f, imageWidth, imageHeight);



//调用函数，主要是读取纹理内容

convolutionRowsGPU (d_sorce2DImage_f, d_fDataArray,floatTex, imageHeight, imageWidth);

//测试2，测试从纹理中读取的内容

cudaMemcpyWrapper(h_sorce2DImage_f, d_sorce2DImage_f, imageSize * sizeof(float), cudaMemcpyDeviceToHost);

WiteDataToFile("filted2.txt", h_sorce2DImage_f, imageWidth, imageHeight);



delete [] h_sorce2DImage_f ;

注：其中函数cudaMallocArrayWrapper()，cudaMemcpyToArrayWrapper()，cudaMemcpyWrapper()只是为了方便进行了封装，不过只是在调用对应CUDA的函数后加上宏 CUDA_SAFE_CALL ()而已。

#2--绑定cudaMallocPitch分配的线性内存到纹理内存的情况
这里面的原因是因为cudaMalloc分配的pitch不是2的倍数，在绑定内存时会出错。

#2.1--下面的代码是绑定cudaMallocPitch分配的线性内存到纹理内存及相关调用



static void convolutionRowsGPU(float* d_DstData, const float* d_bindLinearMem, unsigned int imageHeight, unsigned int imageWidth, size_t pitch)

{

        dim3 myBlock(16, 12); 

        dim3 myGrid((imageWidth+(blockDim.x -1))/blockDim.x, (imageHeight +(blockDim.y -1))/ blockDim.y );

        size_t offset; 

       //绑定纹理内存

        cudaChannelFormatDesc chennelDesc = cudaCreateChannelDesc(); 

        cudaBindTexture2D(&offset, texSrc, d_bindLinearMem, chennelDesc, imageWidth, imageHeight, pitch);

       //调用核函数

        convolutionRows_kernel <<<myGrid, myBlock>>(d_DstData, imageHeight, imageWidth);     

        cudaGetLastErrorWrapper("convolutionRows_kernel"); 

        //解绑定

        CUDA_SAFE_CALL(cudaUnbindTexture(texSrc)); 

 }

#2.2--下面是 cudaMallocPitch分配的线性内存相关的纹理使用



unsigned int imageHeight = 640; 

 unsigned int imageWidth = 480; 

float* h_sorce2DImage_f = new float[ imageHeight  * imageWidth ];



//cudaMallocPitchWrapper分配内存，并赋值

float* d_linearGloabalMem; 

size_t pitch; 

cudaMallocPitchWrapper(((void**)& d_linearGloabalMem), &pitch, imageWidth * sizeof(float), imageHeight);

myError = cudaMemcpy2D(d_linearGloabalMem, pitch, d_sorce2DImage_f, imageWidth * sizeof(float), imageWidth * sizeof(float), imageHeight, cudaMemcpyDeviceToDevice);



//测试3，测试cudaMallocPitch()分配的内存内容

myError = cudaMemcpy2D(h_sorce2DImage_f, imageWidth * sizeof(float), d_linearGloabalMem, pitch, imageWidth * sizeof(float), imageHeight, cudaMemcpyDeviceToHost);

WiteDataToFile("filted3.txt", h_sorce2DImage_f, imageWidth, imageHeight);



//调用函数，主要是读取纹理内容

convolutionRowsGPU (d_sorce2DImage_f, d_linearGloabalMem, imageHeight, imageWidth, pitch);



//测试4，测试从纹理中读取的内容

cudaMemcpyWrapper(h_sorce2DImage_f, d_sorce2DImage_f, imageSize * sizeof(float), cudaMemcpyDeviceToHost);

WiteDataToFile("filted4.txt", h_sorce2DImage_f, imageWidth, imageHeight);

delete [] h_sorce2DImage_f ;