Bloody Structs

Accelerated Computing CUDA CUDA Programming and Performance
1

Ok, I’m having an issue with structs. The following code will compile under emulation mode, but not in normal mode:

[codebox]

#include <stdio.h>

struct bog

{

int bilbo;

};

struct MyBob

{

__device__ void operator()(int j)

{

	;// do nothing

}

};

template

global void bob(Fex f)

{

struct bog Joe;

struct bog *john;

john = &Joe;

john->bilbo = 0;

f(john->bilbo);

}

int main (int argc, char * const argv)

{

printf(“Hello\n”);

MyBob bubba;

bob<<<1,1>>>(bubba);

return 0;

}

[/codebox]

The specific error that nvcc is spitting out is: “nvcc error : ‘cudafe’ died due to signal 11 (Invalid memory reference)”

Anyone know what the heck is going on here?

Thanks much

Paul

2

Works just fine here with CUDA 2.3.

However compilation crashes with CUDA 2.0.

I bet you have a version inbetween those two.

As an aside, you might simplify your functors and avoid the passing of useless arguments by using some temporary functor instantiation, like

template<typename Fex>

__global__ void bob()

{

struct bog Joe;

struct bog *john;

john = &Joe;

john->bilbo = 0;

Fex()(john->bilbo);

}

int main (int argc, char * const argv[])

{

printf("Hello\n");

bob<Fex><<<1,1>>>();

return 0;

}

I use functors like this a lot. Static member functions are even more useful… the method name helps you document. And you can define as many methods as you like, forming a packaged suite of tools, not just one.

You can pick either method of calling:

Fex::myfunc()

Fex().myfunc()
3

Yep, 2.2. External Media

I just installed 2.3 on Mac OS X, and now, with the same code I posted originally, I’m getting the error: “/tmp/tmpxft_00001608_00000000-1_functortest2.cudafe1.stub.c:16: error: ‘MyBob’ is not a template type”

I’m using the command “nvcc functortest2.cu -arch=sm_13 -o test” to compile. Can you give details on how exactly you got it to compile under 2.3? (i.e. platform, command line command, etc)

Also, I think you mentioned the dummy variables bit before in another thread. It’s a very good idea and I’ll most likely use it in the future. For this program that I’m working on though (converting the LSODA stiff/non-stiff ODE solver to run on CUDA), I think it’s better to pass the dummy variables anyway, since it keeps the interface closer to the original fortran, and makes it a little bit easier to keep track of exactly what’s going on. So for the sake of easy reading, it’s more useful this way. At THIS time, anyway. :D In the future, once the conversion is functionally complete, and I start going back to clean it up and optimize it, I may sneak that version of things in.

Thanks much,

Paul