diff --git a/src/core/device.cc b/src/core/device.cc
index 5216d37..5d1889e 100644
--- a/src/core/device.cc
+++ b/src/core/device.cc
@@ -26,6 +26,8 @@
  */
 #include "astaroth_device.h"
 
+#include <string.h> // memcpy
+
 #include "errchk.h"
 #include "math_utils.h"
 
@@ -289,6 +291,10 @@ acDeviceLoadScalarUniform(const Device device, const Stream stream, const AcReal
                           const AcReal value)
 {
     cudaSetDevice(device->id);
+    
+    if (param >= NUM_REAL_PARAMS)
+        return AC_FAILURE;
+
     const size_t offset = (size_t)&d_mesh_info.real_params[param] - (size_t)&d_mesh_info;
     ERRCHK_CUDA(cudaMemcpyToSymbolAsync(&d_mesh_info, &value, sizeof(value), offset,
                                         cudaMemcpyHostToDevice, device->streams[stream]));
@@ -300,6 +306,10 @@ acDeviceLoadVectorUniform(const Device device, const Stream stream, const AcReal
                           const AcReal3 value)
 {
     cudaSetDevice(device->id);
+
+    if (param >= NUM_REAL3_PARAMS || !NUM_REAL3_PARAMS)
+        return AC_FAILURE;
+
     const size_t offset = (size_t)&d_mesh_info.real3_params[param] - (size_t)&d_mesh_info;
     ERRCHK_CUDA(cudaMemcpyToSymbolAsync(&d_mesh_info, &value, sizeof(value), offset,
                                         cudaMemcpyHostToDevice, device->streams[stream]));
@@ -311,6 +321,10 @@ acDeviceLoadIntUniform(const Device device, const Stream stream, const AcIntPara
                        const int value)
 {
     cudaSetDevice(device->id);
+
+    if (param >= NUM_INT_PARAMS)
+        return AC_FAILURE;
+
     const size_t offset = (size_t)&d_mesh_info.int_params[param] - (size_t)&d_mesh_info;
     ERRCHK_CUDA(cudaMemcpyToSymbolAsync(&d_mesh_info, &value, sizeof(value), offset,
                                         cudaMemcpyHostToDevice, device->streams[stream]));
@@ -322,6 +336,10 @@ acDeviceLoadInt3Uniform(const Device device, const Stream stream, const AcInt3Pa
                         const int3 value)
 {
     cudaSetDevice(device->id);
+
+    if (param >= NUM_INT3_PARAMS)
+        return AC_FAILURE;
+
     const size_t offset = (size_t)&d_mesh_info.int3_params[param] - (size_t)&d_mesh_info;
     ERRCHK_CUDA(cudaMemcpyToSymbolAsync(&d_mesh_info, &value, sizeof(value), offset,
                                         cudaMemcpyHostToDevice, device->streams[stream]));
@@ -334,10 +352,13 @@ acDeviceLoadScalarArray(const Device device, const Stream stream, const ScalarAr
 {
     cudaSetDevice(device->id);
 
+    if (handle >= NUM_SCALARARRAY_HANDLES || !NUM_SCALARARRAY_HANDLES)
+        return AC_FAILURE;
+
     ERRCHK((int)(start + num) <= max(device->local_config.int_params[AC_mx],
                                      max(device->local_config.int_params[AC_my],
                                          device->local_config.int_params[AC_mz])));
-
+    ERRCHK_ALWAYS(handle < NUM_SCALARARRAY_HANDLES);
     ERRCHK_CUDA(cudaMemcpyAsync(&device->vba.profiles[handle][start], data, sizeof(data[0]) * num,
                                 cudaMemcpyHostToDevice, device->streams[stream]));
     return AC_SUCCESS;
diff --git a/src/core/errchk.h b/src/core/errchk.h
index 9646a83..592289a 100644
--- a/src/core/errchk.h
+++ b/src/core/errchk.h
@@ -86,8 +86,8 @@ cuda_assert(cudaError_t code, const char* file, int line, bool abort = true)
     #undef WARNCHK
     #define ERRCHK(params)
     #define WARNCHK(params)
-    #define ERRCHK_CUDA(params) params;
-    #define WARNCHK_CUDA(params) params;
+    #define ERRCHK_CUDA(params) params
+    #define WARNCHK_CUDA(params) params
     #define ERRCHK_CUDA_KERNEL() {}
 #else
     #define ERRCHK_CUDA(params) { cuda_assert((params), __FILE__, __LINE__); }