diff --git a/Makefile b/Makefile index 9d302d693..fa2bc4500 100644 --- a/Makefile +++ b/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 diff --git a/builds/make.host.frontier b/builds/make.host.frontier index c225b3655..69f715871 100644 --- a/builds/make.host.frontier +++ b/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) diff --git a/src/dust/dust_cuda.cu b/src/dust/dust_cuda.cu index 524b58cd0..50356c3c5 100644 --- a/src/dust/dust_cuda.cu +++ b/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); } diff --git a/src/grid/initial_conditions.cpp b/src/grid/initial_conditions.cpp index 5d98d8367..8c8c5b3f9 100644 --- a/src/grid/initial_conditions.cpp +++ b/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. */ diff --git a/src/hydro/hydro_cuda.cu b/src/hydro/hydro_cuda.cu index add8af96a..3b60f97f9 100644 --- a/src/hydro/hydro_cuda.cu +++ b/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]; } diff --git a/src/utils/reduction_utilities.h b/src/utils/reduction_utilities.h index 9aef9600d..39089ac2e 100644 --- a/src/utils/reduction_utilities.h +++ b/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);