求助CUDA三维纹理使用问题

global void rgb2gray(floatgray)
{
unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
// read from 3D texturea
float u=(float)x;
float v=(float)y;
float4 rgb=tex3D(texImage,u,v,0);//z的坐标怎么确定
//float4 rgb1=tex3D(texImage,x,y,1);//z的坐标怎么确定
//float4 rgb2=tex3D(texImage,x,y,2);//z的坐标怎么确定
gray[y
1024+x]=rgb.x;

问题是一个图像数据大小为10247683,现在我要把大分为三个二位数组,用tex3D以后始终不对,得不到正确的数据,不知道是不是坐标问题,再网上找了很久的答案,找不到,麻烦哥哥们给我讲解哈。

LZ您好,我对三维纹理不甚熟悉,暂时无法回答您的问题。

有待其他熟悉该内容的斑竹/网友等补充回答了。

祝您好运~

楼主可否发一下cudaArray的建立,复制,以及绑定、纹理引用的具体声明等过程?
楼主您的写法是无问题的。但可能问题出现在上述方面。请补充代码并给出如何不对(例如您的不对指读取的都是0或者都是边界值之类的)

此外,我看到您的元素类型是float4, 请问是您是一大张10247683的,每个点都是float4的图么?(还是每个点一共32位?而不是128位的float4?)如果只有普通的32位的图,那么用float4显然是不对的。您此时也许需要char4. 当然。如果的确每个元素是128-bit的。那么float4是正确的,请无视。

此外,如果您是3张独立的1024768的图片,而不是一张立体上有关联的3D的1024768*3的大图,建议使用2D Layered Texture(也叫纹理数组), 他们的cache策略不同(或者说,在cudaArray里的数据具体存放格式不同)。可能对性能有提升。用法基本一样。

以及,如果楼主是3张独立的1024*658,

谢谢:P。。。。。。。。。。。。。。。。

include “cuda_runtime.h”
include “device_launch_parameters.h”

include <stdio.h>
include <stdlib.h>
include “imageDenoising.h”
uchar4 *h_Src;
int imageW, imageH;
typedef unsigned char uchar;

texture<uchar4, 3, cudaReadModeNormalizedFloat> texImage;
cudaChannelFormatDesc uchar4tex = cudaCreateChannelDesc();
cudaArray *a_Src;

global void rgb2gray(floatgray)
{
unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
// read from 3D texturea
float u=(float)x;
float v=(float)y;
float4 rgb=tex3D(texImage,u,v,0);//z的坐标怎么确定
//float4 rgb1=tex3D(texImage,x,y,1);//z的坐标怎么确定
//float4 rgb2=tex3D(texImage,x,y,2);//z的坐标怎么确定
gray[y
1024+x]=rgb.x;

if ((x < 1024) && (y < 768))
{
	
	// write output color
	
	//gray[y*1024+x]=1;
}

}

cudaError_t CUDA_Bind2TextureArray()
{
return cudaBindTextureToArray(texImage, a_Src);
}
cudaError_t CUDA_UnbindTexture()
{
return cudaUnbindTexture(texImage);
}

cudaError_t CUDA_MallocArray(uchar4 **h_Src, int imageW, int imageH)
{
cudaError_t error;

error = cudaMallocArray(&a_Src, &uchar4tex, imageW, imageH);
error = cudaMemcpyToArray(a_Src, 0, 0,
	*h_Src, imageW * imageH * sizeof(uchar4),
	cudaMemcpyHostToDevice
	);

return error;

}

cudaError_t CUDA_FreeArray()
{
return cudaFreeArray(a_Src);
}

int main()
{
printf(“\n");
printf(" SR Processing Beginning !\n");
printf("
\n”);

const char filename[]="E:\\cuda\\color-cam0-f000.bmp";
LoadBMPFile(&h_Src, &imageW, &imageH, filename);

//将uchar4 拷贝到设备端
float *gray=NULL;
float *d_gray=NULL;
cudaMalloc((void**)&d_gray,1024*768*sizeof(float));

gray = (float*)malloc(1024768sizeof(float));
CUDA_MallocArray(&h_Src,imageW,imageH);
CUDA_Bind2TextureArray();

//设置纹理运行的参数;
texImage.normalized=0;
texImage.filterMode=cudaFilterModePoint;


dim3 dimBlock(16,16);  
dim3 dimGrid((768+dimBlock.x-1)/(dimBlock.x), (1024+dimBlock.y-1)/(dimBlock.y));

rgb2gray<<<dimGrid,dimBlock>>>(d_gray);

cudaMemcpy((void*)(gray), (void*)(d_gray), 1024*768*sizeof(float), cudaMemcpyDeviceToHost);

for (int r = 0; r < 1024; r++)  
{  
	for (int c = 0; c < 768; c++)  
	{  
		printf("%5.2f",gray[c*1024+r]);
	}  
	printf("\n");  
}  

CUDA_UnbindTexture();
CUDA_FreeArray();
cudaFree(d_gray);
free(gray);

}

数组是一张10247683的,,,每个点都是float4,关键是tex3D不知道怎么用,纹理绑定应该么有问题吧。。

这不是没有问题,而是相当有问题,楼主您看如下方面:

楼主您写到“数组是一张10247683的,每个点都是float4”, 但是您的代码:
cudaMallocArray(…); //您确定你不需要3D Array么?
cudaMemcpyToArray(…); //您确定您不需要Memcpy3D么?
以及
texture<uchar4, 3, cudaReadModeNormalizedFloat> texImage; //您确定您不需要float4么?请同时参考我的上个回复。仔细想象你是需要char4之类的,还是需要float4之类的。您虽然已经回复我说您需要float4, 但请再想想。
cudaChannelFormatDesc uchar4tex = cudaCreateChannelDesc(); //还有这里?
等等处

和您对您的1024 * 768 * 3的图像的描述,和您的float4每个元素的描述都矛盾。

考虑过您已经阐述过您的意图,并且在5#再次强调您认为大小,格式了。我将不再次询问。
建议您仔细学习下网上或者examples里的3D Texture的例子,然后您将明白如何声明一个元素为float4(或者说channel为float4),3层的纹理引用,以及如何分配大小为多少多少多少的cudaArray,以及如何复制数据过去。

以及,您可能还有更多其他问题,先改了上文再说,不要浮躁,先学习,上文例子足够您学习了。不要随便看到别人的代码就乱改,然后不乱折腾,最后似乎弄出来了,还不知道怎么回事。静心吧。唉。

我也才开始学习,,cuda中文资料太少了,英文的有看不懂,,,,谢谢版主,我在改改,,,

我也才开始学习,,cuda中文资料太少了,英文的有看不懂,,,,谢谢版主,我在改改,,,

此外,如果您的代码是抄袭的或者是引用的师兄的,那么请注意他的本来意图,他使用了cudaReadModeNormalizedFloat,将原来每个是1B的char4或者uchar4之类的值转换为0.0-1.0的float, 也许这才是他的本意,而不是真的原始图像就是float4(当然也有可能)。供您参考。

哦,,,谢谢