二维纹理拾取查找表出现了错误

加速计算 CUDA
1

版主：您好！
以下代码是实现高速FFT的内核函数，输入序列存储在shared memory，用texture memory中查找表的方法实现旋转因子，二维纹理函数拾取旋转因子的过程中出现了如下错误：
Wn=tex2D(texRef2D,tid_in_block,stage++); //出现错误的地方
cu(537): error : no operator “=” matches these operands

我知道赋值“=”号两边不匹配，但是具体的原因不知道出现在什么地方？
请版主指教，在此谢谢了！

static global void FFT_T(Complex* const DataIn,Complex* DataOut,const unsigned int N)
{
extern shared Complex sdata;
const unsigned int tid_in_block=threadIdx.x; //线程在线程中的位置
if (tid_in_block<N)
{
sdata[tid_in_block]=DataIn[tid_in_block];
sdata[tid_in_block+N/2]=DataIn[tid_in_block+N/2]; //将数据从global memory读入shared memory
__syncthreads(); //线程块中的线程同步
if (tid_in_block<N/2)
{
int p,q;
Complex Xp,XqWn,Wn;
float stage=0.0;
for (int Ns=1;Ns<N;Ns*=2)
{
p=tid_in_block/NsNs2+tid_in_block%Ns;
q=p+Ns;
Wn=tex2D(texRef2D,tid_in_block,stage++);
XqWn=ComplexMul1(sdata[q],Wn);
Xp=sdata[p];
sdata[p]=ComplexAdd(Xp,XqWn);
sdata[q]=ComplexSub(Xp,XqWn);
__syncthreads();
}
DataOut[p]=sdata[p];
DataOut[q]=sdata[q];
}
}
}

2

楼主您好，无法直接拾取为Complex(或者类似的自定义结构的），
不过您可以这样处理:
float2 tmp = tex2D(…); //(假设您的texture的元素类型为float2或者能normalize出float2)
然后:（假设您的Complex结构含有r和i这2个成员）
Wn.r = tmp.x;
Wn.i = tmp.y;
这样即可。

（注意这样不会导致额外的代码生成的）

3

谢谢千军版主，明天改下程序尝试一下。
您的意思就是，不能拾取自定义的数据结构，只能拾取已定义的元组数据。。

4

楼主您好：

是这样的，tex*返回的数据类型就那么几种。（例如可以返回float2)。
而当返回的类型和您的目标类型无法匹配（以及不存在适当的转换方式）的时候，则会报错。

例如假设您的tex2D返回了float2,但是却无法转换到您的Complex结构，则会出现此问题。此时可以手工转换。（这种情况不会给您带来额外的效率下降的）

5

千军版主：您好！
以下代码实现的高速FFT（一段代码），采样点为N=32（需要5级蝶形运算，每级蝶形运算需要16个旋转因子），将数据从全局存储器读入到共享存储器，而使用的旋转因子事先计算好存储在纹理存储器中，将CUDA数组与纹理绑定，然后利用纹理拾取函数读取其中一行的旋转因子。
程序生成解决方案的时候没有出现错误，但是debug完了之后，没有产生数据，花了好长时间还是没有改出来，请问版主，程序错误的具体原因在什么地方？
谢谢了，版主！

//初始化变换核，旋转因子
Complex W = (Complex)malloc(sizeof(Complex) * K);
for(unsigned int i3=0;i3<K;i3++)
{
W[i3].x=cos(2PI/Ki3);
W[i3].y=-1sin(2PI/Ki3);
}
//将旋转因子存储在二维数组中，之后读取
Complex power[5][16]=
{{W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0],W[0]},
{W[0],W[8],W[0],W[8],W[0],W[8],W[0],W[8],W[0],W[8],W[0],W[8],W[0],W[8],W[0],W[8]},
{W[0],W[4],W[8],W[12],W[0],W[4],W[8],W[12],W[0],W[4],W[8],W[12],W[0],W[4],W[8],W[12]},
{W[0],W[2],W[4],W[6],W[8],W[10],W[12],W[14],W[0],W[2],W[4],W[6],W[8],W[10],W[12],W[14]},
{W[0],W[1],W[2],W[3],W[4],W[5],W[6],W[7],W[8],W[9],W[10],W[11],W[12],W[13],W[14],W[15]}
};
Complex host2D = (Complex)calloc(widthheight, sizeof(Complex)); /构造查找表¨ª
cudaArray cuArray; // CUDA数组
for(int row = 0; row < height; ++row) //初始化内存（查找表）数据
{
for(int col = 0; col < width; ++col)
{
host2D[rowwidth + col] = power[row][width];
}
}
//每个像元由一个float2型数据构成
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc();
checkCudaErrors(cudaMallocArray(&cuArray, &channelDesc, width, height)); //申请显存空间
checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, host2D, sizeof(float)widthheight, cudaMemcpyHostToDevice)); //将内存数据拷贝到显存中，数据已经初始化
//设置纹理参数
texRef2D.addressMode[0]=cudaAddressModeClamp;
texRef2D.addressMode[1]=cudaAddressModeClamp;
texRef2D.filterMode=cudaFilterModePoint;
texRef2D.normalized=false;
checkCudaErrors(cudaBindTextureToArray(texRef2D, cuArray)); // 将显存数据与纹理绑定

//内核函数，其中有二维纹理拾取
static global void FFT_T(Complex* const DataIn,Complex* DataOut,const unsigned int N)
{
extern shared Complex sdata;
const unsigned int tid_in_block=threadIdx.x; //线程在线程块中的位置
if (tid_in_block<N)
{
sdata[tid_in_block]=DataIn[tid_in_block];
sdata[tid_in_block+N/2]=DataIn[tid_in_block+N/2]; //将global memory中的数据读入shared memory
__syncthreads(); //线程块中的线程同步
if (tid_in_block<N/2)
{
unsigned int p,q;
Complex Xp,XqWn;
float2 Wn;
float stage=0.0;
for (unsigned int Ns=1;Ns<N;Ns*=2)
{
p=tid_in_block/NsNs2+tid_in_block%Ns;
q=p+Ns;
Wn=tex2D(texRef2D,tid_in_block,stage++);
XqWn=ComplexMul1(sdata[q],Wn);
Xp=sdata[p];
sdata[p]=ComplexAdd(Xp,XqWn);
sdata[q]=ComplexSub(Xp,XqWn);
__syncthreads();
}
DataOut[p]=sdata[p];
DataOut[q]=sdata[q];
}
}
}

谢谢您了。。。

6

经过ICE的辛苦阅读，未能发现问题。

顺便我酱油路过，您的texRef2D是怎么定义的？（您的代码里没有）

以及，您的代码有无出错？谢谢！

7

LZ您好，下面的一句代码中：

checkCudaErrors(cudaMemcpyToArray(cuArray, 0, 0, host2D, sizeof(float)widthheight, cudaMemcpyHostToDevice)); //将内存数据拷贝到显存中，数据已经初始化

您使用了sizeof(float)，但是您前面纹元是float2的，请您检查下这里是否是问题所在。

祝您编码顺利~

8

ice版主，谢谢您了，我尝试一下，如果还没有成功，我把全部代码粘贴过来。。。

9

LZ您好，希望您能马到成功，真心的…:3_48:

10

ice版主，按照您说的方法，我改过来了，但是还是不行的。等明天上午我把程序粘贴上来。。麻烦您了。。

11

千军版主，我的纹理参照系的定义如下：
texture<float2,2>texRef2D;
应该这个地方没有问题。。。