stateless_resource_allocator generates warnings about using __host__ function from __host__ __device__

[robertmaynard]

When building VTK-m ( https://gitlab.kitware.com/vtk/vtk-m/ ) with CUDA 10.1 we are seeing thrust generate warnings the following warning:

... include/thrust/mr/allocator.h(106): warning #2912-D: calling a __host__ function("thrust::mr::stateless_resource_allocator<long long,  ::thrust::system::cuda::detail::cuda_memory_resource<& ::cudaMalloc, & ::cudaFree,  ::thrust::cuda_cub::pointer<void> > > ::stateless_resource_allocator") from a __host__ __device__ function("thrust::mr::stateless_resource_allocator<long long,  ::thrust::system::cuda::detail::cuda_memory_resource<& ::cudaMalloc, & ::cudaFree,  ::thrust::cuda_cub::pointer<void> > > ::stateless_resource_allocator [subobject]") is not allowed

[robertmaynard]

VTK-m builds with separable compilation enabled.

[griwes]

Yeah, this is a known issue. We need to figure out what's the correct qualifications on all of those, and there's probably several places where we still need to add the pragmas disabling the warning, but I didn't want to do this for this release, because I couldn't yet wrap my head around how to make sure it doesn't hide anything bad.

By "builds", do you mean "builds successfully"?

[robertmaynard]

I just meant in generally we always build with the -rdc enabled ( In case that was required to construct a minimal test case ).

[robertmaynard]

So some follow up @griwes

The warning is being generated by the thrust::system::cuda::allocator, which can be reproduced with the following snippet:

#include <thrust/system/cuda/memory_resource.h>
#include <thrust/system/cuda/memory.h>

namespace {

template <typename T>
struct allocator
    : thrust::mr::stateless_resource_allocator<
        T,
        ::thrust::system::cuda::memory_resource
    >
{
private:
    typedef thrust::mr::stateless_resource_allocator<
        T,
        ::thrust::system::cuda::memory_resource
    > base;

public:
  template <typename U>
  struct rebind
  {
    typedef allocator<U> other;
  };

  __host__ __device__
  inline allocator() {}

  __host__ __device__
 inline allocator(const allocator & other) : base(other) {}

  template <typename U>
  __host__ __device__
  inline allocator(const allocator<U> & other) : base(other) {}

  __host__ __device__
  inline ~allocator() {}
};    // struct allocator


class CMR
{
 allocator<float>  m_allocf;
};

__host__ __device__  void compute(int& input)
{
  CMR t{};
}

The simple solution of marking the relevant functions as host device is currently being attempted by vtkm. We are currently tracking down an infrequent runtime segfault ( 30% of the time ) with VTK-m ( backtrace below ) that for some tests is resolved by not applying our patches for this warning. I am still working on developing a minimal test case as currently I can only reproduce it by building: https://gitlab.kitware.com/vtk/vtk/merge_requests/5589

crashed backtrace

1052: (   0.002s) [main thread     ]             loguru.cpp:587   INFO| arguments: TestVTKMCleanGrid -D /home/robert/Work/vtk/build/ExternalData/Testing -T /home/robert/Work/vtk/build/Testing/Temporary
1052: (   0.002s) [main thread     ]             loguru.cpp:590   INFO| Current dir: /home/robert/Work/vtk/build/Accelerators/Vtkm/Testing/Cxx
1052: (   0.002s) [main thread     ]             loguru.cpp:592   INFO| stderr verbosity: 0
1052: (   0.002s) [main thread     ]             loguru.cpp:593   INFO| -----------------------------------
1052: 2019-05-29 15:29:14.051 (   0.002s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'Serial' to 1
1052: 2019-05-29 15:29:14.057 (   0.007s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'Cuda' to 1
1052: 2019-05-29 15:29:14.057 (   0.007s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'TBB' to 0
1052: 2019-05-29 15:29:14.057 (   0.007s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'OpenMP' to 0
1052: 2019-05-29 15:29:14.057 (   0.007s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'InvalidDeviceId' to 0
1052: 2019-05-29 15:29:14.057 (   0.008s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'InvalidDeviceId' to 0
1052: 2019-05-29 15:29:14.057 (   0.008s) [        CFDDD0C0]RuntimeDeviceTracker.cx:93    INFO| Reset device 'InvalidDeviceId' to 0
1052: 2019-05-29 15:29:14.057 (   0.008s) [        CFDDD0C0]       CudaAllocator.cu:307   INFO| CudaAllocator hardware supports managed memory
1052: 2019-05-29 15:29:14.994 (   0.945s) [        CFDDD0C0]                Error.h:65    WARN| Exception raised: scan failed on 2nd step: cudaErrorInvalidValue: invalid argument
1052: 45      0x5620ec4a501a /home/robert/Work/vtk/build/bin/vtkAcceleratorsVTKmCxxTests(+0x13f01a) [0x5620ec4a501a]
1052: 44      0x7f11cf828b97 __libc_start_main + 231
1052: 43      0x5620ec4a59b8 /home/robert/Work/vtk/build/bin/vtkAcceleratorsVTKmCxxTests(+0x13f9b8) [0x5620ec4a59b8]
1052: 42      0x5620ec4a6223 /home/robert/Work/vtk/build/bin/vtkAcceleratorsVTKmCxxTests(+0x140223) [0x5620ec4a6223]

Working DEBUG_SYNC log:

1052: Invoking transform::agent<<<1, 256, 0, 1>>>(), 3 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::init_agent<<<1, 128, 0, 2>>>(), 11 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::scan_agent<<<11, 96, 2128, 2>>>(), 8000 items total, 8 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 4 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking block_sort_agent<<<3, 256, 34832, 2>>>(), 9261 items total, 17 items per thread, 1 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking partition_agent<<<1, 256, 0, 2>>>(), 4 items total, 1 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking partition_agent<<<1, 256, 0, 2>>>(), 4 items total, 1 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking reduce_by_key::init_agent<<<1, 128, 0, 2>>>(), 5 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking reduce_by_keys::reduce_by_key_agent<<<5, 256, 18448, 2>>>(), 9261 items total, 9 items per thread, 4 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 1>>>(), 3 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::init_agent<<<1, 128, 0, 2>>>(), 13 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::scan_agent<<<13, 96, 2128, 2>>>(), 9261 items total, 8 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking block_sort_agent<<<5, 256, 16912, 2>>>(), 9261 items total, 8 items per thread, 1 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking partition_agent<<<1, 256, 0, 2>>>(), 6 items total, 1 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking merge_agent<<<5, 256, 16912, 2>>>(), 9261 items total, 8 items per thread, 2 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking partition_agent<<<1, 256, 0, 2>>>(), 6 items total, 1 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking merge_agent<<<5, 256, 16912, 2>>>(), 9261 items total, 8 items per thread, 2 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking partition_agent<<<1, 256, 0, 2>>>(), 6 items total, 1 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking merge_agent<<<5, 256, 16912, 2>>>(), 9261 items total, 8 items per thread, 2 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking reduce_by_key::init_agent<<<1, 128, 0, 2>>>(), 5 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking reduce_by_keys::reduce_by_key_agent<<<5, 256, 36880, 2>>>(), 9261 items total, 9 items per thread, 3 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 1>>>(), 3 items total, 2 it(   0.988s) [main thread     ]             loguru.cpp:467   INFO| atexit
1052: ems per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::init_agent<<<1, 128, 0, 2>>>(), 13 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::scan_agent<<<13, 96, 2128, 2>>>(), 9261 items total, 8 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<19, 256, 0, 2>>>(), 9261 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking copy_if::init_agent<<<1, 128, 0, 2>>>(), 16 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking copy_if::partition_agent<<<16, 128, 4096, 2>>>(), 8000 items total, 4 items per thread, 10 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 1>>>(), 3 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::init_agent<<<1, 128, 0, 2>>>(), 11 items total, 1 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking scan::scan_agent<<<11, 96, 2128, 2>>>(), 8000 items total, 8 items per thread, 16 SM occupancy, 0 vshmem size, 610 ptx_version 
1052: Invoking transform::agent<<<1, 256, 0, 2>>>(), 1 items total, 2 items per thread, 8 SM occupancy, 0 vshmem size, 610 ptx_version 

[robertmaynard]

I have been able to construct small reproducer that shows the warning is generated by at-least scan_exclusive:
https://github.com/robertmaynard/Sandbox/blob/master/ThrustFixStatelessResourceWarnings/middle.cu

[alliepiper]

Fixed by d43f285.