TOPIO/BigWhoop
TOPIO/BigWhoop
4
1
Fork 0
Code Issues 26 Pull requests 1 Projects 6 Releases Packages Wiki Activity

hip-hacks cmdl tool

Browse source
This commit is contained in:
Gregor Weiss 2024-04-25 07:15:01 +00:00
parent bdff484f1e
commit 3b212bb940
5 changed files with 86 additions and 7 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
include/library/private/saxpy.h
View file

@ -1,10 +1,10 @@
#ifdef __cpluplus
//#ifdef __cpluplus
extern "C" {
#endif
//#endif
void bwc_saxpy();
#ifdef __cpluplus
//#ifdef __cpluplus
}
#endif
//#endif

5
src/library/dwt.c
View file

@ -484,6 +484,8 @@ whole_point_symmetric_extend(bwc_sample *const working_buffer, uint64 res0, uint
! 25.06.2018 Patrick Vogler B87D120 V 0.1.0 function created !
! !
\*----------------------------------------------------------------------------------------------------------*/
// INCLUDE HIP SAXPY
//#include "saxpy.h"
static void
forward_9x7_CDF_wavelet_transform(bwc_sample *const working_buffer, uint64 res0, uint64 res1)
{
@ -536,7 +538,8 @@ forward_9x7_CDF_wavelet_transform(bwc_sample *const working_buffer, uint64 res0,
working_buffer[i].f = KAPPA_L * working_buffer[i].f;
}
//do_saxpy();
// INCLUDE HIP SAXPY
//bwc_saxpy();
}
/*----------------------------------------------------------------------------------------------------------*\

17
src/tools/CMakeLists.txt
View file

@ -48,6 +48,10 @@
add_executable(bwccmd bwccmdl.c
../interfaces/reader/eas3.c)
add_executable(bwccmdhip bwccmdl.hip
../interfaces/reader/eas3.c)
#----------------------------------------------------------#
# Set the target compile definition for the requested file #
# format support. #
@ -56,16 +60,18 @@ MESSAGE(STATUS "EAS3 file format support: ${BUILD_EAS3}")
if(${BUILD_EAS3})
target_compile_definitions(bwccmd PRIVATE -DBWC_EAS3)
target_compile_definitions(bwccmdhip PRIVATE -DBWC_EAS3)
endif()
if(${BUILD_NETCDF})
target_compile_definitions(bwccmd PRIVATE -DBWC_NETCDF)
target_compile_definitions(bwccmdhip PRIVATE -DBWC_NETCDF)
endif()
#----------------------------------------------------------#
# Define the output name for the utility binaries. #
#----------------------------------------------------------#
set_property(TARGET bwccmd PROPERTY OUTPUT_NAME bwc)
set_property(TARGET bwccmdhip PROPERTY OUTPUT_NAME bwchip)
#----------------------------------------------------------#
# Setup up the include directory for the bwc utilities. #
@ -74,12 +80,19 @@ target_include_directories(bwccmd PRIVATE ${CMAKE_SOURCE_DIR}/include/tools)
target_include_directories(bwccmd PRIVATE ${CMAKE_SOURCE_DIR}/include/library/public)
target_include_directories(bwccmd PRIVATE ${CMAKE_SOURCE_DIR}/include/interfaces/reader)
target_include_directories(bwccmdhip PRIVATE ${CMAKE_SOURCE_DIR}/include/tools)
target_include_directories(bwccmdhip PRIVATE ${CMAKE_SOURCE_DIR}/include/library/public)
target_include_directories(bwccmdhip PRIVATE ${CMAKE_SOURCE_DIR}/include/interfaces/reader)
#----------------------------------------------------------#
# Setup the install directories. #
#----------------------------------------------------------#
install(TARGETS bwccmd DESTINATION ${CMAKE_INSTALL_BINDIR})
install(TARGETS bwccmdhip DESTINATION ${CMAKE_INSTALL_BINDIR})
#----------------------------------------------------------#
# Link the bwc utility to the bwc library. #
#----------------------------------------------------------#
target_link_libraries(bwccmd PRIVATE bwclib m)
target_link_libraries(bwccmd PRIVATE bwclib m)
target_link_libraries(bwccmdhip PRIVATE bwclib m)

3
src/tools/bwccmdl.c
View file

@ -802,6 +802,9 @@ parse_arguments(int argc,
}
args->root = args;
// INCLUDE HIP SAXPY
bwc_saxpy();
/*--------------------------------------------------------*\
! Walk through all the command-line arguments passed to !
! main. !

60
src/tools/bwccmdl.hip Normal file
View file

@ -0,0 +1,60 @@
#include <bwc.h>
#include <stdlib.h>
#include "bwccmdl.h"
#include <hip/hip_runtime.h>
#include <iostream>
__constant__ float a = 1.0f;
void init (int n, float *x, float *y)
{
for (std::size_t i = 0; i < n; ++i)
{
x[i] = 1.0;
y[i] = 0.0;
}
}
__global__
void saxpy (int n, float const* x, int incx, float* y, int incy)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < n)
y[i] += a*x[i];
}
int
main(int argc,
char *argv[])
{
// INCLUDE HIP SAXPY
bwc_saxpy();
int n = 256;
std::size_t size = sizeof(float)*n;
std::cout << "Doing saxpy\n";
float *h_x = new float [n];
float *h_y = new float [n];
init(n, h_x, h_y);
float* d_x;
float *d_y;
hipMalloc(&d_x, size);
hipMalloc(&d_y, size);
hipMemcpy(d_x, h_x, size, hipMemcpyHostToDevice);
hipMemcpy(d_y, h_y, size, hipMemcpyHostToDevice);
int num_groups = 2;
int group_size = 128;
saxpy<<<num_groups, group_size>>>(n, d_x, 1, d_y, 1);
hipDeviceSynchronize();
hipMemcpy(h_y, d_y, size, hipMemcpyDeviceToHost);
std::cout << "Done saxpy\n";
return 0;
}