Merge pull request #24938 from fwyzard/skip_CUDA_tests_104x

Skip tests if no CUDA devices are present
cms-sw · Oct 24, 2018 · 3bad4f9 · 3bad4f9
2 parents 7b827de + 16a49e9
commit 3bad4f9
Show file tree

Hide file tree

Showing 2 changed files with 226 additions and 224 deletions.
diff --git a/DataFormats/Math/test/cudaAtan2Test.cu b/DataFormats/Math/test/cudaAtan2Test.cu
@@ -18,28 +18,25 @@ echo $f; $f
 end
 */
 
-#include "DataFormats/Math/interface/approx_atan2.h"
-
-#include "cuda/api_wrappers.h"
-
-#include <iostream>
-#include <iomanip>
-#include <memory>
 #include <algorithm>
+#include <cassert>
 #include <chrono>
+#include <iomanip>
+#include <iostream>
+#include <memory>
 #include <random>
 
-#include <cassert>
+#include "cuda/api_wrappers.h"
+
+#include "DataFormats/Math/interface/approx_atan2.h"
 
 constexpr float xmin=-100.001;  // avoid 0
 constexpr float incr = 0.04;
 constexpr int Nsteps = 2.*std::abs(xmin)/incr;
 
-
 template<int DEGREE>
-__global__ void diffAtan(int * diffs) {
-
-
+__global__ void diffAtan(int * diffs)
+{
   auto mdiff = &diffs[0];
   auto idiff = &diffs[1];
   auto sdiff = &diffs[2];
@@ -48,7 +45,7 @@ __global__ void diffAtan(int * diffs) {
   int j = blockDim.y * blockIdx.y + threadIdx.y;
 
   auto x = xmin +incr*i;
-  auto y = xmin	+incr*j;
+  auto y = xmin +incr*j;
 
   auto approx = unsafe_atan2f<DEGREE>(y,x);
   auto iapprox = unsafe_atan2i<DEGREE>(y,x);
@@ -59,75 +56,76 @@ __global__ void diffAtan(int * diffs) {
   atomicMax(idiff, std::abs(phi2int(std)-iapprox));
   short dd = std::abs(phi2short(std)-sapprox);
   atomicMax(sdiff,int(dd));
-
 }
 
 template<int DEGREE>
 void go() {
   auto start = std::chrono::high_resolution_clock::now();
   auto delta = start - start;
 
-	if (cuda::device::count() == 0) {
-		std::cerr << "No CUDA devices on this system" << "\n";
-		exit(EXIT_FAILURE);
-	}
+  auto current_device = cuda::device::current::get(); 
 
-        auto current_device = cuda::device::current::get(); 
-        // atan2
-        delta -= (std::chrono::high_resolution_clock::now()-start);
+  // atan2
+  delta -= (std::chrono::high_resolution_clock::now()-start);
 
-        auto diff_d = cuda::memory::device::make_unique<int[]>(current_device,3);
+  auto diff_d = cuda::memory::device::make_unique<int[]>(current_device,3);
 
-        int diffs[3];
-        cuda::memory::device::zero(diff_d.get(),3*4);
-
-        // Launch the diff CUDA Kernel
-        dim3 threadsPerBlock(32,32,1);
-        dim3 blocksPerGrid ( (Nsteps + threadsPerBlock.x - 1) / threadsPerBlock.x,
-                             (Nsteps + threadsPerBlock.y - 1) / threadsPerBlock.y,
-                            1
-                               );
-        std::cout
-                << "CUDA kernel 'diff' launch with " << blocksPerGrid.x
-                << " blocks of " << threadsPerBlock.y << " threads\n";
-
-        cuda::launch(
-                diffAtan<DEGREE>,
-                { blocksPerGrid, threadsPerBlock },
-                diff_d.get() );
-
-        cuda::memory::copy(diffs, diff_d.get(), 3*4);
-        delta += (std::chrono::high_resolution_clock::now()-start);
+  int diffs[3];
+  cuda::memory::device::zero(diff_d.get(),3*4);
+
+  // Launch the diff CUDA Kernel
+  dim3 threadsPerBlock(32,32,1);
+  dim3 blocksPerGrid((Nsteps + threadsPerBlock.x - 1) / threadsPerBlock.x,
+                     (Nsteps + threadsPerBlock.y - 1) / threadsPerBlock.y,
+                     1);
+  std::cout
+    << "CUDA kernel 'diff' launch with " << blocksPerGrid.x
+    << " blocks of " << threadsPerBlock.y << " threads\n";
+
+  cuda::launch(
+    diffAtan<DEGREE>,
+    { blocksPerGrid, threadsPerBlock },
+    diff_d.get());
+
+  cuda::memory::copy(diffs, diff_d.get(), 3*4);
+  delta += (std::chrono::high_resolution_clock::now()-start);
 
-        float mdiff = diffs[0]*1.e-7;
-        int idiff = diffs[1];
-        int sdiff = diffs[2];
-
-        std::cout << "for degree " << DEGREE << " max diff is " << mdiff 
-                             << ' ' << idiff << ' ' << int2phi(idiff) 
-                             << ' ' << sdiff << ' ' << short2phi(sdiff) <<  std::endl;
-
-
-        std::cout <<"cuda computation took "
-              << std::chrono::duration_cast<std::chrono::milliseconds>(delta).count()
-              << " ms" << std::endl;
-
-	return;
+  float mdiff = diffs[0]*1.e-7;
+  int idiff = diffs[1];
+  int sdiff = diffs[2];
+
+  std::cout << "for degree " << DEGREE << " max diff is " << mdiff 
+            << ' ' << idiff << ' ' << int2phi(idiff) 
+            << ' ' << sdiff << ' ' << short2phi(sdiff) <<  std::endl;
+  std::cout << "cuda computation took "
+            << std::chrono::duration_cast<std::chrono::milliseconds>(delta).count()
+            << " ms" << std::endl;
 }
 
-
 int main() {
-try{
-  go<3>();
-  go<5>();
-  go<7>();
-  go<9>();
-} catch(cuda::runtime_error ex) {
-   std::cout << "cuda error " << ex.what() << std::endl;
-} catch(...) {
-   std::cout <<	"a non cuda error" << std::endl;
-}
-
-
-  return 0;
+  int count = 0;
+  auto status = cudaGetDeviceCount(& count);
+  if (status != cudaSuccess) {
+    std::cerr << "Failed to initialise the CUDA runtime, the test will be skipped." << "\n";
+    exit(EXIT_SUCCESS);
+  }
+  if (count == 0) {
+    std::cerr << "No CUDA devices on this system, the test will be skipped." << "\n";
+    exit(EXIT_SUCCESS);
+  }
+
+  try {
+    go<3>();
+    go<5>();
+    go<7>();
+    go<9>();
+  } catch(cuda::runtime_error ex) {
+    std::cerr << "CUDA error: " << ex.what() << std::endl;
+    exit(EXIT_FAILURE);
+  } catch(...) {
+    std::cerr << "A non-CUDA error occurred" << std::endl;
+    exit(EXIT_FAILURE);
+  }
+
+  return EXIT_SUCCESS;
 }