]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Add tests for CUDAWrappers::Vector.
authorBruno Turcksin <bruno.turcksin@gmail.com>
Mon, 22 Aug 2016 19:02:16 +0000 (15:02 -0400)
committerBruno Turcksin <bruno.turcksin@gmail.com>
Fri, 26 Aug 2016 20:45:33 +0000 (16:45 -0400)
source/lac/cuda_vector.cu
tests/cuda/CMakeLists.txt [new file with mode: 0644]
tests/cuda/cuda_vector_01.cc [new file with mode: 0644]
tests/cuda/cuda_vector_01.output [new file with mode: 0644]
tests/cuda/cuda_vector_02.cc [new file with mode: 0644]
tests/cuda/cuda_vector_02.output [new file with mode: 0644]
tests/trilinos/epetra_vector_02.cc

index b954c3f7ab9df6a80522906122e5c48795c320c2..caf7f6b44ef9b3d9004db76bb5998d16864bfee5 100644 (file)
@@ -206,6 +206,13 @@ namespace LinearAlgebra
                                                               blockIdx.x*(blockDim.x*CHUNK_SIZE);
         const typename Vector<Number>::size_type local_idx = threadIdx.x;
 
+        if (global_idx<N)
+          result_buffer[local_idx] = v[global_idx];
+        else
+          result_buffer[local_idx] = Operation::null_value();
+
+        __syncthreads();
+
         reduce<Number,Operation> (result, result_buffer, local_idx, global_idx, N);
       }
 
@@ -422,6 +429,11 @@ namespace LinearAlgebra
                 res_buf[local_idx] += v1[idx]*v3[idx];
               }
           }
+
+        __syncthreads();
+
+        reduce<Number, DotProduct<Number>> (res, res_buf, local_idx,
+                                            global_idx, N);
       }
     }
 
