boostorg / boostorg/fusion

Overzealously defining __device__ functions when compiling with nvcc

Open
#209 3 comments 0 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
50
Forks
68
PR merge metrics
No merged PRs in 30d

Description

I'm not entirely sure if this issue belongs here or on the boost umbrella project. The actual error I'm facing could be argued to originate either in boost::fusion or in boost::accumulators, but the overall strategy boost employs to support cuda is brittle to begin with. The overarching problem is that portions of boost cannot be used in `.cu` files even if they are not going to be used in actual device code. This is entirely due to how the sections of boost that support cuda attempt to do so, not due to any intrinsic limitations of host side code in a `.cu` file.

The crux my particular issue is that boost fusion seems to mark many (or all?) functions with `BOOST_FUSION_GPU_ENABLED`, which gets it's value from `BOOST_GPU_ENABLED`, which has the value of `__host__ __device__` when compiling with nvcc (and empty otherwise). This is problematic because this library is heavily templated, and you cannot call a `__host__` only function from a `__host__ __device__` function. Boost fusion calls functions it has no control over (injected via templates), and if any of those functions are not decorated with `__host__ __device__` then there will be problems. Even if the end user code is *only* going to invoke the host side call paths, the mere attempt to define the ill-formed device side call paths causes errors.

A minimal example is included below. It's a trivial cuda program that doesn't even call any GPU code. The mere fact that if you compile it with nvcc, boost::fusion will try and define device functions that call host-only boost::accumulator code. This makes it impossible to use portions of boost code in .cu files, even in portions of the code that will only ever be run on the host.

I can save the following code as "main.cu", and compile with `nvcc` version 10.0.117 on Ubuntu 18.04. The result is a lot of compilation errors about "identifier is undefined in device code". If hack things and redefine `BOOST_GPU_ENABLED` to be just `__host__` then there are no compilation issues.
```
#include
#include

using features =
boost::accumulators::features<
boost::accumulators::tag::count,
boost::accumulators::tag::mean,
boost::accumulators::tag::variance
>;

int main() {
boost::accumulators::accumulator_set accum;
return 0;
}
```

Contributor guide

No contributing guide indexed for this repository

Research direction

Reproduce the failure from the provided main.cu example using the listed Boost.Accumulators headers and nvcc, then inspect Boost.Fusion's BOOST_FUSION_GPU_ENABLED and BOOST_GPU_ENABLED definitions. Done means the host-only accumulator example compiles without device-code errors when compiled as a .cu file, without relying on the macro redefinition workaround.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
compilers
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.