From 3b212bb9400c716d0aeb2d3c1cf3fe8696b41544 Mon Sep 17 00:00:00 2001 From: Gregor Weiss Date: Thu, 25 Apr 2024 07:15:01 +0000 Subject: [PATCH] hip-hacks cmdl tool --- include/library/private/saxpy.h | 8 ++--- src/library/dwt.c | 5 ++- src/tools/CMakeLists.txt | 17 ++++++++-- src/tools/bwccmdl.c | 3 ++ src/tools/bwccmdl.hip | 60 +++++++++++++++++++++++++++++++++ 5 files changed, 86 insertions(+), 7 deletions(-) create mode 100644 src/tools/bwccmdl.hip diff --git a/include/library/private/saxpy.h b/include/library/private/saxpy.h index 975d791..28b6133 100644 --- a/include/library/private/saxpy.h +++ b/include/library/private/saxpy.h @@ -1,10 +1,10 @@ -#ifdef __cpluplus +//#ifdef __cpluplus extern "C" { -#endif +//#endif void bwc_saxpy(); -#ifdef __cpluplus +//#ifdef __cpluplus } -#endif +//#endif diff --git a/src/library/dwt.c b/src/library/dwt.c index 776fc89..1113c95 100755 --- a/src/library/dwt.c +++ b/src/library/dwt.c @@ -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(); } /*----------------------------------------------------------------------------------------------------------*\ diff --git a/src/tools/CMakeLists.txt b/src/tools/CMakeLists.txt index adf9154..1a3c578 100755 --- a/src/tools/CMakeLists.txt +++ b/src/tools/CMakeLists.txt @@ -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) \ No newline at end of file +target_link_libraries(bwccmd PRIVATE bwclib m) +target_link_libraries(bwccmdhip PRIVATE bwclib m) diff --git a/src/tools/bwccmdl.c b/src/tools/bwccmdl.c index 31dbb06..50afdc0 100644 --- a/src/tools/bwccmdl.c +++ b/src/tools/bwccmdl.c @@ -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. ! diff --git a/src/tools/bwccmdl.hip b/src/tools/bwccmdl.hip new file mode 100644 index 0000000..3fbec5d --- /dev/null +++ b/src/tools/bwccmdl.hip @@ -0,0 +1,60 @@ +#include +#include + +#include "bwccmdl.h" + +#include +#include + +__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<<>>(n, d_x, 1, d_y, 1); + hipDeviceSynchronize(); + hipMemcpy(h_y, d_y, size, hipMemcpyDeviceToHost); + + std::cout << "Done saxpy\n"; + return 0; +}