From what I can tell from Troy's slides on Kamasu, he has separated boost::proto from what has to be compiled for the GPU (I may be wrong about this). So code that needs boost::proto compiles separately with regular g++, and code that needs to run on GPU compiles with nvcc and both get linked finally. Maybe the real Troy should speak up..

Manjunath On Mon, Feb 8, 2010 at 1:47 PM, Eric Jourdanneau wrote:

Manjunath Kudlur wrote:

...

Hi Eric,

I am assuming that you are referring to NVIDIA's CUDA compiler, when you say nvcc. I don't know if nvcc has complete C++ compliance/support for CUDA programs. You are better off posting this question to one of NVIDIA's forums http://forums.nvidia.com/index.php.

Yes, indeed. Since You offer me the chance to make a small change in my first post, I would like to tell that the error is on line 314 for the 1.38 version of boost, and on line 422 for the 1.42 one. Well, there are some slides from Troy D. Straszheim, where he talks about nvcc and proto, but he did not seem to have this problem, or at least did not tell about them. Eric

...

Manjunath

On Mon, Feb 8, 2010 at 12:05 PM, Eric Jourdanneau wrote:

...

Hi everybody, Nice to meet you :)

I would like to test proto with nvcc, and when I compile the simple Hello world example of 1.42, I get the following error message: /usr/include/boost/proto/detail/decltype.hpp:314: error: invalid use of template-name ‘boost::type’ without an argument list.

Of course it works with g++ I just don't know what to do, can somebody help me on that point? Thanks Eric

_______________________________________________ Boost-users mailing list Boost-users@lists.boost.org http://lists.boost.org/mailman/listinfo.cgi/boost-users

_______________________________________________ Boost-users mailing list Boost-users@lists.boost.org http://lists.boost.org/mailman/listinfo.cgi/boost-users

_______________________________________________ Boost-users mailing list Boost-users@lists.boost.org http://lists.boost.org/mailman/listinfo.cgi/boost-users

Joel, On Thu, Feb 11, 2010 at 4:49 AM, Joel Falcou wrote:

Hi,

I created the following ticket:

https://svn.boost.org/trac/boost/ticket/3919

that gathers the basic files needed to actually make nvcc a proper compiler target for boost.

THe patch contain a macro called BOOST_GPU_ENABLED that resolves as :

__host__ __device__

This macro can be used on any function or method declaration in Boost so they can be used as device AND host function.

Making boost fully GPU compatible (ie compiling boost component inside a device function) may require to add this on EVERY possible boost function and method. Considering he amount of work needed I prefer first put this up and see what people make out of it.

Concernign proto itself, i finally understood why the previous patch was needed. Seems a BOOST_NO_XXX is still missing. I'll investigate and propose a proper patch for that.

Side note: If anyone is interested into pursuing efforts in this direction, we'll gladly accept help (evn just for testing) and comments so we can actually tackle a proper patching process on interesting components.

Regards _______________________________________________ Boost-users mailing list Boost-users@lists.boost.org http://lists.boost.org/mailman/listinfo.cgi/boost-users

Thanks for your efforts. I was curious as to what version of CUDA you have tested these changes with. Also, I would rather call it something like BOOST_CUDA_DEVICE_ENABLED, rather than BOOST_GPU_ENABLED, but that's just my taste. Manjunath

On Thu, Feb 11, 2010 at 10:27 AM, joel falcou wrote:

Manjunath Kudlur wrote:

...

Thanks for your efforts. I was curious as to what version of CUDA you have tested these changes with.

IIRC it was 2.3 Eric can confirm if needed

Which platform was this tested on? I wrote the following simple program : #include int main(void) { } and I am getting the following from nvcc : nvcc error : 'cudafe++' died due to signal 11 (Invalid memory reference) This is preceded by lot of warnings warnings of the form /home/boost_1_42_0/boost/proto/make_expr.hpp(873): warning: class template "boost::proto::detail::implicit_expr_2" defines no constructor to initialize the following: reference member "boost::proto::detail::implicit_expr_2::a0" reference member "boost::proto::detail::implicit_expr_2::a1" Manjunath

Hello, Manjunath Kudlur wrote:

On Thu, Feb 11, 2010 at 10:27 AM, joel falcou wrote:

...

Manjunath Kudlur wrote:

...

Thanks for your efforts. I was curious as to what version of CUDA you have tested these changes with.

IIRC it was 2.3 Eric can confirm if needed

