Working on HIP port (928f6bdb) · Commits · Berrill, Mark A / RayTrace-miniapp

build/configure

+5 −2

Original line number	Diff line number	Diff line
		@@ -8,14 +8,17 @@ cmake \
		-D CMAKE_BUILD_TYPE=Release \
		-D CMAKE_CXX_COMPILER=mpic++ \
		-D CMAKE_CXX_STANDARD=11 \
		-D USE_MPI=0 \
		-D USE_OPENACC=0 \
		-D USE_OPENMP=0 \
		-D USE_KOKKOS=0 \
		-D KOKKOS_DIRECTORY=${KOKKOS_DIR} \
		-D KOKKOS_WRAPPER=${KOKKOS_DIR}/nvcc_wrapper \
		-D USE_CUDA=0 \
		-D CUDA_FLAGS="-arch sm_30" \
		-D CUDA_FLAGS="-arch sm_30 -std=c++11" \
		-D USE_HIP=1 \
		-D HIP_NVCC_OPTIONS="-arch sm_30" \
		-D HIP_HIPCC_FLAGS="-std=c++11" \
		-D HIP_HCC_FLAGS="-std=c++11" \
		-D HIP_NVCC_OPTIONS="-arch sm_30 -std=c++11" \
		../src

src/CMakeLists.txt

+13 −9

Original line number	Diff line number	Diff line
		@@ -153,12 +153,14 @@ IF ( USE_HIP )
		MESSAGE( "HIP Found")
		MESSAGE( " HIP version: ${HIP_VERSION_STRING}")
		MESSAGE( " HIP platform: ${HIP_PLATFORM}")
		MESSAGE( " HIP Include Path: ${HIP_INCLUDE_DIRS}")
		MESSAGE( " HIP Include Path: ${HIP_INCLUDE_DIRS} ${HIP_HIPCC_INCLUDE_ARGS}")
		MESSAGE( " HIP Libraries: ${HIP_LIBRARIES}")
		MESSAGE( " CUDA Libraries: ${CUDA_LIBRARIES}")
		ADD_DEFINITIONS( -DUSE_HIP )
		INCLUDE_DIRECTORIES( "/sw/summit/hip/hip2.6-cuda10.1.168/hip/roc-2.6.0/include" )
		ADD_DEFINITIONS( -D__HIP_PLATFORM_NVCC__ )
		ENDIF()


		# Include some paths
		INCLUDE_DIRECTORIES( "${RAYTRACE_INSTALL_DIR}/include" )
		INCLUDE_DIRECTORIES( "${RAYTRACE_SOURCE_DIR}" )
		@@ -195,25 +197,27 @@ IF ( USE_KOKKOS )
		#SET( SOURCES ${SOURCES} RayTraceImageKokkos.cpp )
		ENDIF()
		IF (USE_HIP )
		SET( SOURCES ${SOURCES} RayTraceImageHIP.cpp )
		SET( SOURCES ${SOURCES} RayTraceImageHIP.cu )
		ENDIF()
		IF ( USE_CUDA )
		SET( SOURCES ${SOURCES} RayTraceImageCuda.cu )
		ENDIF()
		ADD_LIBRARY( RayTrace SHARED ${SOURCES} ${CUBINS} )
		HIP_ADD_LIBRARY( RayTrace SHARED ${SOURCES} ${CUBINS} )
		IF ( USE_HIP )
		SET( HIP_SEPERABLE_COMPILATION ON )
		SET_SOURCE_FILES_PROPERTIES( RayTraceImageHIP.hip PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1 )
		HIP_ADD_LIBRARY( RayTraceHIP RayTraceImageHIP.hip SHARED HIPCC_OPTIONS ${HIP_HIPCC_OPTIONS} HCC_OPTIONS ${HIP_HCC_OPTIONS} NVCC_OPTIONS ${HIP_NVCC_OPTIONS} ${HIP_NVCC_FLAGS} )
		TARGET_LINK_LIBRARIES( RayTraceHIP /opt/rocm-3.3.0/lib/libhip_hcc.so )
		SET_SOURCE_FILES_PROPERTIES( RayTraceImageHIP.cu PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1 )
		HIP_ADD_LIBRARY( RayTraceHIP RayTraceImageHIP.cu SHARED HIPCC_OPTIONS ${HIP_HIPCC_OPTIONS} HCC_OPTIONS ${HIP_HCC_OPTIONS} NVCC_OPTIONS ${HIP_NVCC_OPTIONS} ${HIP_NVCC_FLAGS} )
		#TARGET_LINK_LIBRARIES( RayTraceHIP /opt/rocm-3.3.0/lib/libhip_hcc.so )
		TARGET_LINK_LIBRARIES( RayTrace RayTraceHIP )
		ADD_DEPENDENCIES( RayTraceHIP copy-include )
		ENDIF()


		# Add the applications
		ADD_EXECUTABLE( CreateImage CreateImage.cpp )
		TARGET_LINK_LIBRARIES( CreateImage RayTrace RayTraceHIP ${KOKKOS_LIB} ${TIMER_LIBS} ${LDFLAGS} ${LDLIBS} )
		IF ( USE_HIP )
		HIP_ADD_EXECUTABLE( CreateImage CreateImage.cpp )
		ENDIF()
		TARGET_LINK_LIBRARIES( CreateImage RayTrace RayTraceHIP ${KOKKOS_LIB} ${TIMER_LIBS} ${LDFLAGS} ${LDLIBS} ${CUDA_LIBRARIES} )
		INSTALL( TARGETS CreateImage DESTINATION ${RAYTRACE_INSTALL_DIR}/bin )

