Commit 207b5404 authored by Berrill, Mark's avatar Berrill, Mark
Browse files

Fixing bugs with HIP

parent 0230b135

src/CMakeLists.txt

+6 −6
Original line number Diff line number Diff line
@@ -200,18 +200,18 @@ ENDIF() 
IF ( USE_KOKKOS )

    #SET( SOURCES ${SOURCES} RayTraceImageKokkos.cpp )

ENDIF()

IF (USE_HIP )

    SET( SOURCES ${SOURCES} RayTraceImageHIP.cpp )

ENDIF()

#IF (USE_HIP )

#    SET( SOURCES ${SOURCES} RayTraceImageHIP.cpp )

#ENDIF()

IF ( USE_CUDA )

    SET( SOURCES ${SOURCES} RayTraceImageCuda.cu )

ENDIF()

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()

src/RayTraceImageHIP.hip→src/RayTraceImageHIP.cu

+2 −2
Original line number Diff line number Diff line
@@ -109,7 +109,7 @@ __global__ __launch_bounds__( 128, 8 ) // Set bounds to limit the number of regi 
                // Note: The sign of the angle is reversed with respect to the euv_beam

                ray2.a = -ray2.a;

                ray2.b = -ray2.b;

                if ( ray2.y < 0.0 && y[0] >= 0.0 ) {

                if ( ray2.y < 0.0 && y_range[0] >= 0.0 ) {

                    // We need to change the sign of y

                    ray2.y = -ray2.y;

                }
@@ -169,7 +169,7 @@ void RayTraceImageHIPLoop( int N, const RayTrace::EUV_beam_struct& beam, 
            RAY_WARNING("maxThreadsPerBlock=0 - Setting to default 128");

            maxThreadsPerBlock = 128;

        }

    }

    }*/

    // place the ray gain and seed structures on the device

    const int nx = beam.nx;

    const int ny = beam.ny;

src/RayTraceStructures.cpp

+16 −3
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
@@ -1375,6 +1383,8 @@ const RayTrace::ray_seed_struct *RayTrace::ray_seed_struct::copy_device() const 
    return dev_instance;

#elif defined( USE_CUDA )

    return ray_seed_struct_copy_device_cuda( *this );

#elif defined( USE_HIP )

    return ray_seed_struct_copy_device_hip( *this );

#elif defined( USE_KOKKOS )

    RAY_ERROR( "Not Finished" );

    return NULL;
@@ -1405,9 +1415,10 @@ void RayTrace::ray_seed_struct::free_device( 
    // now we can delete the structure

    acc_delete( instance, sizeof( *host_seed ) );





#elif defined( USE_CUDA )

    ray_seed_struct_free_device_cuda( device_seed );

#elif defined( USE_HIP )

    ray_seed_struct_free_device_hip( device_seed );

#elif defined( USE_KOKKOS )

    RAY_ERROR( "Not Finished" );

#else
@@ -1980,6 +1991,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 =
@@ -2017,9 +2030,9 @@ void RayTrace::ray_gain_struct::free_device( 
{

    if ( device_arr == nullptr )

        return;

#if defined( USE_CUDA )

#if defined( USE_HIP )

    NULL_USE( host_arr );

    ray_gain_struct_free_device_cuda( N, device_arr );

    ray_gain_struct_free_device_hip( N, device_arr );

#elif defined( USE_OPENACC ) && defined( _OPENACC )



    ray_gain_struct *instances =