  for(size_t l=_ti;l<tsize;l+=_gd){arr0[l]+=(arr1[l]*val);}
  return;
 }
-//----------------------------------------------------------------------------------------------------------
+//-------------------------------------------------------------------------------------------------------------------
 // ARRAY DOT-PRODUCT (R4):
-__global__ void gpu_array_dot_product_r4__(size_t tsize, const float *arr1, const float *arr2, float *dprod)
+__global__ void gpu_array_dot_product_r4__(size_t tsize, const float *arr1, const float *arr2, volatile float *dprod)
 {
  extern __shared__ float dprs_r4[]; //volume = blockDim.x
  size_t l;

// Infrastructure for functions <gpu_array_2norm2_XX> (blocking):
__device__ float gpu_blck_norms2_r4[MAX_CUDA_BLOCKS];
__device__ double gpu_blck_norms2_r8[MAX_CUDA_BLOCKS];
static float blck_norms2_r4[MAX_CUDA_BLOCKS];
static double blck_norms2_r8[MAX_CUDA_BLOCKS];
// Infrastructure for kernels <gpu_array_dot_product_XX__>:
__device__ int dot_product_wr_lock=0; //write lock for all simultaneously running kernels <gpu_array_dot_product_XX__>
#ifndef NO_BLAS
// Infrastructure for CUBLAS:
static cublasHandle_t cublas_handle[MAX_GPUS_PER_NODE]; //each GPU present on a node obtains its own context handle
#endif

//------------------------------------------------------------------------------------------------------

/** Tensor Algebra Library for NVidia GPUs (CUDA).
REVISION: 2015/01/27
Copyright (C) 2015 Dmitry I. Lyakh (email: quant4me@gmail.com)
Copyright (C) 2015 Oak Ridge National Laboratory (UT-Battelle)

