Merge pull request #360 from bcaddy/dev-iss241

Remove Deprecated CUDA Macro

builds/make.inc.template

@@ -1,11 +1,6 @@
 #POISSON_SOLVER ?= -DPFFT
 #DFLAGS += $(POISSON_SOLVER)
 
-
-#To use GPUs, CUDA must be turned on here
-#Optional error checking can also be enabled
-DFLAGS += -DCUDA
-
 #To use MPI, DFLAGS must include -DMPI_CHOLLA
 DFLAGS += -DMPI_CHOLLA
 

builds/make.type.basic_scalar

@@ -1,6 +1,5 @@
 #-- Default hydro build with BASIC_SCALAR
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.cloudy

@@ -6,7 +6,6 @@ OUTPUT    ?=  -DOUTPUT -DHDF5
 
 MPI_GPU   ?=
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPPMP

builds/make.type.cooling

@@ -6,7 +6,6 @@ OUTPUT    ?=  -DOUTPUT -DHDF5
 
 MPI_GPU   ?=
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPPMP

builds/make.type.disk

@@ -21,7 +21,6 @@ DFLAGS += -DGRAVITY_5_POINTS_GRADIENT
 
 #DFLAGS += -DSTATIC_GRAV
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.dust

@@ -6,7 +6,6 @@ OUTPUT    ?=  -DOUTPUT -DHDF5
 
 MPI_GPU   ?=
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.hydro

@@ -1,6 +1,5 @@
 #-- Default hydro only build
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.mhd

@@ -6,7 +6,6 @@ OUTPUT    ?=  -DOUTPUT -DHDF5
 
 MPI_GPU   ?=
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.rot_proj

@@ -1,6 +1,5 @@
 #-- Default hydro only build with rotated projection
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

builds/make.type.static_grav

@@ -1,6 +1,5 @@
 #-- Default hydro only build with static_grav
 
-DFLAGS    += -DCUDA
 DFLAGS    += -DMPI_CHOLLA
 DFLAGS    += -DPRECISION=2
 DFLAGS    += -DPLMC

src/cooling/cooling_cuda.cu

@@ -1,19 +1,18 @@
 /*! \file cooling_cuda.cu
  *  \brief Functions to calculate cooling rate for a given rho, P, dt. */
 
-#ifdef CUDA
-  #ifdef COOLING_GPU
+#ifdef COOLING_GPU
 
-    #include <math.h>
+  #include <math.h>
 
-    #include "../cooling/cooling_cuda.h"
-    #include "../global/global.h"
-    #include "../global/global_cuda.h"
-    #include "../utils/gpu.hpp"
+  #include "../cooling/cooling_cuda.h"
+  #include "../global/global.h"
+  #include "../global/global_cuda.h"
+  #include "../utils/gpu.hpp"
 
-    #ifdef CLOUDY_COOL
-      #include "../cooling/texture_utilities.h"
-    #endif
+  #ifdef CLOUDY_COOL
+    #include "../cooling/texture_utilities.h"
+  #endif
 
 cudaTextureObject_t coolTexObj = 0;
 cudaTextureObject_t heatTexObj = 0;
