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不知道怎么用,纹理绑定应该么有问题吧。。