hip saxpy

include/library/private/dwt.h

@@ -166,4 +166,4 @@
   //==========|==========================|======================|======|=======|====================
   uchar        inverse_wavelet_transform  (bwc_field                    *const  field, 
                                            bwc_parameter                *const  parameter);
-#endif
+#endif

include/library/private/saxpy.h

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

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
+public_header.close

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)

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;
-}
+}

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";
+}
+