@@ -64,10 +63,10 @@ __global__ void cooling_kernel(Real *dev_conserved, int nx, int ny, int nz, int
   Real cool;  // cooling rate per volume, erg/s/cm^3
   // #ifndef DE
   Real vx, vy, vz, p;
-    // #endif
-    #ifdef DE
+  // #endif
+  #ifdef DE
   Real ge;
-    #endif
+  #endif
 
   mu = 0.6;
   // mu = 1.27;
@@ -94,29 +93,29 @@ __global__ void cooling_kernel(Real *dev_conserved, int nx, int ny, int nz, int
     vz = dev_conserved[3 * n_cells + id] / d;
     p  = (E - 0.5 * d * (vx * vx + vy * vy + vz * vz)) * (gamma - 1.0);
     p  = fmax(p, (Real)TINY_NUMBER);
-    // #endif
-    #ifdef DE
+  // #endif
+  #ifdef DE
     ge = dev_conserved[(n_fields - 1) * n_cells + id] / d;
     ge = fmax(ge, (Real)TINY_NUMBER);
-    #endif
+  #endif
 
     // calculate the number density of the gas (in cgs)
     n = d * DENSITY_UNIT / (mu * MP);
 
     // calculate the temperature of the gas
     T_init = p * PRESSURE_UNIT / (n * KB);
-    #ifdef DE
+  #ifdef DE
     T_init = d * ge * (gamma - 1.0) * PRESSURE_UNIT / (n * KB);
-    #endif
+  #endif
 
     // calculate cooling rate per volume
     T = T_init;
-    // call the cooling function
-    #ifdef CLOUDY_COOL
+  // call the cooling function
+  #ifdef CLOUDY_COOL
     cool = Cloudy_cool(n, T, coolTexObj, heatTexObj);
-    #else
+  #else
     cool = CIE_cool(n, T);
-    #endif
+  #endif
 
     // calculate change in temperature given dt
     del_T = cool * dt * TIME_UNIT * (gamma - 1.0) / (n * KB);
@@ -129,12 +128,12 @@ __global__ void cooling_kernel(Real *dev_conserved, int nx, int ny, int nz, int
       T -= cool * dt_sub * TIME_UNIT * (gamma - 1.0) / (n * KB);
       // how much time is left from the original timestep?
       dt -= dt_sub;
-    // calculate cooling again
-    #ifdef CLOUDY_COOL
+  // calculate cooling again
+  #ifdef CLOUDY_COOL
       cool = Cloudy_cool(n, T, coolTexObj, heatTexObj);
-    #else
+  #else
       cool = CIE_cool(n, T);
-    #endif
+  #endif
       // calculate new change in temperature
       del_T = cool * dt * TIME_UNIT * (gamma - 1.0) / (n * KB);
     }
@@ -145,23 +144,23 @@ __global__ void cooling_kernel(Real *dev_conserved, int nx, int ny, int nz, int
     // adjust value of energy based on total change in temperature
     del_T = T_init - T;  // total change in T
     E -= n * KB * del_T / ((gamma - 1.0) * ENERGY_UNIT);
-    #ifdef DE
+  #ifdef DE
     ge -= KB * del_T / (mu * MP * (gamma - 1.0) * SP_ENERGY_UNIT);
-    #endif
+  #endif
 
-    // calculate cooling rate for new T
-    #ifdef CLOUDY_COOL
+  // calculate cooling rate for new T
+  #ifdef CLOUDY_COOL
     cool = Cloudy_cool(n, T, coolTexObj, heatTexObj);
-    #else
+  #else
     cool = CIE_cool(n, T);
-    // printf("%d %d %d %e %e %e\n", xid, yid, zid, n, T, cool);
-    #endif
+  // printf("%d %d %d %e %e %e\n", xid, yid, zid, n, T, cool);
+  #endif
 
     // and send back from kernel
     dev_conserved[4 * n_cells + id] = E;
-    #ifdef DE
+  #ifdef DE
     dev_conserved[(n_fields - 1) * n_cells + id] = d * ge;
-    #endif
+  #endif
   }
 }
 
@@ -317,7 +316,7 @@ __device__ Real CIE_cool(Real n, Real T)
   return cool;
 }
 
-    #ifdef CLOUDY_COOL
+  #ifdef CLOUDY_COOL
 /* \fn __device__ Real Cloudy_cool(Real n, Real T, cudaTextureObject_t
  coolTexObj, cudaTextureObject_t heatTexObj)
  * \brief Uses texture mapping to interpolate Cloudy cooling/heating
@@ -353,7 +352,6 @@ __device__ Real Cloudy_cool(Real n, Real T, cudaTextureObject_t coolTexObj, cuda
   // printf("DEBUG Cloudy L350: %.17e\n",cool);
   return cool;
 }
-    #endif  // CLOUDY_COOL
+  #endif  // CLOUDY_COOL
 
-  #endif  // COOLING_GPU
-#endif    // CUDA
+#endif  // COOLING_GPU

src/cooling/cooling_cuda.h

@@ -1,14 +1,13 @@
 /*! \file cooling_cuda.h
  *  \brief Declarations of cooling functions. */
 
-#ifdef CUDA
-  #ifdef COOLING_GPU
-    #pragma once
+#ifdef COOLING_GPU
+  #pragma once
 
-    #include <math.h>
+  #include <math.h>
 
