From 545588bcc09ebad7e6f00fc176bc2960120ddc1e Mon Sep 17 00:00:00 2001
From: Sarod Yatawatta <sarod@dop373.nfra.nl>
Date: Mon, 27 Nov 2017 20:30:39 +0100
Subject: [PATCH] bug fix

---
 src/lib/predict_withbeam_gpu.c | 16 +++++++++++++++-
 1 file changed, 15 insertions(+), 1 deletion(-)

diff --git a/src/lib/predict_withbeam_gpu.c b/src/lib/predict_withbeam_gpu.c
index 743ad9d..135bb4e 100644
--- a/src/lib/predict_withbeam_gpu.c
+++ b/src/lib/predict_withbeam_gpu.c
@@ -173,6 +173,8 @@ precalcoh_threadfn(void *data) {
 
   err=cudaSetDevice(card);
   checkCudaError(err,__FILE__,__LINE__);
+  err=cudaSetDeviceFlags(cudaDeviceLmemResizeToMax);
+  checkCudaError(err,__FILE__,__LINE__);
 
   /* make sure enough heap memory is available for shapelet computations */
   size_t plim;
@@ -592,6 +594,7 @@ precalcoh_threadfn(void *data) {
 
   /* reset error state */
   err=cudaGetLastError(); 
+
   return NULL;
 
 }
@@ -765,6 +768,7 @@ precalculate_coherencies_withbeam_gpu(double *u, double *v, double *w, complex d
  free(threaddata1);
  pthread_attr_destroy(&attr);
  free(th_array);
+
  return 0;
 }
 
@@ -788,6 +792,8 @@ predictvis_threadfn(void *data) {
 
   err=cudaSetDevice(card);
   checkCudaError(err,__FILE__,__LINE__);
+  err=cudaSetDeviceFlags(cudaDeviceLmemResizeToMax);
+  checkCudaError(err,__FILE__,__LINE__);
 
   /* make sure enough heap memory is available for shapelet computations */
   size_t plim;
@@ -1209,6 +1215,7 @@ predictvis_threadfn(void *data) {
   }
 
 
+
   /* reset error state */
   err=cudaGetLastError(); 
   return NULL;
@@ -1358,6 +1365,8 @@ residual_threadfn(void *data) {
 
   err=cudaSetDevice(card);
   checkCudaError(err,__FILE__,__LINE__);
+  err=cudaSetDeviceFlags(cudaDeviceLmemResizeToMax);
+  checkCudaError(err,__FILE__,__LINE__);
 
   /* make sure enough heap memory is available for shapelet computations */
   size_t plim;
@@ -1792,6 +1801,7 @@ residual_threadfn(void *data) {
   }
 
 
+
   /* reset error state */
   err=cudaGetLastError(); 
   return NULL;
@@ -1815,6 +1825,8 @@ correct_threadfn(void *data) {
 
   err=cudaSetDevice(card);
   checkCudaError(err,__FILE__,__LINE__);
+  err=cudaSetDeviceFlags(cudaDeviceLmemResizeToMax);
+  checkCudaError(err,__FILE__,__LINE__);
 
   double *xd;
   baseline_t *barrd;
@@ -1831,7 +1843,7 @@ correct_threadfn(void *data) {
   /* copy with right offset */
   err=cudaMemcpy(barrd, &(t->barr[t->boff]), t->Nb*sizeof(baseline_t), cudaMemcpyHostToDevice);
   checkCudaError(err,__FILE__,__LINE__);
-  err=cudaMemcpy(pd, t->pinv, t->N*8*t->Nf*sizeof(double), cudaMemcpyHostToDevice);
+  err=cudaMemcpy(pd, t->pinv, t->N*8*t->nchunk*sizeof(double), cudaMemcpyHostToDevice);
   checkCudaError(err,__FILE__,__LINE__);
   
   for (cf=0; cf<t->Nf; cf++) {
@@ -1855,6 +1867,7 @@ correct_threadfn(void *data) {
   err=cudaFree(pd);
   checkCudaError(err,__FILE__,__LINE__);
 
+
   /* reset error state */
   err=cudaGetLastError();
   return NULL;
@@ -2071,5 +2084,6 @@ double ph_ra0, double ph_dec0, double ph_freq0, double *longitude, double *latit
   pthread_attr_destroy(&attr);
 
   destroy_task_hist(&thst);
+
   return 0;
 }