Which platform was this tested on? I tested the Tarray example of the proto documentation. It compiles and run fine on the host. The example you give, too. uname -a gives : Linux machine 2.6.31-17-generic #54-Ubuntu SMP Thu Dec 10 17:01:44 UTC 2009 x86_64 GNU/Linux

Eric J

I wrote the following simple program : #include

int main(void) { }

and I am getting the following from nvcc :

nvcc error : 'cudafe++' died due to signal 11 (Invalid memory reference)

This is preceded by lot of warnings warnings of the form

/home/boost_1_42_0/boost/proto/make_expr.hpp(873): warning: class template "boost::proto::detail::implicit_expr_2" defines no constructor to initialize the following: reference member "boost::proto::detail::implicit_expr_2::a0" reference member "boost::proto::detail::implicit_expr_2::a1"

Manjunath _______________________________________________ Boost-users mailing list Boost-users@lists.boost.org http://lists.boost.org/mailman/listinfo.cgi/boost-users

K. Noel Belcourt wrote:

I think I can offer up some testing cycles on our 256 node system but unfortunately not much time to do anything other than start running the tests. Would that help?

I you have nvcc handy on your system, just running the unti and regression test could lead to a list of failure that we didn't catch already and give us grain to grind for pursuing this. -- ___________________________________________ Joel Falcou - Assistant Professor PARALL Team - LRI - Universite Paris Sud XI Tel : (+33)1 69 15 66 35

Eric Niebler wrote:

Whoa. Even inline functions and methods? This has to be investigated.

I can't speak for anybody else, but I'm pretty sure I'm not ok with that. Nobody will ;) That's the unfun part. I mean, one can manage that if needed in a local copy of bosot (I think it's where we'll end up) but it can't be decently set-up as a global policy.

Anyway, a first step could be to check which libraries makes sense to run on a GPU and other exotic platform (thread, iostreams, asio make no sense for example, same for filesystem or mpi).

You should move this discussion to development list with a more appropriate subject. Up its way

On Feb 8, 2010, at 6:07 PM, Eric Niebler wrote:

