How to use an array of classes on the GPU

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

I’ve been trying to use some classes on the GPU lately. It works if I just use one, but if I try to use an array of classes it doesn’t. Here’s the code that works:

typedef GPU_AbstractSubsystem** GPU_AbstractSubsystemPtr;

global void
TakeAScan(float4* newPos, float4* newVel,float4* oldPos, float4* oldVel,float* volts,
int numIons,float* randomNumbers,float* randomGaussians,int offSet1,int offSet2,
float4* acquisitionBuffer,GPU_AbstractSubsystemPtr g_subsystem)
{
// get the index of this ion
int index = __mul24(blockIdx.x,blockDim.x) + threadIdx.x;
// fetch the current position and velocity from for this ion
float4 pos = oldPos[index];
float4 vel = oldVel[index];
float4 acquire = {0.0,0.0,0.0,0.0};
// if ion is within boundaries
if(1)//RecordData(pos,vel,acquire))
{
(*g_subsystem)->AdvanceParticle(pos,vel,volts);
}
// store new values
newPos[index] = pos;
newVel[index] = vel;
acquisitionBuffer[index] = acquire;
}

I thought I should be able to simply change the function to use GPU_AbstractSubsystemPtr* instead of GPU_AbstractSubsystemPtr, and then call the function with: (*g_subsystem[0])->AdvanceParticle(pos,vel,volts);
And then I pass in a reference to the array of classes &d_subsystem[0] for g_subsystem.

Eventually I would iterate through each subsystem, you know. Is this basic idea not correct? I’m not sure if the syntax is screwing me up, or some other error. Anyone have an idea? I appreciate any advice.

Thanks!

2

if my tv is fuzzy i give it a little smack upside the screen…perhaps hitting it could help External Media

3

I hope all that is connected to a power-supply. :-)

4

hi pojken

Why don’t you show us the code that gives you error? And what kind of error exactly? How is your g_subsystem initialized?

5

I wish I could show the error, unfortunately I have built a GUI with wxWidgets, and this CUDA error causes wxWidgets to quit before I can see the message in my log window. I need to figure out how to write log window to a file. Maybe the first step is to figure out what the error says. As far as how g_subsystem was initialized, I had been looking at the newdelete example in the SDK. They pass in a ** of the class to a helper functions like below.

global void
SubSystemCreate(GPU_AbstractSubsystemPtr g_subsystem,interpolation_t interpRadial,interpolation_t interpAxial,integrator_t integrator,acceleration_t acceleration,
float gasMassAmu,float gasTempK,float gasPressureMTorr,float offsetX,float offsetY,float offsetZ)
{
*g_subsystem = new GPU_Subsystem(interpRadial,interpAxial,integrator,acceleration,gasMassAmu,gasTempK,gasPressureMTorr,offsetX,offsetY,offsetZ);
}
extern"C++"
void
GPU_SubSystemCreate(GPU_AbstractSubsystemPtr g_subsystem,interpolation_t interpRadial,interpolation_t interpAxial,integrator_t integrator,acceleration_t acceleration,float gasMassAmu,float gasTempK,
float gasPressureMTorr,float offsetX,float offsetY,float offsetZ)
{
SubSystemCreate<<<1,1>>>(g_subsystem,INTERP_IONTRAP2D,INTERP_IONTRAP3D,INTEGRATOR_EULER,ACCEL_COMBINE2D3D,gasMassAmu,gasTempK,gasPressureMTorr,offsetX,offsetY,offsetZ);
}

// host code
// each subsystem has a host class with this data member
GPU_AbstractSubsystemPtr d_subsystem_;
// allocate memory for device class
cutilSafeCall(cudaMalloc((void**)&d_subsystem_, sizeof(GPU_AbstractSubsystemPtr)));
// create class on device
GPU_SubSystemCreate(d_subsystem_,INTERP_IONTRAP2D,INTERP_IONTRAP3D,INTEGRATOR_EULER,ACCEL_COMBINE2D3D,gasMass,gasTemp,gasPressure,globalOffsets_.x,globalOffsets_.y,globalOffsets_.z);

My plan was to collect all the subsystem classes with:
GPU_AbstractSubsystemPtr arrayOfSubSystems[MAXIMUM_SUBSYSTEMS];
// then do this for each subsystem
arrayOfSubSystems[0] = instruments_[0]->GetSubSystem();

and then call the TakeAScan function by passing arrayOfSubSystems.

Playing around with classes in a CPU only project, I see it’s possible to just pass an array of classes to a function with
someClass* class[10];
class[0] = new someClass();

void function(someClass** class)
{
class[0]->action();
}

That’s what I’d want to do on the device, but so far I can’t get it to work. This need to use ** syntax and dereference with (*class)->action() as in the newdelete SDK example seems kind of cumbersome. Oh well, I’ll experiment around here and post if I figure it out.

6

The cuda error returned after calling TakeAScan is the good old “unspecified launch failure” :)

7

some problem with your naming:

typedef GPU_AbstractSubsystem** GPU_AbstractSubsystem_Array;
typedef GPU_AbstractSubsystem* GPU_AbstractSubsystemPtr;

perhaps will help you a bit.

From what I see, your SubSystemCreate only initializes a single subsystem class.

8

If you have

typedef GPU_AbstractSubsystem** GPU_AbstractSubsystemPtr;

Then you can’t write

*g_pointer = new ...

because you need to initialize g_pointer first:

g_pointer = new (GPU_AbstractSubsystem*)