CUDA 二维纹理 的获取和使用

[YCMyTot]

最近，在使用CUDA 纹理内存的时候碰到了，比较难搞的事情（PS : 第一次使用 纹理内存。）
使用的是二维uchar 型 纹理，归一化寻址方式， 在构建了 高斯金字塔之后， 进行 高斯 差分 操作。
相关代码如下：

// 声明纹理内存

// 采用 2维float纹理的方式 进行内存上的初步优化
// cudaReadModeNormalizedFloat 代表的是 读取的纹理内存的返回值 是 浮点型 
texture < uchar , 2 , cudaReadModeNormalizedFloat > texRef_1 ;
texture < uchar , 2 , cudaReadModeNormalizedFloat > texRef_2 ;

// 内核函数

__global__ void DogBlur(unsigned char *d_dog_data,int *d_img_width2 , int * d_image_height2 )
{
	int i=blockIdx.x;
	int j=blockDim.x*blockIdx.y+threadIdx.x;
	int height=gridDim.x;
	int width=*d_img_width2;
	int img_height = *d_image_height2 ;
	
	// 设定寻址方式
	float x = j / (float) width ;
	float y = i / (float) img_height ;

	if(j<width)
	{

		// 使用 归一化的 纹理寻址 方式

		float temp_data = fabs( tex2D( texRef_1 , x , y ) - tex2D( texRef_2 , x , y ) )  ;

		d_dog_data[ i * width + j ] = (uchar)(temp_data * 255) ;

		if(d_dog_data[i*width+j]<16)
		{
			d_dog_data[i*width+j]*=d_dog_data[i*width+j];
		}
		else
		{
			d_dog_data[i*width+j]=255;
		}

		if(d_dog_data[i*width+j]<16)
		{
			d_dog_data[i*width+j]=0;
		}

	}
}

//主机端 调用内核函数，对纹理进行设置，绑定

	unsigned char *d_dog_data ;
	int * d_img_width2 ;
	int * d_img_height2 ;
	// 定义 两个 GPU 端的数组 
	cudaArray * cuArray_1 , *cuArray_2 ;			// 分配的内存的大小
			size_t size = scale_space[i][j]->height * scale_space[i][j]->width * sizeof( uchar ) ;

			int img_width = scale_space[i][j]->widthStep ;

			int image_width = scale_space[i][j]->width ;
			int image_height = scale_space[i][j]->height ;


			// 开辟 内存 空间(GPU端)

			// 对纹理内存进行设置并调用

			// 将 纹理内存的 类型 改成 uchar 类型
			cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc <uchar> () ;
			cudaMallocArray ( &cuArray_1 , &channelDesc , scale_space[i][j]->width , scale_space[i][j]->height ) ;
			cudaMallocArray ( &cuArray_2 , &channelDesc , scale_space[i][j]->width , scale_space[i][j]->height ) ;
			
			// 分配 空间 d_img_width2 以及 d_dog_data , 以及 d_img_height2 ;
			cudaMalloc( (void**)&d_img_width2 , sizeof(int) ) ;
			cudaMalloc( (void**)&d_dog_data , size ) ; 
			cudaMalloc( (void**)&d_img_height2 , sizeof(int) ) ;

			// 绑定内存 Array -> Texture
			cudaBindTextureToArray ( texRef_1 , cuArray_1 );
			cudaBindTextureToArray ( texRef_2 , cuArray_2 );


			// 将 CPU 端的 数据 拷贝 到 cuda 数组中
			cudaMemcpyToArray( cuArray_1 , 0 , 0 , data1 , 
				sizeof( uchar ) * scale_space[i][j]->width * scale_space[i][j]->height , 
				cudaMemcpyHostToDevice ) ;

			cudaMemcpyToArray( cuArray_2 , 0 , 0 , data2 ,
				sizeof( uchar ) * scale_space[i][j]->height * scale_space[i][j]->width , 
				cudaMemcpyHostToDevice ) ;

			// 拷贝 内存到 GPU : d_img_width2 以及 d_dog_data
			cudaMemcpy( d_img_width2 , &img_width , sizeof(int) , cudaMemcpyHostToDevice ) ;
			cudaMemcpy( d_dog_data , dog_data , size , cudaMemcpyHostToDevice ) ;
			cudaMemcpy( d_img_height2 , &image_height , sizeof(int) , cudaMemcpyHostToDevice ) ;

还望 各位 大牛们帮小弟看一下这个问题！！

[YCMyTot]

这就是发生 错误的 结果，不晓得 是哪里错了！

[YCMyTot]

哦！多谢提醒！ // 内核函数的 调用

			// 线程的分配 
			if( image_width %768 != 0  )
			{
				dim3 Grid ( image_height ,
					( image_width / ( image_width%768 )) + 1 ) ;
				dim3 Block ( image_width % 768 ) ;

				// 高斯差分内核的调用
				DogBlur <<< Grid,Block >>> ( d_dog_data , d_img_width , d_img_height2) ;

				// 拷贝 内存 ( GPU -> CPU )
				cudaMemcpy( dog_data , d_dog_data , size , cudaMemcpyDeviceToHost );

			}

			//// 对 线程的另一种 情况 进行 分类
			else 
			{
				// 线程的分配
				dim3 Grid ( image_height , image_width/(768)  + 1 ) ;
				dim3 Block ( 768 ) ;

				// 高斯差分 内核的调用
				DogBlur <<< Grid , Block >>> ( d_dog_data , d_img_width , d_img_height2 ) ;

				// 内存的拷贝 ( GPU->CPU )
				cudaMemcpy( dog_data , d_dog_data , size , cudaMemcpyDeviceToHost ) ;

			}

多谢了！！

[adagio_chen]

内核函数调用在哪里？

[adagio_chen]

这个是什么意思？

       if(d_dog_data[i*width+j]<16)
        {
            d_dog_data[i*width+j]*=d_dog_data[i*width+j];
        }
        else
        {
            d_dog_data[i*width+j]=255;
        }
 
        if(d_dog_data[i*width+j]<16)
        {
            d_dog_data[i*width+j]=0;
        }

[YCMyTot]

为了 让 点现实的更加清楚，就将图片 变成 二值图像！！