From bdff484f1e6eead86160dcf1825b1b4c59fbd860 Mon Sep 17 00:00:00 2001
From: Gregor Weiss <gregor.weiss@hlrs.de>
Date: Thu, 25 Apr 2024 07:14:32 +0000
Subject: [PATCH] hip saxpy

---
 include/library/private/dwt.h   |  2 +-
 include/library/private/saxpy.h | 10 +++++++
 public_header.py                |  4 +--
 src/library/CMakeLists.txt      |  2 +-
 src/library/dwt.c               |  4 ++-
 src/library/saxpy.hip           | 53 +++++++++++++++++++++++++++++++++
 6 files changed, 70 insertions(+), 5 deletions(-)
 create mode 100644 include/library/private/saxpy.h
 create mode 100644 src/library/saxpy.hip

diff --git a/include/library/private/dwt.h b/include/library/private/dwt.h
index d9fa0eb..f5c95b0 100755
--- a/include/library/private/dwt.h
+++ b/include/library/private/dwt.h
@@ -166,4 +166,4 @@
   //==========|==========================|======================|======|=======|====================
   uchar        inverse_wavelet_transform  (bwc_field                    *const  field, 
                                            bwc_parameter                *const  parameter);
-#endif
\ No newline at end of file
+#endif
diff --git a/include/library/private/saxpy.h b/include/library/private/saxpy.h
new file mode 100644
index 0000000..975d791
--- /dev/null
+++ b/include/library/private/saxpy.h
@@ -0,0 +1,10 @@
+
+#ifdef __cpluplus
+extern "C" {
+#endif
+
+void bwc_saxpy();
+
+#ifdef __cpluplus
+}
+#endif
diff --git a/public_header.py b/public_header.py
index d5c5f45..82613c9 100644
--- a/public_header.py
+++ b/public_header.py
@@ -91,7 +91,7 @@ destination   = current_path.joinpath('include/library/public')
 if os.path.isdir(destination) == False:
   os.mkdir(destination)
 
-include_files = ['macros.h', 'constants.h', 'dwt.h', 'tagtree.h', 'mq_types.h', 'mq.h',
+include_files = ['macros.h', 'constants.h', 'saxpy.h', 'dwt.h', 'tagtree.h', 'mq_types.h', 'mq.h',
                 'bitstream.h', 'codestream.h', 'tier1.h', 'tier2.h', 'types.h', 'libbwc.h']
 exclude_files = ["prim_types_double.h", "prim_types_single.h"]
 all_files = [f for f in os.listdir(source) if os.path.isfile(os.path.join(source, f))]
@@ -380,4 +380,4 @@ public_header.write("\n" + tab + "#ifdef __cplusplus\n" +
                                   tab + "#endif\n")
 
 public_header.write("#endif")
-public_header.close
\ No newline at end of file
+public_header.close
diff --git a/src/library/CMakeLists.txt b/src/library/CMakeLists.txt
index ad118b9..b99cb57 100755
--- a/src/library/CMakeLists.txt
+++ b/src/library/CMakeLists.txt
@@ -79,8 +79,8 @@ add_library(bwclib ${BWC_LINK} bitstream.c
                                libbwc.c
                                codestream.c
                                dwt.c
-			       dwt.hip
                                mq.c
+			       saxpy.hip
                                tier1.c
                                tier2.c
                                tagtree.c)
diff --git a/src/library/dwt.c b/src/library/dwt.c
index 5b94bc6..776fc89 100755
--- a/src/library/dwt.c
+++ b/src/library/dwt.c
@@ -535,6 +535,8 @@ forward_9x7_CDF_wavelet_transform(bwc_sample *const working_buffer, uint64 res0,
       working_buffer[i].f += DELTA * (working_buffer[i - 1].f + working_buffer[i + 1].f);
       working_buffer[i].f  = KAPPA_L * working_buffer[i].f;
    }
+
+    //do_saxpy();
 }
 
 /*----------------------------------------------------------------------------------------------------------*\
@@ -2645,4 +2647,4 @@ inverse_wavelet_transform(bwc_field *const field, bwc_parameter *const parameter
    free(memory);
 
    return 0;
-}
\ No newline at end of file
+}
diff --git a/src/library/saxpy.hip b/src/library/saxpy.hip
new file mode 100644
index 0000000..486b200
--- /dev/null
+++ b/src/library/saxpy.hip
@@ -0,0 +1,53 @@
+
+#include <hip/hip_runtime.h>
+
+#include <iostream>
+
+#include "saxpy.h"
+
+__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];
+}
+
+void bwc_saxpy()
+{
+    int n = 256;
+    std::size_t size = sizeof(float)*n;
+    std::cout << "2 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";
+}
+