Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-16 17:22:55 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge topic 'allow-disabling/enabling-cuda-managed-memory'

Browse Source 
e34301eca Allow disabling/enabling of CUDA managed memory via an env variable

Acked-by: Kitware Robot <kwrobot@kitware.com>
Acked-by: Robert Maynard <robert.maynard@kitware.com>
Merge-request: !1359
This commit is contained in:
Haocheng LIU 2018-08-17 17:13:53 +00:00 committed by Kitware Robot
commit 7d22132253
2 changed files with 34 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
docs/changelog/allow-disabling-enabling-cuda-managed-memory.md Normal file
View File

@ -0,0 +1,6 @@
Allow disabling/enabling of CUDA managed memory through a environment variable
By setting the environment variable "VTKM_MANAGEDMEMO_DISABLED" to be 1,
users are able to disable CUDA managed memory even though the hardware is capable
of doing so.

44
vtkm/cont/cuda/internal/CudaAllocator.cu
View File

@ -18,8 +18,11 @@
// this software.
//============================================================================
#include <cstdlib>
#include <mutex>
#include <vtkm/cont/cuda/ErrorCuda.h>
#include <vtkm/cont/cuda/internal/CudaAllocator.h>
#define NO_VTKM_MANAGED_MEMORY "NO_VTKM_MANAGED_MEMORY"
#include <mutex>
#include <vector>
@ -32,12 +35,13 @@ VTKM_THIRDPARTY_POST_INCLUDE
namespace
{
#if CUDART_VERSION >= 8000
// Has CudaAllocator::Initialize been called?
static bool IsInitialized = false;
// Has CudaAllocator::Initialize been called by any thread?
static std::once_flag IsInitialized;
#endif
// True if all devices support concurrent pagable managed memory.
static bool ManagedMemorySupported = false;
// True if concurrent pagable managed memory is not disabled by user via a system
// environment variable and all devices support it.
static bool ManagedMemoryEnabled = false;
// Avoid overhead of cudaMemAdvise and cudaMemPrefetchAsync for small buffers.
// This value should be > 0 or else these functions will error out.
@ -56,7 +60,7 @@ namespace internal
bool CudaAllocator::UsingManagedMemory()
{
CudaAllocator::Initialize();
return ManagedMemorySupported;
return ManagedMemoryEnabled;
}
bool CudaAllocator::IsDevicePointer(const void* ptr)
@ -82,7 +86,7 @@ bool CudaAllocator::IsDevicePointer(const void* ptr)
bool CudaAllocator::IsManagedPointer(const void* ptr)
{
if (!ptr || !ManagedMemorySupported)
if (!ptr || !ManagedMemoryEnabled)
{
return false;
}
@ -111,7 +115,7 @@ void* CudaAllocator::Allocate(std::size_t numBytes)
}
void* ptr = nullptr;
if (ManagedMemorySupported)
if (ManagedMemoryEnabled)
{
VTKM_CUDA_CALL(cudaMallocManaged(&ptr, numBytes));
}
@ -222,15 +226,13 @@ void CudaAllocator::PrepareForInPlace(const void* ptr, std::size_t numBytes)
void CudaAllocator::Initialize()
{
#if CUDART_VERSION >= 8000
if (!IsInitialized)
{
std::call_once(IsInitialized, []() {
bool managedMemorySupported = true;
int numDevices;
VTKM_CUDA_CALL(cudaGetDeviceCount(&numDevices));
if (numDevices == 0)
{
ManagedMemorySupported = false;
IsInitialized = true;
return;
}
@ -245,11 +247,21 @@ void CudaAllocator::Initialize()
managed = managed && prop.concurrentManagedAccess;
}
ManagedMemorySupported = managed;
IsInitialized = true;
}
#else
ManagedMemorySupported = false;
managedMemorySupported = managed;
// Check if users want to disable managed memory
#pragma warning(push)
// getenv is not thread safe on windows but since it's inside a call_once block so
// it's fine to suppress the warning here.
#pragma warning(disable : 4996)
const char* buf = std::getenv(NO_VTKM_MANAGED_MEMORY);
#pragma warning(pop)
if (buf != nullptr)
{
ManagedMemoryEnabled = (std::string(buf) != "1");
}
ManagedMemoryEnabled = ManagedMemoryEnabled && managedMemorySupported;
});
#endif
}
}