Merge branch 'emdant/dg-cherry-pick-8f1f09f5' into devel

src/libgmshdg/kernels_cuda.cu

@@ -138,19 +138,6 @@ launch_deriv_kernel(const entity_data_gpu& edg,
 
     using KS = kernel_gpu_sizes<K>;
 
-#ifdef USE_BLOCKED_GPU_KERNELS
-    size_t num_blocks = edg.num_all_elems / (KS::cells_per_dblock * KS::parallel_dblocks);
-    if (num_blocks % (KS::cells_per_dblock * KS::parallel_dblocks))
-        num_blocks += 1;
-
-    dim3 grid_size(1, num_blocks);
-    dim3 block_size(KS::dblock_size, KS::parallel_dblocks);
-    if (edg.g_order == 1)
-        gpu_deriv_planar_blocked<K><<<grid_size, block_size, stream>>>(f, J,
-            Dtex, df_dx, df_dy, df_dz, num_elems, orients, edg.dof_base);
-    //else
-    //    compute_field_derivatives_kernel_curved<1>(ed, f, df_dx, df_dy, df_dz);
-#else
     auto num_blocks = edg.num_bf*edg.num_all_elems/KS::deriv_threads;
     if (edg.num_bf*edg.num_all_elems % KS::deriv_threads)
         num_blocks += 1;
@@ -160,7 +147,6 @@ launch_deriv_kernel(const entity_data_gpu& edg,
             Dtex, df_dx, df_dy, df_dz, alpha, num_elems, orients, edg.dof_base);
     //else
     //    compute_field_derivatives_kernel_curved<1>(ed, f, df_dx, df_dy, df_dz);
-#endif
 }
 
 void

tests/test_differentiation_gpu.cpp

@@ -53,7 +53,7 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
     std::vector<double> sizes({ 0.32, 0.16, 0.08, 0.04 });
     std::vector<double> errors;
 
-    std::cout << cyanfg << "Testing geometric order " << geometric_order;
+    std::cout << std::endl << cyanfg << "Testing geometric order " << geometric_order;
     std::cout << ", approximation order = " << approximation_order << nofg;
     std::cout << std::endl;
 
@@ -101,13 +101,8 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
         vecxd Pdf_dz    = vecxd::Zero(model_num_dofs);
 
         std::vector<entity_data_gpu> edgs;
-
-#ifdef USE_BLOCKED_GPU_KERNELS
-        std::vector<entity_data_cpu> eds;
-        size_t model_num_dofs_gpu = 0;
-#else
         size_t model_num_dofs_gpu = model_num_dofs;
-#endif
+
         for (auto& e : mod)
         {
             e.project(f, Pf);
@@ -117,24 +112,13 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
             entity_data_cpu ed;
             e.populate_entity_data(ed, mod);
             entity_data_gpu edg(ed);
-#ifdef USE_BLOCKED_GPU_KERNELS
-            edg.dof_base = model_num_dofs_gpu;
-            model_num_dofs_gpu += gpu_dblocks_dofs(ed);
-            eds.push_back( std::move(ed) );
-#endif
             edgs.push_back( std::move(edg) );
         }
 
          /* Prepare I/O vectors and call kernel */
-#ifdef USE_BLOCKED_GPU_KERNELS
-        assert(eds.size() == edgs.size());
-        vecxd Pf_reshaped = vecxd::Zero(model_num_dofs_gpu);
-        for (size_t i = 0; i < eds.size(); i++)
-            reshape_dofs(eds[i], edgs[i], Pf, Pf_reshaped, true);
-        texture_allocator<double> Pf_gpu(Pf_reshaped.data(), Pf_reshaped.size());
-#else
+
         device_vector<double> Pf_gpu(Pf.data(), Pf.size());
-#endif
+
         device_vector<double> df_dx_gpu(model_num_dofs_gpu);
         device_vector<double> df_dy_gpu(model_num_dofs_gpu);
         device_vector<double> df_dz_gpu(model_num_dofs_gpu);
@@ -161,30 +145,6 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
                 std::cout << flops/(1e9*time) << " GFlops/s" << std::endl;
             }
         }
-
-#ifdef USE_BLOCKED_GPU_KERNELS
-        vecxd Cdf_dx_exp = vecxd::Zero(model_num_dofs_gpu);
-        vecxd Cdf_dy_exp = vecxd::Zero(model_num_dofs_gpu);
-        vecxd Cdf_dz_exp = vecxd::Zero(model_num_dofs_gpu);
-
-        df_dx_gpu.copyout( Cdf_dx_exp.data() );
-        df_dy_gpu.copyout( Cdf_dy_exp.data() );
-        df_dz_gpu.copyout( Cdf_dz_exp.data() );
-
-        vecxd Cdf_dx = vecxd::Zero(model_num_dofs);
-        vecxd Cdf_dy = vecxd::Zero(model_num_dofs);
-        vecxd Cdf_dz = vecxd::Zero(model_num_dofs);
-
-        assert(eds.size() == edgs.size());
-        Pf = vecxd::Zero(model_num_dofs);
-        for (size_t i = 0; i < eds.size(); i++)
-        {
-            reshape_dofs(eds[i], edgs[i], Cdf_dx_exp, Cdf_dx, false);
-            reshape_dofs(eds[i], edgs[i], Cdf_dy_exp, Cdf_dy, false);
-            reshape_dofs(eds[i], edgs[i], Cdf_dz_exp, Cdf_dz, false);
-            reshape_dofs(eds[i], edgs[i], Pf_reshaped, Pf, false);
-        }
-#else
         vecxd Cdf_dx = vecxd::Zero(model_num_dofs_gpu);
         vecxd Cdf_dy = vecxd::Zero(model_num_dofs_gpu);
         vecxd Cdf_dz = vecxd::Zero(model_num_dofs_gpu);
@@ -192,7 +152,6 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
         df_dx_gpu.copyout( Cdf_dx.data() );
         df_dy_gpu.copyout( Cdf_dy.data() );
         df_dz_gpu.copyout( Cdf_dz.data() );
-#endif
         
         double err_x = 0.0;
         double err_y = 0.0;
@@ -247,6 +206,8 @@ int test_differentiation_convergence(int geometric_order, int approximation_orde
         errs_x.push_back( std::sqrt(err_x) );
         errs_y.push_back( std::sqrt(err_y) );
         errs_z.push_back( std::sqrt(err_z) );
+
+        std::cout << std::endl;
     }
 
     double rate_x = 0.0;