Separate object on host and device, no more global objects

[slizzered]

There is no longer a global __device__ object, which simplifies reasoning about the behavior of mallocMC. Instead, there is an object explicitly created on the host, which internally creates the device allocator and holds a pointer to it. There are several side-effects:

• fixes calling a device object from host? #113
• using multiple allocators will be straightforward and no longer hidden behind MACRO-magic
• the whole macro magic class mallocMC_overwrites.hpp was removed
• it is no longer possible to directly replace standard malloc with mallocMC. Instead, you must pass the pointer to the device allocator into the kernel and use it:
• previously:

using namespace mallocMC;
using MallocMC_Type = ... // alias definition with all the policies
MAMC_SET_ALLOCATOR_TYPE(MallocMC_Type) // ugly MACRO-magic

__global__
void
exampleOLD( )
{
    int* x = malloc( sizeof( int ) * 100 );
}

int
main( )
{
    initHeap( 1024 );
    exampleOLD<<< ... >>>( );
    finalizeHeap();
}

• now:

using namespace mallocMC;
using MallocMC_Type = ... // alias definition with all the policies

__global__
void
exampleNEW(
    MallocMC_Type::AllocatorHandle mMC
)
{
    int* x = mMC.malloc( sizeof( int ) * 100 );
}

int
main( )
{
    MallocMC_Type mMC( 1024 );
    exampleNEW<<< ... >>>( mMC );
    mMC.finalizeHeap();
}

• as you can see, the mallocMC::initHeap() function (actually, it was a macro in client code) is no longer needed. Instead, it uses a constructor that takes the size directly.
• also, the mallocMC::finalizeHeap() macro is no longer used. Instead, it is encouraged to use the finalizeHeap() member function directly.

There remain several small problems, that need to be addressed (will create issues):

• the finalizeHeap() functionality could be done through the destructor of the host object. This will probably be postponed until C++11, where the host object is forbidden to be copied (only moving will be allowed).
• inside the constructor and finalizeHeap(), the hostclass currently uses CUDA functions directly to create/destroy the device object. This could either be integrated into the ReservePool-Policy or a new policy could be established to encapsulate this behaviour. Other suggestions welcome!

[psychocoderHPC]

Is it possible to use a method the get the device object instead to access a member direct?

[psychocoderHPC]

The device allocator should not be pointer, an object (handle) which can contain a pointer should be better.

[slizzered]

So, would you prefer a getter method (on the host) that returns an object like the following? (should be passed to the kernel as a pass-by-value object)

template <typename T_DevAllocator>
struct AllocatorHandle
{
    T_DevAllocator* devAllocator;

    MAMC_ACCELERATOR
    void*
    malloc(
        size_t size
    )
    {
        return devAllocator -> malloc( size );
    }

    MAMC_ACCELERATOR
    void
    free(
        void* p
    )
    {
        devAllocator -> free( p );
    }

};

[psychocoderHPC]

Yes this looks more save than use a pointer. Than the handle is something like an interface for the allocator.

[ax3l]

@Flamefire you can also have a look if you want :)

[ax3l]

@slizzered thank you for the fix and refactoring!

when this is reviewed,we should rather release a 2.3.0 and skip the patch level 2.2.1 because of the changed API, I guess.

[Flamefire]

Why not make this functions non-static? The instance could save the heap in initHeap and reset it in finalizeHeap
Then further calls do not need to pass the heap pointer

[Flamefire]

And actually T_Obj& should be Scatter_Impl&. It shouldn't be possible to pass another type in here that just happens to have a similar interface...

[Flamefire]

It seems that the policies methods, that are used from the hostclass_host are all static. This makes it necessary that the device object is passed to all this static functions which would be a point against them.
But if all used functions are static, the class should NOT inherit from the policies! It adds false properties (invalid, never set and must never be used) to the instance!

[Flamefire]

And last: I'd also provide an implicit conversion from the Allocator to the handle so one could pass the Allocator into a kernel (which causes the conversion and saves the allocator.getAllocatorHandle call)

[psychocoderHPC]

I will test soon

[psychocoderHPC]

After long long time I will merge it in and add the new changes to PIConGPU asap

[ax3l]

@psychocoderHPC can you please apply the latest mallocMC, including this feature to PIConGPU? It stops parallel/separable compilation.