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[y1024+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[y1024+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(当然也有可能）。供您参考。
哦，，，谢谢