]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Fix a bug in CUDA matix-free when there is no constraint 4910/head
authorBruno Turcksin <bruno.turcksin@gmail.com>
Sat, 19 Aug 2017 22:45:12 +0000 (18:45 -0400)
committerBruno Turcksin <bruno.turcksin@gmail.com>
Sat, 19 Aug 2017 22:45:12 +0000 (18:45 -0400)
include/deal.II/matrix_free/cuda_matrix_free.templates.cuh

index e954a2d4c53c270fb43f5245571aaec13c0c501b..56e1cc94702e6527d6787b50d71f4a6340f31e69 100644 (file)
@@ -519,36 +519,40 @@ namespace CUDAWrappers
     // Constrained indices
     n_constrained_dofs = constraints.n_constraints();
 
-    const unsigned int constraint_n_blocks = std::ceil(static_cast<double>(n_constrained_dofs) /
-                                                       static_cast<double>(BLOCK_SIZE));
-    const unsigned int constraint_x_n_blocks = std::round(std::sqrt(constraint_n_blocks));
-    const unsigned int constraint_y_n_blocks = std::ceil(static_cast<double>(constraint_n_blocks) /
-                                                         static_cast<double>(constraint_x_n_blocks));
+    if (n_constrained_dofs != 0)
+      {
 
-    constraint_grid_dim = dim3(constraint_x_n_blocks, constraint_y_n_blocks);
-    constraint_block_dim = dim3(BLOCK_SIZE);
+        const unsigned int constraint_n_blocks = std::ceil(static_cast<double>(n_constrained_dofs) /
+                                                           static_cast<double>(BLOCK_SIZE));
+        const unsigned int constraint_x_n_blocks = std::round(std::sqrt(constraint_n_blocks));
+        const unsigned int constraint_y_n_blocks = std::ceil(static_cast<double>(constraint_n_blocks) /
+                                                             static_cast<double>(constraint_x_n_blocks));
 
-    std::vector<dealii::types::global_dof_index> constrained_dofs_host(n_constrained_dofs);
+        constraint_grid_dim = dim3(constraint_x_n_blocks, constraint_y_n_blocks);
+        constraint_block_dim = dim3(BLOCK_SIZE);
 
-    unsigned int i_constraint = 0;
-    const unsigned int n_dofs = dof_handler.n_dofs();
-    for (unsigned int i=0; i<n_dofs; ++i)
-      {
-        if (constraints.is_constrained(i))
+        std::vector<dealii::types::global_dof_index> constrained_dofs_host(n_constrained_dofs);
+
+        unsigned int i_constraint = 0;
+        const unsigned int n_dofs = dof_handler.n_dofs();
+        for (unsigned int i=0; i<n_dofs; ++i)
           {
-            constrained_dofs_host[i_constraint] = i;
-            ++i_constraint;
+            if (constraints.is_constrained(i))
+              {
+                constrained_dofs_host[i_constraint] = i;
+                ++i_constraint;
+              }
           }
-      }
 
-    cuda_error = cudaMalloc(&constrained_dofs, n_constrained_dofs *
-                            sizeof(dealii::types::global_dof_index));
-    AssertCuda(cuda_error);
+        cuda_error = cudaMalloc(&constrained_dofs, n_constrained_dofs *
+                                sizeof(dealii::types::global_dof_index));
+        AssertCuda(cuda_error);
 
-    cuda_error = cudaMemcpy(constrained_dofs, &constrained_dofs_host[0],
-                            n_constrained_dofs * sizeof(dealii::types::global_dof_index),
-                            cudaMemcpyHostToDevice);
-    AssertCuda(cuda_error);
+        cuda_error = cudaMemcpy(constrained_dofs, &constrained_dofs_host[0],
+                                n_constrained_dofs * sizeof(dealii::types::global_dof_index),
+                                cudaMemcpyHostToDevice);
+        AssertCuda(cuda_error);
+      }
   }
 
 

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.