From 580dcac03555e6f4b2105bd1cfd1a84ed7ed12ed Mon Sep 17 00:00:00 2001 From: Matteo Cicuttin <datafl4sh@toxicnet.eu> Date: Tue, 27 Apr 2021 18:16:47 +0200 Subject: [PATCH] Differentiation matrix in texture. --- include/kernels_gpu.h | 10 ++++++++++ src/kernels_cuda.cu | 7 ++++--- 2 files changed, 14 insertions(+), 3 deletions(-) diff --git a/include/kernels_gpu.h b/include/kernels_gpu.h index b7f95c1..f3a4c59 100644 --- a/include/kernels_gpu.h +++ b/include/kernels_gpu.h @@ -51,6 +51,16 @@ struct kernel_gpu_sizes<4> static const size_t parallel_dblocks = 1; }; +template<> +struct kernel_gpu_sizes<5> +{ + static const size_t num_bf = 56; + static const size_t cells_per_dblock = 12; + static const size_t dblock_bf = num_bf * cells_per_dblock; + static const size_t dblock_size = 128; + static const size_t parallel_dblocks = 1; +}; + struct kernel_gpu_sizes_runtime { size_t num_bf; diff --git a/src/kernels_cuda.cu b/src/kernels_cuda.cu index b37d3e8..80fb688 100644 --- a/src/kernels_cuda.cu +++ b/src/kernels_cuda.cu @@ -100,14 +100,15 @@ gpu_compute_field_derivatives(entity_data_gpu& edg, //else // compute_field_derivatives_kernel_curved<4>(ed, f, df_dx, df_dy, df_dz); break; -#if 0 + case 5: if (edg.g_order == 1) - gpu_compute_field_derivatives_kernel_planar<5>(edg, f, df_dx, df_dy, df_dz); + gpu_deriv_planar<5><<<num_blocks, THREADS_PER_BLOCK>>>(f, J, + Dtex, df_dx, df_dy, df_dz, num_elems, orients); //else // compute_field_derivatives_kernel_curved<5>(ed, f, df_dx, df_dy, df_dz); break; - +#if 0 case 6: if (edg.g_order == 1) gpu_compute_field_derivatives_kernel_planar<6>(edg, f, df_dx, df_dy, df_dz); -- GitLab