Adding vtkm::cont::RuntimeDeviceInformation.

The RuntimeDeviceInformation class allows developers to check if a given
device is supported on a machine at runtime. This allows developers to properly
check for CUDA support before running any worklets.

examples/multi_backend/MultiBackend.cxx

@@ -21,6 +21,8 @@
 
 #include <vtkm/Math.h>
 #include <vtkm/cont/ArrayHandle.h>
+#include <vtkm/cont/RuntimeDeviceInformation.h>
+
 #include <vtkm/worklet/DispatcherMapField.h>
 #include <vtkm/worklet/WorkletMapField.h>
 
@@ -104,13 +106,23 @@ struct CanRun<true>
                   vtkm::cont::ArrayHandle< vtkm::Vec< vtkm::UInt8, 4 > > outColors,
                   DeviceAdapterTag)
   {
-  typedef vtkm::worklet::DispatcherMapField<GenerateSurfaceWorklet,
-                                            DeviceAdapterTag> DispatcherType;
 
-  GenerateSurfaceWorklet worklet( 0.05f );
+  //even though we have support for this device adapter we haven't determined
-  DispatcherType(worklet).Invoke( inHandle,
+  //if we actually have run-time support. This is a significant issue with
-                                  outCoords,
+  //the CUDA backend
-                                  outColors);
+  vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtime;
+  const bool haveSupport = runtime.Exists();
+
+  if(haveSupport)
+    {
+    typedef vtkm::worklet::DispatcherMapField<GenerateSurfaceWorklet,
+                                              DeviceAdapterTag> DispatcherType;
+
+    GenerateSurfaceWorklet worklet( 0.05f );
+    DispatcherType(worklet).Invoke( inHandle,
+                                    outCoords,
+                                    outColors);
+    }
   }
 };
 
@@ -149,11 +161,9 @@ int main(int, char**)
   vtkm::cont::ArrayHandle< Uint8Vec4 > color;
 
   //Run the algorithm on all backends that we have compiled support for.
-  run_if_valid(in, out, color, SerialTag());
   run_if_valid(in, out, color, CudaTag());
   run_if_valid(in, out, color, TBBTag());
-
+  run_if_valid(in, out, color, SerialTag());
-
 }
 
 

vtkm/cont/DeviceAdapterAlgorithm.h

@@ -496,6 +496,33 @@ public:
   }
 };
 
+/// \brief Class providing a device-specific runtime support detector.
+///
+/// The class provide the actual implementation used by
+/// vtkm::cont::RuntimeDeviceInformation.
+///
+/// A default implementation is provided but device adapters which require
+/// physical hardware or other special runtime requirements should provide
+/// one (in conjunction with DeviceAdapterAlgorithm) where appropriate.
+///
+template<class DeviceAdapterTag>
+class DeviceAdapterRuntimeDetector
+{
+public:
+
+  /// Returns true if the given device adapter is supported on the current
+  /// machine.
+  ///
+  /// The default implementation is to return the value of
+  /// vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag>::Valid
+  ///
+  VTKM_CONT_EXPORT bool Exists() const
+  {
+    typedef vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag> DeviceAdapterTraits;
+    return DeviceAdapterTraits::Valid;
+  }
+};
+
 }
 } // namespace vtkm::cont
 

vtkm/cont/RuntimeDeviceInformation.h

@@ -0,0 +1,57 @@
+//============================================================================
+//  Copyright (c) Kitware, Inc.
+//  All rights reserved.
+//  See LICENSE.txt for details.
+//  This software is distributed WITHOUT ANY WARRANTY; without even
+//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+//  PURPOSE.  See the above copyright notice for more information.
+//
+//  Copyright 2014 Sandia Corporation.
+//  Copyright 2014 UT-Battelle, LLC.
+//  Copyright 2014 Los Alamos National Security.
+//
+//  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
+//  the U.S. Government retains certain rights in this software.
+//
+//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+//  Laboratory (LANL), the U.S. Government retains certain rights in
+//  this software.
+//============================================================================
+#ifndef vtk_m_cont_RuntimeDeviceInformation_h
+#define vtk_m_cont_RuntimeDeviceInformation_h
+
+#include <vtkm/cont/DeviceAdapter.h>
+
+namespace vtkm {
+namespace cont {
+
+/// A class that can be used to determine if a given device adapter
+/// is supported on the current machine at runtime. This is very important
+/// for device adapters that a physical hardware requirements such as a GPU
+/// or a Accelerator Card.
+///
+///
+template<class Device = VTKM_DEFAULT_DEVICE_ADAPTER_TAG>
+class RuntimeDeviceInformation
+{
+public:
+  VTKM_CONT_EXPORT
+  RuntimeDeviceInformation() : RuntimeImplementation() {  }
+
+  /// Returns true if the given device adapter is supported on the current
+  /// machine.
+  ///
+  VTKM_CONT_EXPORT
+  bool Exists() const
+  {
+    return this->RuntimeImplementation.Exists();
+  }
+
+private:
+  vtkm::cont::DeviceAdapterRuntimeDetector<Device> RuntimeImplementation;
+};
+
+}
+} // namespace vtkm::cont
+
+#endif //vtk_m_cont_RuntimeDeviceInformation_h

