From 20a71b1d23dca6feb347749da621b887fccd2bef Mon Sep 17 00:00:00 2001
From: sjplimp <sjplimp@f3b2605a-c512-4ea7-a41b-209d697bcdaa>
Date: Tue, 21 Oct 2014 16:50:00 +0000
Subject: [PATCH] git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@12640
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
lib/cuda/pair_cg_cmm_coul_cut_cuda.cu | 81 ---------------------
lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h | 26 -------
lib/cuda/pair_cg_cmm_coul_debye_cuda.cu | 81 ---------------------
lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h | 26 -------
lib/cuda/pair_cg_cmm_coul_long_cuda.cu | 81 ---------------------
lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h | 26 -------
lib/cuda/pair_cg_cmm_cuda.cu | 87 -----------------------
lib/cuda/pair_cg_cmm_cuda_cu.h | 26 -------
lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu | 49 -------------
9 files changed, 483 deletions(-)
delete mode 100644 lib/cuda/pair_cg_cmm_coul_cut_cuda.cu
delete mode 100644 lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h
delete mode 100644 lib/cuda/pair_cg_cmm_coul_debye_cuda.cu
delete mode 100644 lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h
delete mode 100644 lib/cuda/pair_cg_cmm_coul_long_cuda.cu
delete mode 100644 lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h
delete mode 100644 lib/cuda/pair_cg_cmm_cuda.cu
delete mode 100644 lib/cuda/pair_cg_cmm_cuda_cu.h
delete mode 100644 lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu
diff --git a/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu b/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu
deleted file mode 100644
index bafb83705a..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu
+++ /dev/null
@@ -1,81 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include <stdio.h>
-
-#define _lj1 MY_AP(coeff1)
-#define _lj2 MY_AP(coeff2)
-#define _lj3 MY_AP(coeff3)
-#define _lj4 MY_AP(coeff4)
-#define _cg_type MY_AP(coeff5)
-
-
-#include "pair_cg_cmm_coul_cut_cuda_cu.h"
-#include <time.h>
-
-
-
-
-void Cuda_PairCGCMMCoulCutCuda_Init(cuda_shared_data* sdata)
-{
- Cuda_Pair_Init_AllStyles(sdata, 5, true, false);
-
-}
-
-
-
-
-void Cuda_PairCGCMMCoulCutCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom)
-{
-
- // initialize only on first call
- static short init = 0;
-
- if(! init) {
- init = 1;
- Cuda_PairCGCMMCoulCutCuda_Init(sdata);
- }
-
- dim3 grid, threads;
- int sharedperproc;
-
- Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128);
-
- cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
-
- if(sdata->pair.use_block_per_atom)
- Pair_Kernel_BpA<PAIR_CG_CMM, COUL_CUT, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
- else
- Pair_Kernel_TpA<PAIR_CG_CMM, COUL_CUT, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
-
- Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
-}
-
-#undef _lj1
-#undef _lj2
-#undef _lj3
-#undef _lj4
-#undef _cg_type
-
diff --git a/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h
deleted file mode 100644
index fff7db7a64..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h
+++ /dev/null
@@ -1,26 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include "cuda_shared.h"
-
-extern "C" void Cuda_PairCGCMMCoulCutCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom);
diff --git a/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu b/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu
deleted file mode 100644
index ee18b81899..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu
+++ /dev/null
@@ -1,81 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include <stdio.h>
-
-#define _lj1 MY_AP(coeff1)
-#define _lj2 MY_AP(coeff2)
-#define _lj3 MY_AP(coeff3)
-#define _lj4 MY_AP(coeff4)
-#define _cg_type MY_AP(coeff5)
-
-
-#include "pair_cg_cmm_coul_debye_cuda_cu.h"
-#include <time.h>
-
-
-
-
-void Cuda_PairCGCMMCoulDebyeCuda_Init(cuda_shared_data* sdata)
-{
- Cuda_Pair_Init_AllStyles(sdata, 5, true, false);
-
-}
-
-
-
-
-void Cuda_PairCGCMMCoulDebyeCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom)
-{
-
- // initialize only on first call
- static short init = 0;
-
- if(! init) {
- init = 1;
- Cuda_PairCGCMMCoulDebyeCuda_Init(sdata);
- }
-
- dim3 grid, threads;
- int sharedperproc;
-
- Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128);
-
- cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
-
- if(sdata->pair.use_block_per_atom)
- Pair_Kernel_BpA<PAIR_CG_CMM, COUL_DEBYE, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
- else
- Pair_Kernel_TpA<PAIR_CG_CMM, COUL_DEBYE, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
-
- Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
-}
-
-#undef _lj1
-#undef _lj2
-#undef _lj3
-#undef _lj4
-#undef _cg_type
-
diff --git a/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h
deleted file mode 100644
index 5b0b6af597..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h
+++ /dev/null
@@ -1,26 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include "cuda_shared.h"
-
-extern "C" void Cuda_PairCGCMMCoulDebyeCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom);
diff --git a/lib/cuda/pair_cg_cmm_coul_long_cuda.cu b/lib/cuda/pair_cg_cmm_coul_long_cuda.cu
deleted file mode 100644
index 67c6ebaa7f..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_long_cuda.cu
+++ /dev/null
@@ -1,81 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include <stdio.h>
-
-#define _lj1 MY_AP(coeff1)
-#define _lj2 MY_AP(coeff2)
-#define _lj3 MY_AP(coeff3)
-#define _lj4 MY_AP(coeff4)
-#define _cg_type MY_AP(coeff5)
-
-
-#include "pair_cg_cmm_coul_long_cuda_cu.h"
-#include <time.h>
-
-
-
-
-void Cuda_PairCGCMMCoulLongCuda_Init(cuda_shared_data* sdata)
-{
- Cuda_Pair_Init_AllStyles(sdata, 5, true, false);
-
-}
-
-
-
-
-void Cuda_PairCGCMMCoulLongCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom)
-{
-
- // initialize only on first call
- static short init = 0;
-
- if(! init) {
- init = 1;
- Cuda_PairCGCMMCoulLongCuda_Init(sdata);
- }
-
- dim3 grid, threads;
- int sharedperproc;
-
- Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128);
-
- cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
-
- if(sdata->pair.use_block_per_atom)
- Pair_Kernel_BpA<PAIR_CG_CMM, COUL_LONG, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
- else
- Pair_Kernel_TpA<PAIR_CG_CMM, COUL_LONG, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
-
- Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
-}
-
-#undef _lj1
-#undef _lj2
-#undef _lj3
-#undef _lj4
-#undef _cg_type
-
diff --git a/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h
deleted file mode 100644
index 86befd78b8..0000000000
--- a/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h
+++ /dev/null
@@ -1,26 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include "cuda_shared.h"
-
-extern "C" void Cuda_PairCGCMMCoulLongCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom);
diff --git a/lib/cuda/pair_cg_cmm_cuda.cu b/lib/cuda/pair_cg_cmm_cuda.cu
deleted file mode 100644
index ccc4782ec0..0000000000
--- a/lib/cuda/pair_cg_cmm_cuda.cu
+++ /dev/null
@@ -1,87 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include <stdio.h>
-
-#define _lj1 MY_AP(coeff1)
-#define _lj2 MY_AP(coeff2)
-#define _lj3 MY_AP(coeff3)
-#define _lj4 MY_AP(coeff4)
-#define _cg_type MY_AP(coeff5)
-
-enum {CG_NOT_SET = 0, CG_LJ9_6, CG_LJ12_4, CG_LJ12_6, NUM_CG_TYPES,
- CG_COUL_NONE, CG_COUL_CUT, CG_COUL_DEBYE, CG_COUL_LONG
- };
-
-#include "pair_cg_cmm_cuda_cu.h"
-#include "pair_cg_cmm_cuda_kernel_nc.cu"
-#include <time.h>
-
-
-
-
-void Cuda_PairCGCMMCuda_Init(cuda_shared_data* sdata)
-{
- Cuda_Pair_Init_AllStyles(sdata, 5, false, false);
-
-}
-
-
-
-
-void Cuda_PairCGCMMCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom)
-{
-
- // initialize only on first call
- static short init = 0;
-
- if(! init) {
- init = 1;
- Cuda_PairCGCMMCuda_Init(sdata);
- }
-
- dim3 grid, threads;
- int sharedperproc;
-
- int maxthreads = 128;
-
- Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, false, maxthreads);
-
- cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
-
- if(sdata->pair.use_block_per_atom)
- Pair_Kernel_BpA<PAIR_CG_CMM, COUL_NONE, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
- else
- Pair_Kernel_TpA<PAIR_CG_CMM, COUL_NONE, DATA_NONE>
- <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
-
- Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
-}
-
-#undef _lj1
-#undef _lj2
-#undef _lj3
-#undef _lj4
-#undef _cg_type
-
diff --git a/lib/cuda/pair_cg_cmm_cuda_cu.h b/lib/cuda/pair_cg_cmm_cuda_cu.h
deleted file mode 100644
index 739c0ae28f..0000000000
--- a/lib/cuda/pair_cg_cmm_cuda_cu.h
+++ /dev/null
@@ -1,26 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-#include "cuda_shared.h"
-
-extern "C" void Cuda_PairCGCMMCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom);
diff --git a/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu b/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu
deleted file mode 100644
index a19903b41a..0000000000
--- a/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu
+++ /dev/null
@@ -1,49 +0,0 @@
-/* ----------------------------------------------------------------------
- LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
-
- Original Version:
- http://lammps.sandia.gov, Sandia National Laboratories
- Steve Plimpton, sjplimp@sandia.gov
-
- See the README file in the top-level LAMMPS directory.
-
- -----------------------------------------------------------------------
-
- USER-CUDA Package and associated modifications:
- https://sourceforge.net/projects/lammpscuda/
-
- Christian Trott, christian.trott@tu-ilmenau.de
- Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
- Theoretical Physics II, University of Technology Ilmenau, Germany
-
- See the README file in the USER-CUDA directory.
-
- This software is distributed under the GNU General Public License.
-------------------------------------------------------------------------- */
-
-__device__ inline F_CFLOAT PairCGCMMCuda_Eval(const F_CFLOAT &rsq, const int ij_type, F_CFLOAT &factor_lj, int &eflag, ENERGY_CFLOAT &evdwl) //0.11 of 0.4
-{
- const F_CFLOAT r2inv = F_F(1.0) / rsq;
- const int cg_type = _cg_type[ij_type];
- const F_CFLOAT r4inv = r2inv * r2inv;
- const F_CFLOAT rNinv_first = cg_type != CG_LJ9_6 ? r4inv : _RSQRT_(rsq);
- const F_CFLOAT rNinv_second = cg_type != CG_LJ12_4 ? -r2inv : -F_F(1.0);
- const F_CFLOAT forcelj = r4inv * (_lj1[ij_type] * r4inv * rNinv_first + _lj2[ij_type] * rNinv_second);
-
- if(eflag) evdwl += factor_lj * (r4inv * (_lj3[ij_type] * r4inv * rNinv_first + _lj4[ij_type] * rNinv_second) - _offset[ij_type]);
-
- return factor_lj * forcelj * r2inv;
-}
-
-/*__device__ inline F_CFLOAT PairCGCMMCuda_Eval(const F_CFLOAT& rsq,const int ij_type,F_CFLOAT& factor_lj,int& eflag, ENERGY_CFLOAT& evdwl)
-{
- const int cg_type = tex1Dfetch(_coeff5_gm_tex,ij_type);
- const F_CFLOAT r2inv = F_F(1.0)/rsq;
- const F_CFLOAT r4inv = r2inv*r2inv;
- const F_CFLOAT rNinv_first = cg_type!=CG_LJ9_6?r4inv:_RSQRT_(rsq);
- const F_CFLOAT rNinv_second = cg_type!=CG_LJ12_4?r2inv:F_F(1.0);
- const F_CFLOAT forcelj = r4inv * (tex1Dfetch(_coeff1_gm_tex,ij_type)*r4inv*rNinv_first - tex1Dfetch(_coeff2_gm_tex,ij_type)*rNinv_second);
-
- if(eflag) evdwl += factor_lj*(r4inv*(tex1Dfetch(_coeff3_gm_tex,ij_type)*r4inv*rNinv_first-tex1Dfetch(_coeff4_gm_tex,ij_type)*rNinv_second));
- return factor_lj*forcelj*r2inv;
-}*/
--
GitLab