Change a few comments and make dust compile on crusher

Makefile

@@ -134,7 +134,7 @@ ifdef HIPCONFIG
   DFLAGS    += -DO_HIP
   CXXFLAGS  += $(HIPCONFIG)
   GPUCXX    ?= hipcc
-  GPUFLAGS  += -Wall
+  #GPUFLAGS  += -Wall
   LD        := $(CXX)
   LDFLAGS   := $(CXXFLAGS) -L$(ROCM_PATH)/lib
   LIBS      += -lamdhip64

builds/make.host.frontier

@@ -11,8 +11,8 @@ CFLAGS_OPTIMIZE   = -g -O2
 CXXFLAGS_DEBUG    = -g -O0 -std=c++17
 CXXFLAGS_OPTIMIZE = -g -Ofast -std=c++17 -Wno-unused-result
 
-GPUFLAGS_OPTIMIZE = -std=c++17 --offload-arch=gfx90a -Wno-unused-result
-GPUFLAGS_DEBUG    = -g -O0 -std=c++17 --offload-arch=gfx90a -Wno-unused-result
+GPUFLAGS_OPTIMIZE = -std=c++17 --offload-arch=gfx90a -Wall -Wno-unused-result
+GPUFLAGS_DEBUG    = -g -O0 -std=c++17 --offload-arch=gfx90a -Wall -Wno-unused-result
 HIPCONFIG	  = -I$(ROCM_PATH)/include $(shell hipconfig -C) # workaround for Rocm 5.2 warnings
 #HIPCONFIG	  = $(shell hipconfig -C)
 

src/dust/dust_cuda.cu

@@ -111,7 +111,7 @@ __global__ void Dust_Kernel(Real *dev_conserved, int nx, int ny, int nz, int n_g
 }
 
 // McKinnon et al. (2017)
-__device__ Real calc_tau_sp(Real n, Real T) {
+__device__ __host__ Real calc_tau_sp(Real n, Real T) {
   Real YR_IN_S = 3.154e7;
   Real a1 = 1; // dust grain size in units of 0.1 micrometers
   Real d0 = n / (6e-4); // gas density in units of 10^-27 g/cm^3
@@ -125,7 +125,7 @@ __device__ Real calc_tau_sp(Real n, Real T) {
 }
 
 // McKinnon et al. (2017)
-__device__ Real calc_dd_dt(Real d_dust, Real tau_sp) {
+__device__ __host__ Real calc_dd_dt(Real d_dust, Real tau_sp) {
     return -d_dust / (tau_sp/3);
 }
 

src/grid/initial_conditions.cpp

@@ -1,5 +1,5 @@
 /*! \file initial_conditions.cpp
-/*  \brief Definitions of initial conditions for different tests.
+ *  \brief Definitions of initial conditions for different tests.
            Note that the grid is mapped to 1D as i + (x_dim)*j + (x_dim*y_dim)*k.
            Functions are members of the Grid3D class. */
 

src/hydro/hydro_cuda.cu

@@ -604,6 +604,7 @@ Real Calc_dt_GPU(Real *dev_conserved, int nx, int ny, int nz, int n_ghost, int n
   }
   CudaCheckError();
 
+  // Note: dev_dti[0] is DeviceVector syntactic sugar for returning a value via cudaMemcpy
   return dev_dti[0];
 }
 

src/utils/reduction_utilities.h

@@ -267,8 +267,6 @@
          */
         __inline__ __device__ void gridReduceMax(Real val, Real* out)
         {
-            // __syncthreads();  // Wait for all threads to calculate val;
-            // __syncthreads();  // Wait for all threads to calculate val;
 
             // Reduce the entire block in parallel
             val = blockReduceMax(val);