(Manjunath, please don't top-post. Rearranging things ...)

On 2/9/2010 8:35 AM, Manjunath Kudlur wrote:

...

On Mon, Feb 8, 2010 at 12:05 PM, Eric Jourdanneau wrote:

...

I would like to test proto with nvcc, and when I compile the simple Hello world example of 1.42, I get the following error message: /usr/include/boost/proto/detail/decltype.hpp:314: error: invalid use of template-name ‘boost::type’ without an argument list.

Of course it works with g++ I just don't know what to do, can somebody help me on that point?

I am assuming that you are referring to NVIDIA's CUDA compiler, when you say nvcc. I don't know if nvcc has complete C++ compliance/ support for CUDA programs. You are better off posting this question to one of NVIDIA's forums http://forums.nvidia.com/index.php.

I've never tested proto with nvcc, so I'm not surprised it doesn't work. I would need someone familiar with the quirks of this compiler (Troy?) to help get proto working with it. And there's no guarantee it's possible -- proto requires a fairly compliant compiler.

Boost is not compilable with nvcc. The CUDA compiler (nvcc) is, in essence, a C compiler not C++. You'll need to read the documentation but C++ code is compiled by a host compiler (g++, etc...) while nvcc is used to compile device (GPU) specific C-like code. -- Noel

On Feb 8, 2010, at 7:37 PM, Emil Dotchevski wrote:

On Mon, Feb 8, 2010 at 5:21 PM, Belcourt, Kenneth wrote:

...

The CUDA compiler (nvcc) is, in essence, a C compiler not C++.

This is incorrect.

...

You'll need to read the documentation but C++ code is compiled by a host compiler (g++, etc...) while nvcc is used to compile device (GPU) specific C-like code.

The device-specific compiler has a C++ front-end, complete with templates, specializations, etc. + NVIDIA-specific extensions. The complication is that code that targets both the "device" (GPU) and the "host" (CPU) needs to be compatible with at least 2 front-ends (NVCC and GCC/MSVC), but that's not unusual for Boost in general.

Oh, I stand corrected. Guess things have changed quite a bit since I last worked with CUDA. Apologies. -- Noel

On Feb 8, 2010, at 7:37 PM, Emil Dotchevski wrote:

On Mon, Feb 8, 2010 at 5:21 PM, Belcourt, Kenneth wrote:

...

The CUDA compiler (nvcc) is, in essence, a C compiler not C++.

This is incorrect.

...

You'll need to read the documentation but C++ code is compiled by a host compiler (g++, etc...) while nvcc is used to compile device (GPU) specific C-like code.

The device-specific compiler has a C++ front-end, complete with templates, specializations, etc. + NVIDIA-specific extensions. The complication is that code that targets both the "device" (GPU) and the "host" (CPU) needs to be compatible with at least 2 front-ends (NVCC and GCC/MSVC), but that's not unusual for Boost in general.

Oh, I guess things have changed a bit since I last used CUDA. Sorry about the bad info. -- Noel

I've never tested proto with nvcc, so I'm not surprised it doesn't work. I would need someone familiar with the quirks of this compiler (Troy?) to help get proto working with it. And there's no guarantee it's possible -- proto requires a fairly compliant compiler.

Today with Eric we took a look at what happened. The problem is twofold: 1/ firs, we need to run the config script so it geenrae a proper config/user.hpp file for NVCC 2/ the BOOST_PROTO_DECLTYPE_ macro: The original version is something like: BOOST_TYPEOF_NESTED_TYPEDEF_TPL(BOOST_PP_CAT(nested_and_hidden_, NESTED), EXPR) static int const sz = sizeof(boost::proto::detail::check_reference(EXPR)); struct NESTED : boost::mpl::if_c< 1==sz , typename BOOST_PP_CAT(nested_and_hidden_, NESTED)::type & , typename BOOST_PP_CAT(nested_and_hidden_, NESTED)::type > {}; # define BOOST_PROTO_DECLTYPE_(EXPR, TYPE) BOOST_PROTO_DECLTYPE_NESTED_TYPEDEF_TPL_(BOOST_PP_CAT(nested_, TYPE), (EXPR)) typedef typename BOOST_PP_CAT(nested_, TYPE)::type TYPE; # endif NVCC chokes on that by sayign that nested_type::type is used w/o template qualifiers. I can't make head or tail of thi message BUT. By changing it to: #define BOOST_PROTO_DECLTYPE_NESTED_TYPEDEF_TPL_(NESTED, EXPR) BOOST_TYPEOF_NESTED_TYPEDEF_TPL(BOOST_PP_CAT(hidden_,NESTED), EXPR) BOOST_STATIC_CONSTANT(int, sz = sizeof(boost::proto::detail::check_reference(EXPR))); typedef typename boost::mpl::if_c< 1==sz , typename BOOST_PP_CAT(hidden_, NESTED)::type & , typename BOOST_PP_CAT(hidden_, NESTED)::type >::type # define BOOST_PROTO_DECLTYPE_(EXPR, TYPE) BOOST_PROTO_DECLTYPE_NESTED_TYPEDEF_TPL_(BOOST_PP_CAT(nested_, TYPE), (EXPR)) TYPE; # endif ie removing the NESTED struct and computing the return of if_c directly into TYPE, NVCC happily compiles proto code in .cu file. Now, we can compile proto and fusion code in .cu for HOS function (ie code executed by the cpu). For making it works in DEVICE function (ie run by the GPU itself), it need a substantial amount of work that includes prefixing ALL proto, mpl, fusion etc function with __host__ __device__ so they can be called inside kernels. If anyone is interested, the user.hpp for nvcc CUDA is available on request. We ponder to do a global BOOST_GPU_ENABLED macro that actually evaluates as __host__ __device__ if compiled with nvccc and to nothing otherwise and do a local patch for that in w/e boost components it makes sense in. Considering the impact and the maybe restricted audience for that, not sure it's worth a global support from boost but it'll be here and available if needed.

Hi Joel and Eric, it is really awesome work that you both do. I am also interested in this user.hpp file. Can you forward a copy to me? That will be really nice. Best regards, Kim Joel Falcou schrieb:

If anyone is interested, the user.hpp for nvcc CUDA is available on request. We ponder to do a global BOOST_GPU_ENABLED macro that actually evaluates as __host__ __device__ if compiled with nvccc and to nothing otherwise and do a local patch for that in w/e boost components it makes sense in. Considering the impact and the maybe restricted audience for that, not sure it's worth a global support from boost but it'll be here and available if needed.

On Tue, Feb 9, 2010 at 12:11 PM, joel falcou wrote:

Kim Kuen Tang wrote:

...

Hi Joel and Eric, it is really awesome work that you both do. I am also interested in this user.hpp file. Can you forward a copy to me? That will be really nice.

Well, as there is like half a million people asking for that, I'll build up a page with the fiel and some directions. Which lead me to wonder: is there any volunteer to explore other parts of boost that need patching for GPU support ?

Sure, count me in. I was thinking about getting spirit/phoenix to work with nvcc. Manjunath