vtkm/cont/cuda/internal/DeviceAdapterAlgorithmCuda.h

@@ -26,11 +26,31 @@
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
 #include <vtkm/cont/ErrorControlInternal.h>
 
+#include <vtkm/Math.h>
+
 // Here are the actual implementation of the algorithms.
 #include <vtkm/cont/cuda/internal/DeviceAdapterAlgorithmThrust.h>
 
 #include <cuda.h>
 
+namespace vtkm {
+namespace cont {
+namespace cuda {
+namespace internal {
+
+static
+__global__
+void DetermineIfValidCudaDevice()
+{
+  //used only to see if we can launch kernels. It is possible to have a
+  //CUDA capable device, but still fail to have CUDA support.
+}
+
+}
+}
+}
+}
+
 namespace vtkm {
 namespace cont {
 
@@ -94,6 +114,80 @@ private:
   cudaEvent_t EndEvent;
 };
 
+/// \brief Class providing a CUDA runtime support detector.
+///
+/// The class provide the actual implementation used by
+/// vtkm::cont::RuntimeDeviceInformation for the CUDA backend.
+///
+/// We will verify at runtime that the machine has at least one CUDA
+/// capable device, and said device is from the 'fermi' (SM_20) generation
+/// or newer.
+///
+template<>
+class DeviceAdapterRuntimeDetector<vtkm::cont::DeviceAdapterTagCuda>
+{
+public:
+  VTKM_CONT_EXPORT DeviceAdapterRuntimeDetector():
+    NumberOfDevices(0),
+    HighestArchSupported(0)
+  {
+    static bool deviceQueryInit = false;
+    static int numDevices = 0;
+    static int archVersion = 0;
+
+    if(!deviceQueryInit)
+      {
+      deviceQueryInit = true;
+
+      //first query for the number of devices
+      cudaGetDeviceCount(&numDevices);
+
+      for (vtkm::Int32 i = 0; i < numDevices; i++)
+      {
+        cudaDeviceProp prop;
+        cudaGetDeviceProperties(&prop, i);
+        const vtkm::Int32 arch = (prop.major * 10) + prop.minor;
+        archVersion = vtkm::Max(arch, archVersion);
+      }
+
+      //Make sure we can actually launch a kernel. This could fail for any
+      //of the following reasons:
+      //
+      // 1. cudaErrorInsufficientDriver, caused by out of data drives
+      // 2. cudaErrorDevicesUnavailable, caused by another process locking the
+      //    device or somebody disabling cuda support on the device
+      // 3. cudaErrorNoKernelImageForDevice we built for a compute version
+      //    greater than the device we are running on
+      // Most likely others that I don't even know about
+      vtkm::cont::cuda::internal::DetermineIfValidCudaDevice <<<1,1>>> ();
+      if(cudaSuccess != cudaGetLastError())
+        {
+        numDevices = 0;
+        archVersion = 0;
+        }
+      }
+
+    this->NumberOfDevices = numDevices;
+    this->HighestArchSupported = archVersion;
+  }
+
+  /// Returns true if the given device adapter is supported on the current
+  /// machine.
+  ///
+  /// Only returns true if we have at-least one CUDA capable device of SM_20 or
+  /// greater ( fermi ).
+  ///
+  VTKM_CONT_EXPORT bool Exists() const
+  {
+    //
+    return this->NumberOfDevices > 0 && this->HighestArchSupported >= 20;
+  }
+
+private:
+  vtkm::Int32 NumberOfDevices;
+  vtkm::Int32 HighestArchSupported;
+};
+
 }
 } // namespace vtkm::cont
 

vtkm/cont/testing/CMakeLists.txt

@@ -53,6 +53,7 @@ set(unit_tests
   UnitTestDeviceAdapterSerial.cxx
   UnitTestDynamicArrayHandle.cxx
   UnitTestDynamicCellSet.cxx
+  UnitTestRuntimeDeviceInformation.cxx
   UnitTestStorageBasic.cxx
   UnitTestStorageImplicit.cxx
   UnitTestStorageListTag.cxx

vtkm/cont/testing/TestingDeviceAdapter.h

@@ -30,6 +30,7 @@
 #include <vtkm/cont/ArrayPortalToIterators.h>
 #include <vtkm/cont/ErrorControlBadAllocation.h>
 #include <vtkm/cont/ErrorExecution.h>
+#include <vtkm/cont/RuntimeDeviceInformation.h>
 #include <vtkm/cont/StorageBasic.h>
 #include <vtkm/cont/Timer.h>
 #include <vtkm/cont/DeviceAdapterAlgorithm.h>
@@ -415,6 +416,18 @@ private:
                      "Timer counted too far or system really busy.");
   }
 
