From 624b2f35ffdc09fc757a63a8be59f336266439ef Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 13 Apr 2016 17:32:38 -0500
Subject: [PATCH] add hcc dialects sample

---
 samples/0_Intro/hcc_dialects/Makefile         | 66 +++++++++++++++++++
 .../hcc_dialects/vadd_amp_arrayview.cpp       | 48 ++++++++++++++
 samples/0_Intro/hcc_dialects/vadd_hc_am.cpp   | 59 +++++++++++++++++
 .../0_Intro/hcc_dialects/vadd_hc_array.cpp    | 53 +++++++++++++++
 samples/0_Intro/hcc_dialects/vadd_hc_array.hc | 33 ++++++++++
 .../hcc_dialects/vadd_hc_arrayview.cpp        | 48 ++++++++++++++
 samples/0_Intro/hcc_dialects/vadd_hip.cpp     | 51 ++++++++++++++
 7 files changed, 358 insertions(+)
 create mode 100644 samples/0_Intro/hcc_dialects/Makefile
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_array.hc
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
 create mode 100644 samples/0_Intro/hcc_dialects/vadd_hip.cpp

diff --git a/samples/0_Intro/hcc_dialects/Makefile b/samples/0_Intro/hcc_dialects/Makefile
new file mode 100644
index 0000000000..108d30201c
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/Makefile
@@ -0,0 +1,66 @@
+HCC_HOME?=/opt/rocm/hcc
+HCC = $(HCC_HOME)/bin/hcc
+
+HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` 
+HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags`
+
+CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I/opt/hcc/include
+CPPAMP_LDFLAGS=  -std=c++amp -L/opt/hcc/lib -Wl,--rpath=/opt/hcc/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive 
+
+HIP_PATH?=/opt/rocm/hip
+HIPCC=$(HIP_PATH)/bin/hipcc
+HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
+
+ifneq (${HIP_PLATFORM}, hcc)
+$(error hcc_dialects requires hcc compiler and only runs on hcc platform)
+endif
+
+
+TARGETS=vadd_hc_arrayview vadd_hc_array vadd_amp_arrayview  vadd_hip  
+
+all: $(TARGETS)
+
+clean: 
+	rm -f $(TARGETS) *.o
+
+run: $(TARGETS)
+	@for t in $(TARGETS); do\
+		echo "Running $$t"; \
+		./$$t; \
+	done
+	
+
+# HCC version:
+vadd_hc_arrayview.o: vadd_hc_arrayview.cpp
+	$(HCC) $(HCC_CFLAGS) -c $< -o $@
+vadd_hc_arrayview: vadd_hc_arrayview.o
+	$(HCC) $(HCC_LDFLAGS) $< -o $@
+
+
+# HCC version, using explicit arrays:
+vadd_hc_array.o: vadd_hc_array.cpp
+	$(HCC) $(HCC_CFLAGS) -c $< -o $@
+vadd_hc_array: vadd_hc_array.o
+	$(HCC) $(HCC_LDFLAGS) $< -o $@
+
+
+# HCC version, using AM (accelerator memory) pointer
+vadd_hc_am.o: vadd_hc_am.cpp
+	$(HCC) $(HCC_CFLAGS) -c $< -o $@
+vadd_hc_am: vadd_hc_am.o
+	$(HCC) $(HCC_LDFLAGS) $< -o $@
+
+
+
+# HIP version:
+vadd_hip.o: vadd_hip.cpp
+	$(HIPCC) -c $< -o $@
+vadd_hip: vadd_hip.o
+	$(HIPCC) $< -o $@
+
+
+# AMP version:
+vadd_amp_arrayview.o: vadd_amp_arrayview.cpp
+	$(HCC) $(CPPAMP_CFLAGS) -c $< -o $@
+vadd_amp_arrayview: vadd_amp_arrayview.o
+	$(HCC) $(CPPAMP_LDFLAGS) $< -o $@
diff --git a/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp b/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
new file mode 100644
index 0000000000..6fdea5d830
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
@@ -0,0 +1,48 @@
+// Simple test showing how to use C++AMP syntax with array_view.
+// The code uses AMP's array_view class, which provides automatic data synchronization
+// of data between the host and the accelerator. As noted below, the HCC runtime
+// will automatically copy data to and from the host, without the user needing
+// to manually perform such copies.  This is an excellent mode for developers
+// new to GPU programming and matches the memory models provided by recent systems where
+// CPU and GPU share the same memory pool.  Advanced programmers may prefer 
+// more explicit control over the data movement - shown in the other vadd_hc_array and 
+// vadd_hc_am examples.
+// This example shows the similarity between C++AMP and and HC for simple cases where
+// implicit data transfer is used - really the only difference is the namespace.
+// Other examples show some of the more advanced controls.
+
+#include <amp.h>
+
+int main(int argc, char *argv[])
+{
+    int sizeElements = 1000000;
+
+    // Allocate auto-managed host/device views of data:
+    concurrency::array_view<float> A(sizeElements);
+    concurrency::array_view<float> B(sizeElements);
+    concurrency::array_view<float> C(sizeElements);
+
+    // Initialize host data
+    for (int i=0; i<sizeElements; i++) {
+        A[i] = 1.618f * i; 
+        B[i] = 3.142f * i;
+    }
+    C.discard_data(); // tell runtime not to copy CPU host data.
+
+
+    // Launch kernel onto default accelerator
+    // The HCC runtime will ensure that A and B are available on the accelerator before launching the kernel.
+    concurrency::parallel_for_each(concurrency::extent<1> (sizeElements),
+      [=] (concurrency::index<1> idx) restrict(amp) { 
+        int i = idx[0];
+        C[i] = A[i] + B[i];
+    });
+
+    for (int i=0; i<sizeElements; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        // Because C is an array_view, the HCC runtime will copy C back to host at first access here:
+        if (C[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+        }
+    };
+}
diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
new file mode 100644
index 0000000000..df6a831fd2
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
@@ -0,0 +1,59 @@
+// Simple test showing how to use HC syntax with AM (accelerator memory).
+// AM provides a set of c-style memory management routines for allocating,
+// freeing, and copying memory.   am_alloc returns a device pointer
+// which can only be used on the device.  The programmer has full control
+// over when data is copied.  
+
+#include <hc.hpp>
+#include <hc_am.hpp>
+
+int main(int argc, char *argv[])
+{
+    int sizeElements = 1000000;
+    size_t sizeBytes = sizeElements * sizeof(float);
+
+    // Allocate host memory
+    float *A_h = (float*)malloc(sizeBytes);
+    float *B_h = (float*)malloc(sizeBytes);
+    float *C_h = (float*)malloc(sizeBytes);
+
+    // Allocate device pointers:
+    // Unlike array_view, these must be explicitly managed by user:
+    hc::accelerator acc;  // grab default accelerator where we want to allocate memory:
+    hc::accelerator_view av = acc.get_default_view();
+
+    float *A_d, *B_d, *C_d;
+    A_d = hc::am_alloc(sizeBytes, acc, 0);
+    B_d = hc::am_alloc(sizeBytes, acc, 0);
+    C_d = hc::am_alloc(sizeBytes, acc, 0);
+
+    // Initialize host data
+    for (int i=0; i<sizeElements; i++) {
+        A_h[i] = 1.618f * i; 
+        B_h[i] = 3.142f * i;
+    }
+
+    av.copy(A_h, A_d); // C++ copy H2D
+    av.copy(B_h, B_d); //C++  copy H2D
+
+    // Launch kernel onto AV.  
+    // Because the kernel PFE and the copies are submitted to same AV, they will execute in order
+    // and we don't need additional synchronization to ensure the copies complete before the PFE begins.
+    hc::parallel_for_each(av,  hc::extent<1> (sizeElements),
+      [&] (hc::index<1> idx) [[hc]] { 
+        int i = idx[0];
+        C_d[i] = A_d[i] + B_d[i];
+    });
+
+   
+    // This copy is in same AV as the kernel and thus will wait for the kernel to finish before executing.
+    av.copy(C_d, C_h); // C++ copy D2H
+
+
+    for (int i=0; i<sizeElements; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        if (C_h[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+        }
+    };
+}
diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp b/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
new file mode 100644
index 0000000000..63f455a615
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
@@ -0,0 +1,53 @@
+// Simple test showing how to use HC syntax with array.
+// Array provides a type-safe C++ mechanism to allocate accelerator memory.  
+// Like array_view, hc::array provides multi-dimensional indexing capability,
+// and is typed.  However, unlike array_view, hc::array does not provide
+// automatic data management capabilities - instead the programmer
+// takes the reins and controls when copies are executed.
+
+#include <hc.hpp>
+
+int main(int argc, char *argv[])
+{
+    int sizeElements = 1000000;
+    size_t sizeBytes = sizeElements * sizeof(float);
+
+    // Allocate host memory
+    float *A_h = (float*)malloc(sizeBytes);
+    float *B_h = (float*)malloc(sizeBytes);
+    float *C_h = (float*)malloc(sizeBytes);
+
+    // Allocate device arrays<>
+    // Unlike array_view, these must be explicitly managed by user:
+    hc::array<float> A_d(sizeElements);
+    hc::array<float> B_d(sizeElements);
+    hc::array<float> C_d(sizeElements);
+
+    // Initialize host data
+    for (int i=0; i<sizeElements; i++) {
+        A_h[i] = 1.618f * i; 
+        B_h[i] = 3.142f * i;
+    }
+
+    hc::copy(A_h, A_d); // C++ copy H2D
+    hc::copy(B_h, B_d); // C++ copy H2D
+
+    // Launch kernel onto default accelerator:
+    // array<> types are not implicitly copied, so we performed copies above.
+    hc::parallel_for_each(hc::extent<1> (sizeElements),
+      [&] (hc::index<1> idx) [[hc]] { 
+        int i = idx[0];
+        C_d[i] = A_d[i] + B_d[i];
+    });
+
+    // HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the PFE to complte.
+    hc::copy(C_d, C_h); // C++ copy D2H
+
+
+    for (int i=0; i<sizeElements; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        if (C_h[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+        }
+    };
+}
diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_array.hc b/samples/0_Intro/hcc_dialects/vadd_hc_array.hc
new file mode 100644
index 0000000000..8a9f8e71e7
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_hc_array.hc
@@ -0,0 +1,33 @@
+#include <hc.hpp>
+
+int main(int argc, char *argv[])
+{
+    int size = 1000000;
+
+    // Allocate auto-managed host/device views of data:
+    hc::array_view<float> A(size);
+    hc::array_view<float> B(size);
+    hc::array_view<float> C(size);
+
+    // Initialize host data
+    for (int i=0; i<size; i++) {
+        A[i] = 1.618f * i; 
+        B[i] = 3.142f * i;
+    }
+    C.discard_data(); // tell runtime not to copy CPU host data.
+
+
+    // Launch kernel onto default accelerator:
+    hc::parallel_for_each(hc::extent<1> (size),
+      [=] (hc::index<1> idx) [[hc]] { 
+        int i = idx[0];
+        C[i] = A[i] + B[i];
+    });
+
+    for (int i=0; i<size; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        if (C[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+        }
+    };
+}
diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp b/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
new file mode 100644
index 0000000000..41d9124ddc
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
@@ -0,0 +1,48 @@
+// Simple test showing how to use HC syntax with array_view.
+// The code uses AMP's array_view class, which provides automatic data synchronization
+// of data between the host and the accelerator. As noted below, the HCC runtime
+// will automatically copy data to and from the host, without the user needing
+// to manually perform such copies.  This is an excellent mode for developers
+// new to GPU programming and matches the memory models provided by recent systems where
+// CPU and GPU share the same memory pool.  Advanced programmers may prefer 
+// more explicit control over the data movement - shown in the other vadd_hc_array and 
+// vadd_hc_am examples.
+// This example shows the similarity between C++AMP and and HC for simple cases where
+// implicit data transfer is used - really the only difference is the namespace.
+// Other examples show some of the more advanced controls.
+
+#include <hc.hpp>
+
+int main(int argc, char *argv[])
+{
+    int sizeElements = 1000000;
+
+    // Allocate auto-managed host/device views of data:
+    hc::array_view<float> A(sizeElements);
+    hc::array_view<float> B(sizeElements);
+    hc::array_view<float> C(sizeElements);
+
+    // Initialize host data
+    for (int i=0; i<sizeElements; i++) {
+        A[i] = 1.618f * i; 
+        B[i] = 3.142f * i;
+    }
+    C.discard_data(); // tell runtime not to copy CPU host data.
+
+
+    // Launch kernel onto default accelerator:
+    // The HCC runtime will ensure that A and B are available on the accelerator before launching the kernel.
+    hc::parallel_for_each(hc::extent<1> (sizeElements),
+      [=] (hc::index<1> idx) [[hc]] { 
+        int i = idx[0];
+        C[i] = A[i] + B[i];
+    });
+
+    for (int i=0; i<sizeElements; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        // Because C is an array_view, the HCC runtime will copy C back to host at first access here:
+        if (C[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref);
+        }
+    };
+}
diff --git a/samples/0_Intro/hcc_dialects/vadd_hip.cpp b/samples/0_Intro/hcc_dialects/vadd_hip.cpp
new file mode 100644
index 0000000000..31751b3419
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/vadd_hip.cpp
@@ -0,0 +1,51 @@
+#include <hip_runtime.h>
+
+__global__ void vadd_hip(hipLaunchParm lp, const float *a, const float *b, float *c, int N)
+{
+    int idx = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+
+    if (idx < N) {
+        c[idx] = a[idx] + b[idx];
+    }
+}
+
+
+int main(int argc, char *argv[])
+{
+    int sizeElements = 1000000;
+    size_t sizeBytes = sizeElements * sizeof(float);
+
+    // Allocate host memory
+    float *A_h = (float*)malloc(sizeBytes);
+    float *B_h = (float*)malloc(sizeBytes);
+    float *C_h = (float*)malloc(sizeBytes);
+
+    // Allocate device memory:
+    float *A_d, *B_d, *C_d;
+    hipMalloc(&A_d, sizeBytes);
+    hipMalloc(&B_d, sizeBytes);
+    hipMalloc(&C_d, sizeBytes);
+
+    // Initialize host data
+    for (int i=0; i<sizeElements; i++) {
+        A_h[i] = 1.618f * i; 
+        B_h[i] = 3.142f * i;
+    }
+
+    hipMemcpy(A_d, A_h, sizeBytes, hipMemcpyHostToDevice);
+    hipMemcpy(B_d, B_h, sizeBytes, hipMemcpyHostToDevice);
+
+    // Launch kernel onto default accelerator:
+    int blockSize = 256;  // pick arbitrary block size
+    int blocks = (sizeElements+blockSize-1)/blockSize; // round up to launch enough blocks
+    hipLaunchKernel(vadd_hip, dim3(blocks), dim3(blockSize), 0, 0, A_d, B_d, C_d, sizeElements);
+
+    hipMemcpy(C_h, C_d, sizeBytes, hipMemcpyDeviceToHost);
+
+    for (int i=0; i<sizeElements; i++) {
+        float ref= 1.618f * i + 3.142f * i;
+        if (C_h[i] != ref) {
+            printf ("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref);
+        }
+    };
+}