updates compilation for cuda/hip code

cenjinzhiyu · Apr 13, 2021 · 72deb8c · 72deb8c
1 parent a992fac
commit 72deb8c
Show file tree

Hide file tree

Showing 14 changed files with 46 additions and 72 deletions.
diff --git a/Makefile.in b/Makefile.in
@@ -263,7 +263,7 @@ GENCODE_AMD_MI100 = --amdgpu-target=gfx908
 @COND_HIP_TRUE@@COND_HIP_PLATFORM_AMD_TRUE@HIP_CFLAG_ENDING = -x hip
 # NVIDIA default Tesla
 @COND_HIP_TRUE@@COND_HIP_PLATFORM_NVIDIA_TRUE@GENCODE_HIP = $(GENCODE_30)
-@COND_HIP_TRUE@@COND_HIP_PLATFORM_AMD_TRUE@HIP_CFLAG_ENDING = -x cu
+@COND_HIP_TRUE@@COND_HIP_PLATFORM_NVIDIA_TRUE@HIP_CFLAG_ENDING =      # no need for ending
 
 # specific targets
 @COND_HIP_TRUE@@COND_HIP_MI8_TRUE@GENCODE_HIP = $(GENCODE_AMD_MI8)    # --with-hip=MI8 ..

diff --git a/configure b/configure
@@ -630,13 +630,13 @@ VTK_INCLUDES
 VTK_MAJOR
 OMP_LIB
 OMP_FCFLAGS
-HIP_LIBS
-HIP_LDFLAGS
-HIP_CPPFLAGS
 COND_HIP_PLATFORM_NVIDIA_FALSE
 COND_HIP_PLATFORM_NVIDIA_TRUE
 COND_HIP_PLATFORM_AMD_FALSE
 COND_HIP_PLATFORM_AMD_TRUE
+HIP_LIBS
+HIP_LDFLAGS
+HIP_CPPFLAGS
 HIPCONFIG_PROG
 HIPCC_PROG
 HIP_LIB
@@ -8488,22 +8488,6 @@ $as_echo "$as_me: HIP_FLAGS: ${HIP_FLAGS}" >&6;} ;;
       GPU_PLATFORM=amd
 
 fi
-     if test x"$GPU_PLATFORM" = xamd; then
-  COND_HIP_PLATFORM_AMD_TRUE=
-  COND_HIP_PLATFORM_AMD_FALSE='#'
-else
-  COND_HIP_PLATFORM_AMD_TRUE='#'
-  COND_HIP_PLATFORM_AMD_FALSE=
-fi
-
-     if test x"$GPU_PLATFORM" = xnvidia; then
-  COND_HIP_PLATFORM_NVIDIA_TRUE=
-  COND_HIP_PLATFORM_NVIDIA_FALSE='#'
-else
-  COND_HIP_PLATFORM_NVIDIA_TRUE='#'
-  COND_HIP_PLATFORM_NVIDIA_FALSE=
-fi
-
 
     # for compilation errors like:
     #   /usr/bin/ld: obj/assemble_MPI_scalar_gpu.hip.o: relocation R_X86_64_32 against `.rodata.str1.1' can not be used
@@ -8699,6 +8683,22 @@ ac_compiler_gnu=$ac_cv_fc_compiler_gnu
 
 
 fi
+ if test x"$GPU_PLATFORM" = xamd; then
+  COND_HIP_PLATFORM_AMD_TRUE=
+  COND_HIP_PLATFORM_AMD_FALSE='#'
+else
+  COND_HIP_PLATFORM_AMD_TRUE='#'
+  COND_HIP_PLATFORM_AMD_FALSE=
+fi
+
+ if test x"$GPU_PLATFORM" = xnvidia; then
+  COND_HIP_PLATFORM_NVIDIA_TRUE=
+  COND_HIP_PLATFORM_NVIDIA_FALSE='#'
+else
+  COND_HIP_PLATFORM_NVIDIA_TRUE='#'
+  COND_HIP_PLATFORM_NVIDIA_FALSE=
+fi
+
 
 
 ###

