Merge branch 'release/6.0.1'

miguelcarcamov · Mar 15, 2019 · 7e81ddc · 7e81ddc
2 parents f1e8126 + f844bbc
commit 7e81ddc
Show file tree

Hide file tree

Showing 13 changed files with 419 additions and 16 deletions.
diff --git a/Makefile b/Makefile
@@ -28,7 +28,7 @@ ARCHFLAG += -gencode arch=compute_$(HIGHEST_SM),code=compute_$(HIGHEST_SM)
 endif
 endif
 
-main:	build/main.o cfits build/MSFITSIO.o build/functions.o build/directioncosines.o build/copyrightwarranty.o build/rngs.o build/rvgs.o build/f1dim.o build/mnbrak.o build/brent.o build/linmin.o  build/frprmn.o build/synthesizer.o build/imageProcessor.o build/chi2.o build/laplacian.o build/gridding.o build/entropy.o build/ioms.o build/totalvariation.o build/quadraticpenalization.o build/error.o
+main:	build/main.o cfits build/MSFITSIO.o build/functions.o build/directioncosines.o build/copyrightwarranty.o build/rngs.o build/rvgs.o build/f1dim.o build/mnbrak.o build/brent.o build/linmin.o  build/frprmn.o build/synthesizer.o build/imageProcessor.o build/chi2.o build/laplacian.o build/gridding.o build/entropy.o build/ioms.o build/totalvariation.o build/quadraticpenalization.o build/error.o build/lbfgs.o
 	@ echo "Linking CUDAMEM"
 	@ mkdir -p bin
 	@ nvcc build/*.o -o bin/gpuvmem $(LDFLAGS) $(CFFLAG) $(FOPENFLAG) $(CUFFTFLAG) $(ARCHFLAG) $(CCFLAG)
@@ -134,6 +134,10 @@ build/frprmn.o: src/frprmn.cu
 	@ echo "Building frprmn"
 	@ nvcc $(CFLAGS) $(INC_DIRS) src/frprmn.cu -o build/frprmn.o $(LDFLAGS) $(CFFLAG) $(ARCHFLAG)
 
+build/lbfgs.o: src/lbfgs.cu
+	@ echo "Building lbfgs"
+	@ nvcc $(CFLAGS) $(INC_DIRS) src/lbfgs.cu -o build/lbfgs.o $(LDFLAGS) $(CFFLAG) $(ARCHFLAG)
+
 cfits:
 	@ mkdir -p lib
 	@ cd cfitsio; make; cp libcfitsio.a ../lib/.
@@ -143,6 +147,7 @@ cleanall:
 	@ rm -rf build/*
 	@ rm -rf bin/*
 	@ rm -f lib/*.a
+	@ rm -r *.fits
 	@ cd cfitsio; make clean
 	@ cd cfitsio; make distclean
 
@@ -156,7 +161,7 @@ conf:
 	@ ./configure
 
 co65:
-	@ ./bin/gpuvmem -i ./tests/co65/co65.ms -o ./tests/co65/co65_out.ms -O ./tests/co65/mod_out.fits -m ./tests/co65/mod_in_0.fits -I ./tests/co65/input.dat -p ./tests/co65/mem/ -X 16 -Y 16 -V 256 -z 0.001 -Z 0.01,0.5 -t 500000000 --verbose --print-images --print-errors
+	@ ./bin/gpuvmem -i ./tests/co65/co65.ms -o ./tests/co65/co65_out.ms -O ./tests/co65/mod_out.fits -m ./tests/co65/mod_in_0.fits -I ./tests/co65/input.dat -p ./tests/co65/mem/ -X 16 -Y 16 -V 256 -z 0.001 -Z 0.01,0.05 -t 500000000 --verbose --print-images --print-errors
 selfcalband9:
 	@ cuda-memcheck ./bin/gpuvmem -i ./tests/selfcalband9/hd142_b9cont_self_tav.ms -o ./tests/selfcalband9/hd142_b9cont_out.ms -O ./tests/selfcalband9/mod_out.fits -m ./tests/selfcalband9/mod_in_0.fits -I ./tests/selfcalband9/input.dat -p ./tests/selfcalband9/mem/ -X 8 -Y 8 -V 256 --verbose -z 0.001 -Z 0.05,0.5 -t 500000000 --print-images
 freq78:

diff --git a/include/framework.cuh b/include/framework.cuh
@@ -119,6 +119,9 @@ void setDS(float *DS){
         cudaFree(device_DS); this->device_DS = DS;
 };
 virtual float calculateSecondDerivate() = 0;
+float getPenalizationFactor(){
+        return this->penalization_factor;
+}
 protected:
 float *device_S;
 float *device_DS;
@@ -258,7 +261,8 @@ public:
 ObjectiveFunction(){
 };
 void addFi(Fi *fi){
-        fis.push_back(fi);
+        if(fi->getPenalizationFactor())
+                fis.push_back(fi);
 };
 //virtual void print() = 0;
 float calcFunction(float *p)

diff --git a/include/functions.cuh b/include/functions.cuh
@@ -107,5 +107,12 @@ __global__ void newPNoPositivity(float *p, float *xi, float xmin, long N, long M
 __global__ void evaluateXtNoPositivity(float *xt, float *pcom, float *xicom, float x, long N, long M, int image);
 __global__ void chainRule2I(float *chain, float *noise, float *I, float nu, float nu_0, float noise_cut, float fg_scale, long N, long M);
 __global__ void DChi2_2I(float *noise, float *chain, float *dchi2, float *dchi2_total, float threshold, int image, long N, long M);
+__global__ void calculateSandY (float *d_y, float *d_s, float *p, float *xi, float *p_old, float *xi_old, int iter, int M, int N, int image);
+__global__ void getR (float *d_r, float *d_q, float scalar, int M, int N, int image);
+__global__ void updateQ (float *d_q, float alpha, float *d_y, int k, int M, int N, int image);
+__global__ void getDot_LBFGS_ff(float *aux_vector, float *vec_1, float *vec_2, int k, int h, int M, int N, int image);
+__global__ void searchDirection_LBFGS(float *xi, long N, long M, int image);
+
+
 
 #endif
diff --git a/include/lbfgs.cuh b/include/lbfgs.cuh
@@ -0,0 +1,26 @@
+#ifndef LBFGS_CUH
+#define LBFGS_CUH
+#include "linmin.cuh"
+
+class LBFGS : public Optimizator
+{
+public:
+__host__ void allocateMemoryGpu();
+__host__ void deallocateMemoryGpu();
+__host__ void minimizate();
+__host__ void LBFGS_recursion(float *d_y, float *d_s, float *xi, int par_M, int lbfgs_it, int M, int N);
+private:
+float ftol = 0;
+float fret = 0;
+float gg, dgg, gam, fp;
+float *device_g, *device_h;
+float *device_gg_vector, *device_dgg_vector;
+int configured = 1;
+float *d_s;
+float *d_y, *xi, *xi_old, *norm_vector, *d_r;
+float norm;
+float *p_old;
+int K = 200;
+};
+
+#endif
diff --git a/src/chi2.cu b/src/chi2.cu
@@ -23,7 +23,7 @@ void Chi2::configure(int penalizatorIndex, int imageIndex, int imageToAdd)
 
         if(penalizatorIndex != -1)
         {
-                if(penalizatorIndex > (nPenalizators - 1))
+                if(penalizatorIndex > (nPenalizators - 1) || penalizatorIndex < 0)
                 {
                         printf("invalid index for penalizator (chi2)\n");
                         exit(-1);

diff --git a/src/entropy.cu b/src/entropy.cu
@@ -45,7 +45,7 @@ void Entropy::configure(int penalizatorIndex, int imageIndex, int imageToAdd)
 
         if(penalizatorIndex != -1)
         {
-                if(penalizatorIndex > (nPenalizators - 1))
+                if(penalizatorIndex > (nPenalizators - 1) || penalizatorIndex < 0)
                 {
                         printf("invalid index for penalizator (entropy)\n");
                         exit(-1);

diff --git a/src/functions.cu b/src/functions.cu
@@ -1521,6 +1521,48 @@ __global__ void searchDirection(float *g, float *xi, float *h, long N)
         xi[N*i+j] = h[N*i+j] = g[N*i+j];
 }
 
+__global__ void searchDirection_LBFGS(float *xi, long N, long M, int image)
+{
+  int j = threadIdx.x + blockDim.x * blockIdx.x;
+	int i = threadIdx.y + blockDim.y * blockIdx.y;
+
+  xi[M*N*image+N*i+j] *= -1.0f;
+}
+
+__global__ void getDot_LBFGS_ff(float *aux_vector, float *vec_1, float *vec_2, int k, int h, int M, int N, int image)
+{
+  int j = threadIdx.x + blockDim.x * blockIdx.x;
+	int i = threadIdx.y + blockDim.y * blockIdx.y;
+
+  aux_vector[N*i+j] = vec_1[M*N*image*k + M*N*image + (N*i+j)]*vec_2[M*N*image*h + M*N*image + (N*i+j)];
+}
+
+__global__ void updateQ (float *d_q, float alpha, float *d_y, int k, int M, int N, int image)
+{
+  int j = threadIdx.x + blockDim.x * blockIdx.x;
+	int i = threadIdx.y + blockDim.y * blockIdx.y;
+
+  d_q[M*N*image+N*i+j] += alpha *d_y[M*N*image + M*N*k + (N*i+j)];
+}
+
+__global__ void getR (float *d_r, float *d_q, float scalar, int M, int N, int image)
+{
+  int j = threadIdx.x + blockDim.x * blockIdx.x;
+	int i = threadIdx.y + blockDim.y * blockIdx.y;
+
+  d_r[M*N*image+N*i+j] = d_q[M*N*image+N*i+j] * scalar;
+}
+
+__global__ void calculateSandY (float *d_y, float *d_s, float *p, float *xi, float *p_old, float *xi_old, int iter, int M, int N, int image)
+{
+  int j = threadIdx.x + blockDim.x * blockIdx.x;
+	int i = threadIdx.y + blockDim.y * blockIdx.y;
+
+  d_y[M*N*image*iter + M*N*image + (N*i+j)] = xi[M*N*image+N*i+j] - (-1.0f*xi_old[M*N*image+N*i+j]);
+  d_s[M*N*image*iter + M*N*image + (N*i+j)] = p[M*N*image+N*i+j] - p_old[M*N*image+N*i+j];
+
+}
+
 __global__ void searchDirection(float* g, float* xi, float* h, long N, long M, int image)
 {
         int j = threadIdx.x + blockDim.x * blockIdx.x;

diff --git a/src/laplacian.cu b/src/laplacian.cu
@@ -45,7 +45,7 @@ void Laplacian::configure(int penalizatorIndex, int imageIndex, int imageToAdd)
 
         if(penalizatorIndex != -1)
         {
-                if(penalizatorIndex > (nPenalizators - 1))
+                if(penalizatorIndex > (nPenalizators - 1) || penalizatorIndex < 0)
                 {
                         printf("invalid index for penalizator (laplacian)\n");
                         exit(-1);