Improve the debug message for kernel launch errors

Print out the kernel name and also some more detail about the dimensions
that are problematic.

src/basic/accel.F90

@@ -178,6 +178,7 @@ module accel_oct_m
     logical                       :: initialized = .false.
     type(accel_kernel_t), pointer :: next
     integer                       :: arg_count
+    character(len=128)            :: kernel_name
   end type accel_kernel_t
 
   type(accel_t), public :: accel
@@ -1320,15 +1321,38 @@ contains
 
 #ifdef HAVE_CUDA
     ! Maximum dimension of a block
-    ASSERT(all(lsizes(1:3) <= accel%max_block_dim(1:3)))
+    if (any(lsizes(1:3) > accel%max_block_dim(1:3))) then
+      message(1) = "Maximum dimension of a block too large in kernel "//trim(kernel%kernel_name)
+      message(2) = "The following conditions should be fulfilled:"
+      write(message(3), "(A, I8, A, I8)") "Dim 1: ", lsizes(1), " <= ", accel%max_block_dim(1)
+      write(message(4), "(A, I8, A, I8)") "Dim 2: ", lsizes(2), " <= ", accel%max_block_dim(2)
+      write(message(5), "(A, I8, A, I8)") "Dim 3: ", lsizes(3), " <= ", accel%max_block_dim(3)
+      message(6) = "This is an internal error, please contact the developers."
+      call messages_fatal(6)
+    end if
+
 
     ! Maximum number of threads per block
-    ASSERT(product(lsizes) <= accel_max_workgroup_size())
+    if (product(lsizes) > accel_max_workgroup_size()) then
+      message(1) = "Maximum number of threads per block too large in kernel "//trim(kernel%kernel_name)
+      message(2) = "The following condition should be fulfilled:"
+      write(message(3), "(I8, A, I8)") product(lsizes), " <= ", accel_max_workgroup_size()
+      message(4) = "This is an internal error, please contact the developers."
+      call messages_fatal(4)
+    end if
 
     gsizes(1:3) = gsizes(1:3)/lsizes(1:3)
 
     ! Maximum dimensions of the grid of thread block
-    ASSERT(all(gsizes(1:3) <= accel%max_grid_dim(1:3)))
+    if (any(gsizes(1:3) > accel%max_grid_dim(1:3))) then
+      message(1) = "Maximum dimension of grid too large in kernel "//trim(kernel%kernel_name)
+      message(2) = "The following conditions should be fulfilled:"
+      write(message(3), "(A, I8, A, I10)") "Dim 1: ", gsizes(1), " <= ", accel%max_grid_dim(1)
+      write(message(4), "(A, I8, A, I10)") "Dim 2: ", gsizes(2), " <= ", accel%max_grid_dim(2)
+      write(message(5), "(A, I8, A, I10)") "Dim 3: ", gsizes(3), " <= ", accel%max_grid_dim(3)
+      message(6) = "This is an internal error, please contact the developers."
+      call messages_fatal(6)
+    end if
 
     call cuda_launch_kernel(kernel%cuda_kernel, gsizes(1), lsizes(1), kernel%cuda_shared_mem, kernel%arguments)
 
@@ -1990,6 +2014,7 @@ contains
 #endif
 
     this%initialized = .true.
+    this%kernel_name = trim(kernel_name)
 
     call profiling_out(prof)