Hello all,
I’m trying to port Radius-CUDA (a GPGPU raytracer) to work on Linux. The original windows code is available here. Radius-CUDA is a CUDA adaptation of Radius (which is a Linux raytracer to begin with), so the conversion itself shouldn’t be too painful… or so I thought. The problem is with a nested anonymous struct that appears in the CUDA code. The relevant code is as follows:
struct vec_t {
float x,y,z;
__device__ vec_t() {}
__device__ vec_t(const float a, const float b, const float c) : x(a), y(b), z(c) {}
__device__ vec_t(const float a) : x(a), y(a), z(a) {}
__device__ vec_t operator+(const vec_t &v) const { return vec_t(x+v.x,y+v.y,z+v.z); }
__device__ vec_t operator-(const vec_t &v) const { return vec_t(x-v.x,y-v.y,z-v.z); }
__device__ vec_t operator-() const { return vec_t(-x,-y,-z); }
__device__ vec_t operator*(const float d) const { return vec_t(x*d,y*d,z*d); }
__device__ vec_t cross(const vec_t &v) const { return vec_t(y*v.z-z*v.y,z*v.x-x*v.z,x*v.y-y*v.x); }
__device__ vec_t normalize() const { return *this * (1.f/sqrtf(magsqr())); }
__device__ float norm() const { return sqrtf(magsqr()); }
__device__ float dot(const vec_t &v) const { return x*v.x+y*v.y+z*v.z; }
__device__ float magsqr() const { return dot(*this); }
__device__ float get_min() const { return fminf(fminf(x,y),z); }
__device__ float get_max() const { return fmaxf(fmaxf(x,y),z); }
__device__ vec_t perm_x() const { return vec_t(x, y, z); }
__device__ vec_t perm_y() const { return vec_t(y, z, x); }
__device__ vec_t perm_z() const { return vec_t(z, x, y); }
};
struct sampler_t {
vec_t top, dx, dy;
__device__ vec_t map(const point_t &screen) const {
return vec_t(top + dx*float(screen.x) + dy*float(screen.y));
}
};
struct camera_t {
struct {
vec_t eye,dir, up, right;
float fovx;
int world_up_index;
sampler_t sampler;
};
__device__ inline void look_at(const vec_t &target, const int up_idx = -1);
__device__ inline void set_fovx(const float degree) { fovx = DEGRAD(degree) * .5f; }
__device__ inline void set_eye(const vec_t &v) { eye = v; }
__device__ inline float get_fovx() const { return RADDEG(fovx)*2.f; }
__device__ inline const vec_t &get_eye() const { return eye; }
__device__ inline const vec_t &get_up() const { return up; }
__device__ inline const vec_t &get_dir() const { return dir; }
__device__ inline const vec_t &get_right() const { return right; }
__device__ inline void set_world_up_index(int idx) { world_up_index = idx; }
__device__ inline int get_world_up_index() const { return world_up_index; }
};
The compiler chokes on the following declaration:
__device__ __constant__ camera_t cu_cam;
The compiler errors are:
I’m using gcc-4.2.4 with the 177.80 drivers and the CUDA-2.0 SDK. However, I’ve also tried compiling with gcc-4.1 and get the same errors. Obviously, this code compiles just swimmingly using MSVC in Visual Studio 2008. Any ideas?
Thanks,
Rob