diff --git a/tests/cuda/CMakeLists.txt b/tests/cuda/CMakeLists.txt
new file mode 100644 (file)
index 0000000..cc91456
--- /dev/null
@@ -0,0 +1,6 @@
+CMAKE_MINIMUM_REQUIRED(VERSION 2.8.9)
+INCLUDE(../setup_testsubproject.cmake)
+PROJECT(testsuite CXX)
+IF(DEAL_II_WITH_CUDA)
+  DEAL_II_PICKUP_TESTS()
+ENDIF()
diff --git a/tests/cuda/cuda_vector_01.cc b/tests/cuda/cuda_vector_01.cc
new file mode 100644 (file)
index 0000000..8abae2c
--- /dev/null
@@ -0,0 +1,112 @@
+// ---------------------------------------------------------------------
+//
+// Copyright (C) 2015 by the deal.II authors
+//
+// This file is part of the deal.II library.
+//
+// The deal.II library is free software; you can use it, redistribute
+// it, and/or modify it under the terms of the GNU Lesser General
+// Public License as published by the Free Software Foundation; either
+// version 2.1 of the License, or (at your option) any later version.
+// The full text of the license can be found in the file LICENSE at
+// the top level of the deal.II distribution.
+//
+// ---------------------------------------------------------------------
+
+
+// Check LinearAlgebra::CUDAWrappers::Vector assignement and import
+
+#include "../tests.h"
+#include <deal.II/base/utilities.h>
+#include <deal.II/lac/cuda_vector.h>
+#include <deal.II/lac/read_write_vector.h>
+#include <fstream>
+#include <iostream>
+
+void test()
+{
+  const unsigned int size = 100;
+  LinearAlgebra::CUDAWrappers::Vector<double> a;
+  LinearAlgebra::CUDAWrappers::Vector<double> b(size);
+  LinearAlgebra::CUDAWrappers::Vector<double> c(b);
+
+  AssertThrow(a.size()==0, ExcMessage("Vector has the wrong size."));
+  AssertThrow(b.size()==size, ExcMessage("Vector has the wrong size."));
+  AssertThrow(c.size()==size, ExcMessage("Vector has the wrong size."));
+
+  a.reinit(size);
+  AssertThrow(a.size()==size, ExcMessage("Vector has the wrong size."));
+
+
+  LinearAlgebra::ReadWriteVector<double> read_write_1(size);
+  LinearAlgebra::ReadWriteVector<double> read_write_2(size);
+  LinearAlgebra::ReadWriteVector<double> read_write_3(size);
+  for (unsigned int i=0; i<size; ++i)
+    {
+      read_write_1[i] = i;
+      read_write_2[i] = 5.+i;
+    }
+
+  a.import(read_write_2, VectorOperation::insert);
+  b.import(read_write_1, VectorOperation::insert);
+  c.import(read_write_2, VectorOperation::insert);
+
+
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(read_write_2[i] == read_write_3[i],
+                ExcMessage("Vector a has been modified."));
+
+  read_write_3.import(b, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(read_write_1[i] == read_write_3[i],
+                ExcMessage("Vector b has been modified."));
+
+  read_write_3.import(c, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(read_write_2[i] == read_write_3[i],
+                ExcMessage("Vector c has been modified."));
+
+  a *= 2.;
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(2.*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in operator *=."));
+
+  c /= 2.;
+  read_write_3.import(c, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(0.5*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in operator /=."));
+
+  b += a;
+  read_write_3.import(b, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(2.*read_write_2[i]+read_write_1[i]==read_write_3[i],
+                ExcMessage("Problem in operator +=."));
+
+  b -= c;
+  read_write_3.import(b, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(1.5*read_write_2[i]+read_write_1[i]==read_write_3[i],
+                ExcMessage("Problem in operator -=."));
+
+  b.import(read_write_1, VectorOperation::insert);
+  c.import(read_write_1, VectorOperation::insert);
+  const double val = b*c;
+  AssertThrow(val==328350., ExcMessage("Problem in operator *."));
+}
+
+int main(int argc, char **argv)
+{
+  std::ofstream logfile("output");
+  deallog.attach(logfile);
+  deallog.depth_console(0);
+  deallog.threshold_double(1.e-10);
+
+  test();
+
+  deallog << "OK" <<std::endl;
+
+  return 0;
+}
diff --git a/tests/cuda/cuda_vector_01.output b/tests/cuda/cuda_vector_01.output
new file mode 100644 (file)
index 0000000..0fd8fc1
--- /dev/null
@@ -0,0 +1,2 @@
+
+DEAL::OK
diff --git a/tests/cuda/cuda_vector_02.cc b/tests/cuda/cuda_vector_02.cc
new file mode 100644 (file)
index 0000000..c17999c
--- /dev/null
@@ -0,0 +1,113 @@
+// ---------------------------------------------------------------------
+//
+// Copyright (C) 2015 by the deal.II authors
+//
+// This file is part of the deal.II library.
+//
+// The deal.II library is free software; you can use it, redistribute
+// it, and/or modify it under the terms of the GNU Lesser General
+// Public License as published by the Free Software Foundation; either
+// version 2.1 of the License, or (at your option) any later version.
+// The full text of the license can be found in the file LICENSE at
+// the top level of the deal.II distribution.
+//
+// ---------------------------------------------------------------------
+
+
+#include "../tests.h"
+#include <deal.II/base/utilities.h>
+#include <deal.II/lac/read_write_vector.h>
+#include <deal.II/lac/cuda_vector.h>
+#include <fstream>
+#include <iostream>
+#include <vector>
+
+// Check LinearAlgebra::CUDAWrappers::Vector add and sadd.
+
+void test()
+{
+  const unsigned int size = 100;
+  LinearAlgebra::CUDAWrappers::Vector<double> a(size);
+  LinearAlgebra::CUDAWrappers::Vector<double> b(size);
+  LinearAlgebra::CUDAWrappers::Vector<double> c(size);
+
+  LinearAlgebra::ReadWriteVector<double> read_write_1(size);
+  LinearAlgebra::ReadWriteVector<double> read_write_2(size);
+  LinearAlgebra::ReadWriteVector<double> read_write_3(size);
+
+  for (unsigned int i=0; i<size; ++i)
+    {
+      read_write_1[i] = i;
+      read_write_2[i] = 5.+i;
+    }
+
+  a.import(read_write_1, VectorOperation::insert);
+  b.import(read_write_2, VectorOperation::insert);
+  c.import(read_write_2, VectorOperation::insert);
+
+  a.add(1.);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(1.+read_write_1[i]==read_write_3[i],
+                ExcMessage("Problem in add(scalar)."));
+
+  a.add(2.,b);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(1.+read_write_1[i]+2.*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in add(scalar,Vector)."));
+
+  LinearAlgebra::CUDAWrappers::Vector<double> d(a);
+  a.add(2.,b,3.,d);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(4.+4.*read_write_1[i]+10.*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in add(scalar,Vector,scalar,Vector)."));
+
+  a.import(read_write_1, VectorOperation::insert);
+  a.sadd(3.,2.,c);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(3.*read_write_1[i]+2.*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in sadd(scalar,scalar,Vector)."));
+
+  a.import(read_write_1, VectorOperation::insert);
+  a.scale(b);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(read_write_1[i]*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in scale."));
+
+  a.equ(2.,c);
+  read_write_3.import(a, VectorOperation::insert);
+  for (unsigned int i=0; i<size; ++i)
+    AssertThrow(2.*read_write_2[i]==read_write_3[i],
+                ExcMessage("Problem in equ."));
+
+  AssertThrow(b.l1_norm()==5450., ExcMessage("Problem in l1_norm."));
+
+  const double eps=1e-3;
+  AssertThrow(std::fabs(b.l2_norm()-616.725222)<eps,
+              ExcMessage("Problem in l2_norm"));
+
+  AssertThrow(b.linfty_norm()==104., ExcMessage("Problem in linfty_norm."));
+
+  a.import(read_write_1, VectorOperation::insert);
+  const double val = a.add_and_dot(2.,a,b);
+  AssertThrow(val==1059300., ExcMessage("Problem in add_and_dot"));
+}
+
+
+int main(int argc, char **argv)
+{
+  std::ofstream logfile("output");
+  deallog.attach(logfile);
+  deallog.depth_console(0);
+  deallog.threshold_double(1.e-10);
+
+  test();
+
+  deallog << "OK" <<std::endl;
+
+  return 0;
+}
diff --git a/tests/cuda/cuda_vector_02.output b/tests/cuda/cuda_vector_02.output
new file mode 100644 (file)
index 0000000..0fd8fc1
--- /dev/null
@@ -0,0 +1,2 @@
+
+DEAL::OK
index e0c5eb1fc683ad140a00bddb29ed908e5b0117eb..cc8a6cfc774cb8ea6f3601338468d700866cfd0f 100644 (file)
@@ -170,7 +170,7 @@ void test()
     {
       for (unsigned int i=5; i<10; ++i)
         AssertThrow(2.*read_write_2[i]==read_write_3[i],
-                    ExcMessage("Problem in scale."));
+                    ExcMessage("Problem in equ."));
     }
 
 

In the beginning the Universe was created. This has made a lot of people very angry and has been widely regarded as a bad move.

Douglas Adams


Typeset in Trocchi and Trocchi Bold Sans Serif.