Hi. I’m having problems with circular references in my .cu file code and was wondering if it’s possible to declare classes in .cu or .cuh(which I’m using) files instead of structs? I’m not too familiar with structs and I think my circular refs are because of this. Please can someone let know as I don’t want to change all my code to find it doesn’t work :(
My circular references are like so:
from main.cu:
[codebox]#include “Whitted.cuh”
#include “Pinhole.cuh”
#include “World.cuh”[/codebox]
in Whitted.cuh:
[codebox]#ifndef WHITTED
#define WHITTED
struct Whitted
{
World* world_ptr;
__host__ __device__ Whitted() {}
__host__ __device__ Whitted(World * world):world_ptr(world) {}
__device__ RGBColour trace_ray(Ray &ray, int depth)
{
RGBColour L;
if (depth > world_ptr->vp.max_depth)
{
L.r = 0; L.g = 0; L.b = 0;
return L;
}
else
{
ShadeRec sr(world_ptr->hit_objects(ray, world_ptr->objects));
if (sr.hit_an_object)
{
sr.depth = depth;
sr.ray = ray;
return (RGBColour(1,0,0));//sr.material_ptr->shade(sr));
}
else
{
return (RGBColour(0.5,0.5,0.5));
}
}
}
};
#endif[/codebox]
In World.cuh:
[codebox]#ifndef WORLD
#define WORLD
struct Whitted;
struct Pinhole;
using namespace std;
struct World{
ViewPlane vp;
Whitted* tracer_ptr;
RGBColour background_colour;
vector<Sphere*> objects;
Pinhole* camera_ptr;
//Light* ambient_ptr;
//vector<Light*> lights;
//vector<Material*> materials;
__host__ __device__ World():
background_colour(black),
camera_ptr(NULL)
//ambient_ptr(new Ambient)
{
tracer_ptr = (Whitted*)malloc(sizeOf(Whitted));
tracer_ptr = Whitted(this);
}
__host__ __device__ Vector3D get_direction(Point2D& p)
{
Vector3D dir;
dir = VectorMultMatrix(dir, cam.rayRotationMatrix); // rotate the ray
dir = (camera_ptr->u * p.x) + (camera_ptr->v * p.y) - (camera_ptr->w * camera_ptr->d);
dir.normalize();
return(dir);
}
__host__ __device__ ShadeRec hit_objects(const Ray& ray)
{
ShadeRec sr;
float t = 0;
Normal normal;
Point3D local_hit_point;
float tmin = 1.0E10;
int num_objects = objects.size();
for (int j = 0; j < num_objects; j++)
{
if (objects[j].hit(ray, t, sr) && (t < tmin))
{
sr.hit_an_object = true;
tmin = t;
//sr.material_ptr = objects[j]->get_material();
sr.hit_point = ray.o + t * ray.d;
normal = sr.normal;
local_hit_point = sr.local_hit_point;
}
if(sr.hit_an_object) {
sr.t = tmin;
sr.normal = normal;
sr.local_hit_point = local_hit_point;
}
}
return(sr);
}
};
host device inline void add_material(Sphere* sphere_ptr)
{
objects.push_back(sphere_ptr);
}
host device void set_camera(Pinhole* c_ptr) {
c_ptr->compute_uvw();
camera_ptr = c_ptr;
}
host device void build()
{
int num_samples = 1;
vp.set_res(512, 512);
vp.set_samples(num_samples);
vp.max_depth = 1;
Pinhole* pinhole_ptr = (Pinhole*)malloc(sizeOf(Pinhole));
pinhole_ptr = Pinhole(*this);
pinhole_ptr->set_eye(0, 0, -20);
pinhole_ptr->set_view_distance(500);
pinhole_ptr->set_lookat(0, 0, 0);
set_camera(pinhole_ptr);
Sphere * sphere_ptr;
sphere_ptr = Sphere(Point3D(-5, 0, 0), 5);
//sphere_ptr->set_material(matte_ptr);
add_object(sphere_ptr);
}
#endif[/codebox]
I’ve tried adding things like “struct World;” to the top of the Whitted.cuh but this results in more errors because “partial types are not allowed”. I’ve tried re-ordering the files but it makes no difference. Someone please help!
structs and classes are virtually identical in CUDA (and C++)… the only real difference is by default struct members are public.
But your problems are much, much, deeper and fundamental.
Device code can’t malloc(). It can’t new(). You can’t use std::vectors.
In that sense you’re very, very, limited… you can’t just take working CPU code and try to make it work on the GPU.
This is kind of a first frustration every new CUDA programmer goes through. The environment looks familiar but suddenly quite fundamental ideas and strategies (like a malloc() ) are not allowed, and you realize you’re living in a different land than you expected. Then you get angry (Gaaaa, CUDA sucks!). Then eventually you realize that if you approach things from other directions they’ll work (though you’ll grumble angrily about it). You try it, and suddenly you get this overwhelming monster power result and you go “ohhh, I see, the limitations are worth it when it all works!”
I can’t use vectors?! I thought I read somewhere that you could. How do you propose I dynamically alter the size of arrays then? I think I will get rid of the circular references by making all teh objects global (but in constant memory) that way everything can access everything. But is far from the object oriented approach I outlined in my project report! Goddamn you CUDA
Using CUDA doesn’t stop you from using the normal host compiler too. I write apps where the CUDA code presents a simple C API which the main OO C++ application calls. All dynamic memory (vectors and whatnot) are handled within the main C++ application. (You can cudaMalloc, cudaFree, cudaMallocArray, cudaMemcpy, … all from the main host application).
So does anyone know how I can sort out my circular struct references? I’ve sorted out everything ese but I’m despairing over the struct definitions because they are very tightly coupled and no matter what ordering, they still say they can’t see one of the other structs. OMG! I want to keep the files separated as much as possible without having to “bung” everything into one massive struct and then pass that around my raytracer
CUDA does not support having multiple source files. You can still have multiple files, but only one can be compiled and you need to include all others in there. eg:
class1.cu (do not compile)
class2.cu (do not compile)
kernel.cu: include class1.cu, include class2.cu
What you should do is the same thing you do in C++. Forward declare all the classes you want to use (circularly). Then after all your classes are declared you can implement them since obviously you cannot use (members of) an only declared class.