cuda的kernel函数中能调用自定义的结构体吗

[gabriel1861]

我在计算中经常要使用复数，以前用C++的时候都是用的复数模板，现在用cuda不知道是否支持，所以我就自建了一个复数结构体并定义了几个简单的复数运算，做了一个程序测试一下。
程序就是计算一个复数数组中的树的平方和，编译的时候编译器提示出现错误终止，不知道原因何在？我的系统是win7 x64旗舰版，在VS2008下用cuda3.2的toolkit编译的，配置安装没有问题，以前做的几个简单的例子都可以用。源代码如下
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>

#define DATA_SIZE 1048576
#define BLOCK_NUM 32
#define THREAD_NUM 256

struct Complex{
float r,i;

Complex(float _x, float _y);
Complex(const Complex &c);
Complex();
~Complex(){}

Complex operator*(const Complex &c)const;
Complex operator=(const Complex &c);
Complex operator+(const Complex &c)const;
Complex operator+=(const Complex &c);
};

Complex::Complex(float _x, float _y)
{
r=_x; i=_y;
}
Complex::Complex(const Complex &c)
{
r=c.r; i=c.i;
}
Complex::Complex()
{
r=0.0; i=0.0;
}

Complex Complex::operator*(const Complex & c)const
{
return Complex(r*c.r-i*c.i,r*c.i+i*c.r);
}

Complex Complex::operator=(const Complex & c)
{
r=c.r; i=c.i;
return *this;
}

Complex Complex::operator+(const Complex &c)const
{
return Complex(r+c.r, r+c.i);
}

Complex Complex::operator+=(const Complex &c)
{
r += c.r; i += c.i;
return *this;
}


void GenerateNumber(Complex* number, int size)
{
for(int i=0; i<size; i++)
number[i] = Complex(i*1.0,(i+1)*1.0);
}

__global__ static void sumOfSquares(Complex *num, Complex* result, clock_t* time)
{
extern __shared__ Complex shared[];
const int tid = threadIdx.x;
const int bid = blockIdx.x;
int i;
int offset;
if(tid == 0)time[bid] = clock();
shared[tid] = Complex(0.0,0.0);
for(i = bid*THREAD_NUM+tid; i<DATA_SIZE;
i += BLOCK_NUM*THREAD_NUM)
shared[tid] += num[i]*num[i];

__syncthreads();
offset = THREAD_NUM / 2;
while(offset>0)
{
if(tid<offset)
shared[tid] += shared[tid+offset];
offset >>= 1;
__syncthreads();
}

if(tid==0)
{
result[bid] = shared[0];
if(tid==0)time[bid+BLOCK_NUM] = clock();
}
}

Complex data[DATA_SIZE];

int main()
{
GenerateNumber(data, DATA_SIZE);

Complex *gpudata, *result;
clock_t *time;
cudaMalloc((void**)&gpudata, sizeof(Complex)*DATA_SIZE);
cudaMalloc((void**)&result, sizeof(Complex)*BLOCK_NUM);
cudaMalloc((void**)&time, sizeof(clock_t)*BLOCK_NUM*2);
cudaMemcpy(gpudata, data, sizeof(Complex)*DATA_SIZE, cudaMemcpyHostToDevice);

sumOfSquares<<<BLOCK_NUM, THREAD_NUM, THREAD_NUM*sizeof(Complex)>>>
(gpudata, result, time);

Complex sum[BLOCK_NUM];
clock_t time_used[BLOCK_NUM*2];
cudaMemcpy(&sum, result, sizeof(Complex)*BLOCK_NUM, cudaMemcpyDeviceToHost);
cudaMemcpy(&time_used, time, sizeof(clock_t)*BLOCK_NUM*2, cudaMemcpyDeviceToHost);
cudaFree(gpudata);
cudaFree(result);
cudaFree(time);

Complex final_sum;
for(int i=0; i<BLOCK_NUM; i++)
final_sum += sum[i];

clock_t min_start, max_end;
min_start = time_used[0];
max_end = time_used[BLOCK_NUM];
for(int i=1; i<BLOCK_NUM; i++)
{
if(min_start > time_used[i])
min_start = time_used[i];
if(max_end < time_used[i+BLOCK_NUM])
max_end = time_used[i+BLOCK_NUM];
}

printf("sum: (%f , %f) time: %d\n", final_sum.r, final_sum.i , max_end-min_start);

return 0;
}

在给kernel函数传参的时候我用的是自定义的Complex结构体，kernel函数中复数的乘法和加法运算，我是调用的外部的定义的重载操作符，不知道是不是这个的问题。

应该如何修改，请赐教

[linxxx3]

你确定这个是struct而不是class么。。
另外关于CUDA是否支持运算符重载不清楚，等高人。
lz先把编译错误提示贴出来看看吧

[gabriel1861]

[Quote=引用 2 楼 linxxx3 的回复:]
你确定这个是struct而不是class么。。
另外关于CUDA是否支持运算符重载不清楚，等高人。
lz先把编译错误提示贴出来看看吧
[/Quote]

如果不使用操作符重载，那如何能是复数运算变得更简洁呢，编译错误如下

“cudafe.exe已停止工作
出现了一个问题，导致程序停止正常工作。如果有可用的解决方案，Windows将关闭程序并通知您”

[zhuxueling]

其实这个很简单。
我想知道楼主在哪里工作。。

[gabriel1861]

[Quote=引用 4 楼 zhuxueling 的回复:]
其实这个很简单。
我想知道楼主在哪里工作。。
[/Quote]

问题已解决，无意中发现应用指南里就有讲如何进行操作符重载的。看来多读文档才是王道

我是做电磁场计算问题的，接触的变量几乎都是复数

[gabriel1861]

[Quote=引用 2 楼 linxxx3 的回复:]
你确定这个是struct而不是class么。。
另外关于CUDA是否支持运算符重载不清楚，等高人。
lz先把编译错误提示贴出来看看吧
[/Quote]

操作符是可以重载的，cuda文档里有讲解
还是很感谢你的帮助，分全给你了

[水月洞天]

能问一下楼主是在哪个资料上看到的运算符重载，求分享呀[e-mail]

[战斗猿]

把所有的操作符重载函数设置成 __device__ 类型，编译通过

[水月洞天]

谢谢。我试试...

[gabriel1861]

我后来没用操作符重载，因为发现cuda里面有一个cuComplex.h头文件里就有关于复数的计算，实际就是用内置变量类型float2来代替复数。帖子是前年发的，后来一直没有关住，这次最后再来说明以下

[cai_niao_yi_zhi]

内核函数中好像只能支持部分的C++特性，而且是SDK3.0以上的,不知道能不能支持C++的结构体！