src/RayTraceImageHIP.hip→src/RayTraceImageHIP.cu

+1 −1

Original line number	Diff line number	Diff line
		@@ -276,7 +276,7 @@ const RayTrace::ray_gain_struct* ray_gain_struct_copy_device_hip( size_t N, cons
		return dev_ptr;
		}
		// Free ray_gain_struct from GPU
		void ray_gain_struct_free_device_cuda( size_t N, const RayTrace::ray_gain_struct* dev_ptr )
		void ray_gain_struct_free_device_hip( size_t N, const RayTrace::ray_gain_struct* dev_ptr )
		{
		RayTrace::ray_gain_struct* host_ptr = new RayTrace::ray_gain_struct[N];
		hipMemcpy(host_ptr,dev_ptr,N*sizeof(RayTrace::ray_gain_struct),hipMemcpyDeviceToHost);

src/RayTraceStructures.cpp

+13 −0

Original line number	Diff line number	Diff line
		@@ -25,6 +25,14 @@ extern const RayTrace::ray_seed_struct *ray_seed_struct_copy_device_cuda(
		extern void ray_gain_struct_free_device_cuda( size_t N, const RayTrace::ray_gain_struct *arr );
		extern void ray_seed_struct_free_device_cuda( const RayTrace::ray_seed_struct *dev_ptr );
		#endif
		#ifdef USE_HIP
		extern const RayTrace::ray_gain_struct *ray_gain_struct_copy_device_hip(
		size_t N, const RayTrace::ray_gain_struct *arr );
		extern const RayTrace::ray_seed_struct *ray_seed_struct_copy_device_hip(
		const RayTrace::ray_seed_struct &seed );
		extern void ray_gain_struct_free_device_hip( size_t N, const RayTrace::ray_gain_struct *arr );
		extern void ray_seed_struct_free_device_hip( const RayTrace::ray_seed_struct *dev_ptr );
		#endif


		// Atomic model headers
		@@ -1988,6 +1996,8 @@ const RayTrace::ray_gain_struct *RayTrace::ray_gain_struct::copy_device(
		{
		#if defined( USE_CUDA )
		return ray_gain_struct_copy_device_cuda( N, arr );
		#elif defined( USE_HIP )
		return ray_gain_struct_copy_device_hip( N, arr );
		#elif defined( USE_OPENACC ) && defined( _OPENACC )
		ray_gain_struct instances = const_cast<ray_gain_struct >( arr ); // bad API const-ness spec.
		ray_gain_struct *dev_instances =
		@@ -2028,6 +2038,9 @@ void RayTrace::ray_gain_struct::free_device(
		#if defined( USE_CUDA )
		NULL_USE( host_arr );
		ray_gain_struct_free_device_cuda( N, device_arr );
		#elif defined( USE_HIP )
		NULL_USE( host_arr );
		ray_gain_struct_free_device_hip( N, device_arr );
		#elif defined( USE_OPENACC ) && defined( _OPENACC )

		ray_gain_struct *instances =