diff --git a/configure.ac b/configure.ac
@@ -562,8 +562,6 @@ AS_IF([test x"$want_hip" != xno], [
     AS_IF([test x"$GPU_PLATFORM" = x],[
       GPU_PLATFORM=amd
     ])
-    AM_CONDITIONAL([COND_HIP_PLATFORM_AMD], [test x"$GPU_PLATFORM" = xamd])
-    AM_CONDITIONAL([COND_HIP_PLATFORM_NVIDIA], [test x"$GPU_PLATFORM" = xnvidia])
 
     # for compilation errors like:
     #   /usr/bin/ld: obj/assemble_MPI_scalar_gpu.hip.o: relocation R_X86_64_32 against `.rodata.str1.1' can not be used
@@ -661,6 +659,8 @@ AS_IF([test x"$want_hip" != xno], [
     AC_SUBST([HIP_LDFLAGS])
     AC_SUBST([HIP_LIBS])
 ])
+AM_CONDITIONAL([COND_HIP_PLATFORM_AMD], [test x"$GPU_PLATFORM" = xamd])
+AM_CONDITIONAL([COND_HIP_PLATFORM_NVIDIA], [test x"$GPU_PLATFORM" = xnvidia])
 
 
 ###

diff --git a/src/gpu/kernels/compute_gradient_kernel.cpp b/src/gpu/kernels/compute_gradient_kernel.cpp
diff --git a/src/gpu/kernels/compute_gradient_kernel.cu → src/gpu/kernels/compute_gradient_kernel.h b/src/gpu/kernels/compute_gradient_kernel.cu → src/gpu/kernels/compute_gradient_kernel.h
@@ -27,9 +27,12 @@
 !=====================================================================
 */
 
+#ifndef COMPUTE_GRADIENT_GPU_H
+#define COMPUTE_GRADIENT_GPU_H
+
 // needed in compute_kernels_acoustic_kernel.cu and compute_kernels_hess_ac_cudakernel.cu
 
-__device__ void compute_gradient_kernel(int ijk,
+__device__ __forceinline__ void compute_gradient_kernel(int ijk,
                                         int ispec,int ispec_irreg,
                                         field* scalar_field,
                                         field* vector_field_loc,
@@ -111,3 +114,4 @@ __device__ void compute_gradient_kernel(int ijk,
 }
 
 
+#endif   // COMPUTE_GRADIENT_GPU_H
diff --git a/src/gpu/kernels/compute_kernels_acoustic_kernel.cpp b/src/gpu/kernels/compute_kernels_acoustic_kernel.cpp
@@ -28,6 +28,6 @@
 */
 
 // includes device function compute_gradient_kernel()
-#include "compute_gradient_kernel.cu"
+#include "compute_gradient_kernel.h"
 
 #include "compute_kernels_acoustic_kernel.cu"
diff --git a/src/gpu/kernels/compute_kernels_acoustic_kernel.cu b/src/gpu/kernels/compute_kernels_acoustic_kernel.cu
@@ -27,6 +27,9 @@
 !=====================================================================
 */
 
+// includes device function compute_gradient_kernel()
+#include "compute_gradient_kernel.h"
+
 
 __global__ void compute_kernels_acoustic_kernel(int* ispec_is_acoustic,
                                                 int* d_ibool,

diff --git a/src/gpu/kernels/compute_kernels_hess_ac_cudakernel.cpp b/src/gpu/kernels/compute_kernels_hess_ac_cudakernel.cpp
@@ -28,6 +28,6 @@
 */
 
 // includes device function compute_gradient_kernel()
-#include "compute_gradient_kernel.cu"
+#include "compute_gradient_kernel.h"
 
 #include "compute_kernels_hess_ac_cudakernel.cu"
diff --git a/src/gpu/kernels/compute_kernels_hess_ac_cudakernel.cu b/src/gpu/kernels/compute_kernels_hess_ac_cudakernel.cu
@@ -27,6 +27,9 @@
 !=====================================================================
 */
 
+// includes device function compute_gradient_kernel()
+#include "compute_gradient_kernel.h"
+
 
 __global__ void compute_kernels_hess_ac_cudakernel(int* ispec_is_acoustic,
                                                    int* d_ibool,

diff --git a/src/gpu/kernels/kernel_cuda.mk b/src/gpu/kernels/kernel_cuda.mk
@@ -13,7 +13,6 @@ cuda_kernels_OBJS := \
 	$O/compute_dynamic_fault_cuda.cuda-kernel.o \
 	$O/compute_elastic_seismogram_kernel.cuda-kernel.o \
 	$O/compute_element_strain_cudakernel.cuda-kernel.o \
-	$O/compute_gradient_kernel.cuda-kernel.o \
 	$O/compute_kernels_acoustic_kernel.cuda-kernel.o \
 	$O/compute_kernels_ani_cudakernel.cuda-kernel.o \
 	$O/compute_kernels_cudakernel.cuda-kernel.o \

diff --git a/src/gpu/kernels/kernel_proto.cu.h b/src/gpu/kernels/kernel_proto.cu.h
@@ -603,11 +603,6 @@ __global__ void compute_element_strain_cudakernel(int* ispec_is_elastic,
                                                   int NSPEC_AB) ;
 
 
-//
-// src/gpu/kernels/compute_gradient_kernel.cu
-//
-
-
 //
 // src/gpu/kernels/compute_kernels_acoustic_kernel.cu
 //

diff --git a/src/gpu/prepare_constants_cuda.h b/src/gpu/prepare_constants_cuda.h
@@ -110,15 +110,15 @@ __device__ __constant__ realw d_wgll_cube[NGLL3]; // needed only for gravity cas
 // If not added dummy kernel, optimizer will delete constant variables,
 // because they have not used in any of the kernel.
 __global__ void dummy_kernel(){
-        d_hprime_xx[0]=1;
-//      d_hprime_yy[0]=1;
-//      d_hprime_zz[0]=1;
-        d_hprimewgll_xx[0]=1;
-//      d_hprimewgll_yy[0]=1;
-//      d_hprimewgll_zz[0]=1;
-        d_wgllwgll_xy[0]=1;
-        d_wgllwgll_xz[0]=1;
-        d_wgllwgll_yz[0]=1;
+  d_hprime_xx[0]=1;
+  //d_hprime_yy[0]=1;
+  //d_hprime_zz[0]=1;
+  d_hprimewgll_xx[0]=1;
+  //d_hprimewgll_yy[0]=1;
+  //d_hprimewgll_zz[0]=1;
+  d_wgllwgll_xy[0]=1;
+  d_wgllwgll_xz[0]=1;
+  d_wgllwgll_yz[0]=1;
 }
 #endif
 

diff --git a/src/gpu/rules.mk b/src/gpu/rules.mk
@@ -144,7 +144,7 @@ ifeq ($(CUDA),yes)
 endif
 
 ## HIP compilation
-HIPCC_CFLAGS := ${HIP_CFLAGS} ${HIP_CFLAG_ENDING}  # adding either -x hip or -x cu depending on platform
+HIPCC_CFLAGS := ${HIP_CFLAGS} ${HIP_CFLAG_ENDING}  # adding -x hip depending on platform
 ifeq ($(HIP), yes)
   BUILD_VERSION_TXT += HIP
   SELECTOR_CFLAG += $(FC_DEFINE)USE_HIP

diff --git a/utils/create_specfem3D_gpu_cuda_method_stubs.pl b/utils/create_specfem3D_gpu_cuda_method_stubs.pl
@@ -57,7 +57,7 @@ END
 
 
 $warning = <<END;
- fprintf(stderr,"ERROR: GPU_MODE enabled without GPU/CUDA Support. To enable GPU support, reconfigure with --with-cuda flag.\\n");
+ fprintf(stderr,"ERROR: GPU_MODE enabled without GPU/CUDA/HIP Support. To enable GPU support, reconfigure with --with-cuda or --with-hip flag.\\n");
  exit(1);
 END
-Original file line number
+Diff line change
@@ Expand Up @@
                                                       int NSPEC_AB) ;
-    //
-    // src/gpu/kernels/compute_gradient_kernel.cu
-    //
     //
     // src/gpu/kernels/compute_kernels_acoustic_kernel.cu
     //
@@ Expand Down @@