+  VTKM_CONT_EXPORT
+  static void TestRuntime()
+  {
+    std::cout << "-------------------------------------------" << std::endl;
+    std::cout << "Testing RuntimeDeviceInformation" << std::endl;
+
+    vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtime;
+    const bool valid_runtime = runtime.Exists();
+
+    VTKM_TEST_ASSERT(valid_runtime, "runtime detection failed for device");
+  }
+
   static VTKM_CONT_EXPORT void TestAlgorithmSchedule()
   {
     std::cout << "-------------------------------------------" << std::endl;
@@ -1550,6 +1563,7 @@ private:
       TestArrayManagerExecution();
       TestOutOfMemory();
       TestTimer();
+      TestRuntime();
 
       TestAlgorithmSchedule();
       TestErrorExecution();

vtkm/cont/testing/UnitTestRuntimeDeviceInformation.cxx

@@ -0,0 +1,86 @@
+//============================================================================
+//  Copyright (c) Kitware, Inc.
+//  All rights reserved.
+//  See LICENSE.txt for details.
+//  This software is distributed WITHOUT ANY WARRANTY; without even
+//  the implied warranty of MERCHANTABILITY or FITNESS FOR A PARTICULAR
+//  PURPOSE.  See the above copyright notice for more information.
+//
+//  Copyright 2014 Sandia Corporation.
+//  Copyright 2014 UT-Battelle, LLC.
+//  Copyright 2014 Los Alamos National Security.
+//
+//  Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation,
+//  the U.S. Government retains certain rights in this software.
+//
+//  Under the terms of Contract DE-AC52-06NA25396 with Los Alamos National
+//  Laboratory (LANL), the U.S. Government retains certain rights in
+//  this software.
+//============================================================================
+
+#include <vtkm/cont/RuntimeDeviceInformation.h>
+
+//include all backends
+#include <vtkm/cont/cuda/DeviceAdapterCuda.h>
+#include <vtkm/cont/tbb/DeviceAdapterTBB.h>
+#include <vtkm/cont/DeviceAdapterSerial.h>
+
+#include <vtkm/cont/testing/Testing.h>
+
+namespace {
+
+template<bool> struct DoesExist;
+
+template<typename DeviceAdapterTag>
+void detect_if_exists(DeviceAdapterTag tag)
+{
+  typedef vtkm::cont::DeviceAdapterTraits<DeviceAdapterTag> DeviceAdapterTraits;
+  DoesExist<DeviceAdapterTraits::Valid>::Exist(tag);
+}
+
+template<>
+struct DoesExist<false>
+{
+  template<typename DeviceAdapterTag>
+  static void Exist(DeviceAdapterTag)
+  {
+    //runtime information for this device should return false
+    vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtime;
+    VTKM_TEST_ASSERT(runtime.Exists() == false,
+      "A backend with zero compile time support, can't have runtime support");
+  }
+};
+
+template<>
+struct DoesExist<true>
+{
+  template<typename DeviceAdapterTag>
+  static void Exist(DeviceAdapterTag)
+  {
+    //runtime information for this device should return true
+    vtkm::cont::RuntimeDeviceInformation<DeviceAdapterTag> runtime;
+    VTKM_TEST_ASSERT(runtime.Exists() == true,
+      "A backend with compile time support, should have runtime support");
+  }
+};
+
+void Detection()
+{
+  typedef ::vtkm::cont::DeviceAdapterTagSerial SerialTag;
+  typedef ::vtkm::cont::DeviceAdapterTagTBB TBBTag;
+  typedef ::vtkm::cont::DeviceAdapterTagCuda CudaTag;
+
+  //Verify that for each device adapter we compile code for, that it
+  //has valid runtime support.
+  detect_if_exists(CudaTag());
+  detect_if_exists(TBBTag());
+  detect_if_exists(SerialTag());
+
+}
+
+} // anonymous namespace
+
+int UnitTestRuntimeDeviceInformation(int, char *[])
+{
+  return vtkm::cont::testing::Testing::Run(Detection);
+}

vtkm/worklet/testing/UnitTestWorkletMapFieldWholeArray.cxx

@@ -114,7 +114,7 @@ struct DoTestWorklet
 
 void TestWorkletMapFieldExecArg()
 {
-  typedef vtkm::cont::internal::DeviceAdapterTraits<
+  typedef vtkm::cont::DeviceAdapterTraits<
                     VTKM_DEFAULT_DEVICE_ADAPTER_TAG> DeviceAdapterTraits;
   std::cout << "Testing Worklet with WholeArray on device adapter: "
             << DeviceAdapterTraits::GetId() << std::endl;