Commit 82f2313c authored by Berrill, Mark's avatar Berrill, Mark
Browse files

Fixing minor bug with HIP

parent 883d0b92

src/RayTraceImageCuda.cu

+2 −2
Original line number Diff line number Diff line
@@ -184,8 +184,8 @@ void RayTraceImageCudaLoop( int N, const RayTrace::EUV_beam_struct& beam, 
    const double dz = beam.dz;

    const double da = beam.da;

    const double db = beam.db;

    const RayTrace::ray_gain_struct* gain = RayTrace::ray_gain_struct::copy_device( N, gain_in );

    const RayTrace::ray_seed_struct* seed = nullptr;

    const auto gain = RayTrace::ray_gain_struct::copy_device( N, gain_in );

    const auto seed = nullptr;

    if ( seed_in!=NULL )

        seed = seed_in->copy_device();

    // Allocate device memory

src/RayTraceImageHIP.cu

+9 −6
Original line number Diff line number Diff line
@@ -7,6 +7,9 @@ 
#include "hip/hip_runtime.h"





#define HIP_CHECK()  do {} while (0)





// Atomic add operation for double

__device__ double atomicAdd2( double* address, double val )

{
@@ -134,8 +137,8 @@ void RayTraceImageHIPLoop( int N, const RayTrace::EUV_beam_struct& beam, 
    const double dz = beam.dz;

    const double da = beam.da;

    const double db = beam.db;

    const RayTrace::ray_gain_struct* gain = RayTrace::ray_gain_struct::copy_device( N, gain_in );

    const RayTrace::ray_seed_struct* seed = nullptr;

    const auto gain = RayTrace::ray_gain_struct::copy_device( N, gain_in );

    const auto seed = nullptr;

    if ( seed_in!=NULL )

        seed = seed_in->copy_device();

    // Allocate device memory
@@ -180,7 +183,7 @@ void RayTraceImageHIPLoop( int N, const RayTrace::EUV_beam_struct& beam, 
// Copy ray_gain_struct to GPU

const RayTrace::ray_gain_struct* ray_gain_struct_copy_device_hip( size_t N, const RayTrace::ray_gain_struct* arr )

{

    RayTrace::ray_gain_struct* host_ptr = new RayTrace::ray_gain_struct[N];

    auto host_ptr = new RayTrace::ray_gain_struct[N];

    for (size_t i=0; i<N; i++) {

        host_ptr[i].Nx = arr[i].Nx;

        host_ptr[i].Ny = arr[i].Ny;
@@ -221,7 +224,7 @@ const RayTrace::ray_gain_struct* ray_gain_struct_copy_device_hip( size_t N, cons 
// Free ray_gain_struct from GPU

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];

    auto host_ptr = new RayTrace::ray_gain_struct[N];

    hipMemcpy(host_ptr,dev_ptr,N*sizeof(RayTrace::ray_gain_struct),hipMemcpyDeviceToHost);

    for (size_t i=0; i<N; i++) {

        hipFree(host_ptr[i].x);
@@ -248,7 +251,7 @@ void ray_gain_struct_free_device_hip( size_t N, const RayTrace::ray_gain_struct* 
// Copy ray_seed_struct to GPU

const RayTrace::ray_seed_struct* ray_seed_struct_copy_device_hip( const RayTrace::ray_seed_struct& seed )

{

    RayTrace::ray_seed_struct* host_ptr = new RayTrace::ray_seed_struct();

    auto host_ptr = new RayTrace::ray_seed_struct();

    host_ptr->f0 = seed.f0;

    for (size_t i=0; i<5; i++) {

        host_ptr->dim[i] = seed.dim[i];
@@ -271,7 +274,7 @@ const RayTrace::ray_seed_struct* ray_seed_struct_copy_device_hip( const RayTrace 
// Free ray_seed_struct from GPU

void ray_seed_struct_free_device_hip( const RayTrace::ray_seed_struct* dev_ptr )

{

    RayTrace::ray_seed_struct* host_ptr = new RayTrace::ray_seed_struct;

    auto host_ptr = new RayTrace::ray_seed_struct;

    hipMemcpy(host_ptr,dev_ptr,sizeof(RayTrace::ray_seed_struct),hipMemcpyDeviceToHost);

    for (size_t i=0; i<5; i++) {

        hipFree(host_ptr->x[i]);