-    #include "../global/global.h"
-    #include "../utils/gpu.hpp"
+  #include "../global/global.h"
+  #include "../utils/gpu.hpp"
 
 extern cudaTextureObject_t coolTexObj;
 extern cudaTextureObject_t heatTexObj;
@@ -48,5 +47,4 @@ __device__ Real CIE_cool(Real n, Real T);
           tables at z = 0 with solar metallicity and an HM05 UV background. */
 __device__ Real Cloudy_cool(Real n, Real T, cudaTextureObject_t coolTexObj, cudaTextureObject_t heatTexObj);
 
-  #endif  // COOLING_GPU
-#endif    // CUDA
+#endif  // COOLING_GPU

src/cooling/load_cloudy_texture.cu

@@ -1,18 +1,17 @@
 /*! \file load_cloudy_texture.cu
  *  \brief Wrapper file to load cloudy cooling table as CUDA texture. */
 
-#ifdef CUDA
-  #ifdef CLOUDY_COOL
+#ifdef CLOUDY_COOL
 
-    #include <stdio.h>
-    #include <stdlib.h>
+  #include <stdio.h>
+  #include <stdlib.h>
 
-    #include "../cooling/cooling_cuda.h"
-    #include "../cooling/load_cloudy_texture.h"
-    #include "../cooling/texture_utilities.h"
-    #include "../global/global.h"
-    #include "../global/global_cuda.h"
-    #include "../io/io.h"  // provides chprintf
+  #include "../cooling/cooling_cuda.h"
+  #include "../cooling/load_cloudy_texture.h"
+  #include "../cooling/texture_utilities.h"
+  #include "../global/global.h"
+  #include "../global/global_cuda.h"
+  #include "../io/io.h"  // provides chprintf
 
 cudaArray *cuCoolArray;
 cudaArray *cuHeatArray;
@@ -284,5 +283,4 @@ void Test_Cloudy_Speed()
   exit(0);
 }
 
-  #endif
-#endif
+#endif  // CLOUDY_COOL

src/cooling/load_cloudy_texture.h

@@ -1,12 +1,11 @@
 /*! \file load_cloudy_texture.h
  *  \brief Wrapper file to load cloudy cooling table as CUDA texture. */
 
-#ifdef CUDA
-  #ifdef CLOUDY_COOL
+#ifdef CLOUDY_COOL
 
-    #pragma once
+  #pragma once
 
-    #include "../global/global.h"
+  #include "../global/global.h"
 
 /* \fn void Load_Cuda_Textures()
  * \brief Load the Cloudy cooling tables into texture memory on the GPU. */
@@ -17,5 +16,4 @@ void Load_Cuda_Textures();
  * arrays. */
 void Free_Cuda_Textures();
 
-  #endif
-#endif
+#endif  // CLOUDY_COOL

src/cooling/texture_utilities.h

@@ -5,13 +5,12 @@
 // would be included into a .cpp file because tex2D is undefined when compiling
 // with gcc.
 
-#ifdef CUDA
-  #pragma once
+#pragma once
 
-  #include <math.h>
+#include <math.h>
 
-  #include "../global/global.h"
-  #include "../utils/gpu.hpp"
+#include "../global/global.h"
+#include "../utils/gpu.hpp"
 
 inline __device__ float lerp(float v0, float v1, float f) { return fma(f, v1, fma(-f, v0, v0)); }
 
@@ -40,5 +39,3 @@ inline __device__ float Bilinear_Texture(cudaTextureObject_t tex, float x, float
   // The outer lerp interpolates along y
   return lerp(lerp(t00, t10, fx), lerp(t01, t11, fx), fy);
 }
-
-#endif  // CUDA

src/global/global_cuda.cu

@@ -1,9 +1,7 @@
 /*! \file global_cuda.cu
  *  \brief Declarations of the cuda global variables. */
 
-#ifdef CUDA
-
-  #include "../global/global.h"
+#include "../global/global.h"
 
 // Declare global variables
 bool memory_allocated;
@@ -16,5 +14,3 @@ Real *eta_x, *eta_y, *eta_z, *etah_x, *etah_y, *etah_z;
 Real *dev_grav_potential;
 Real *temp_potential;
 Real *buffer_potential;
-
-#endif  // CUDA