AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

add hcc dialects sample

Selaa lähdekoodia
This commit is contained in:
Ben Sander
2016-04-13 17:32:38 -05:00
vanhempi 210ba67b84
commit 624b2f35ff
7 muutettua tiedostoa jossa 358 lisäystä ja 0 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
samples/0_Intro/hcc_dialects/Makefile
+66
Näytä tiedosto
@@ -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 $@
samples/0_Intro/hcc_dialects/vadd_amp_arrayview.cpp
+48
Näytä tiedosto
@@ -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);
}
};
}
samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
+59
Näytä tiedosto
@@ -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);
}
};
}
samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
+53
Näytä tiedosto
@@ -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);
}
};
}
samples/0_Intro/hcc_dialects/vadd_hc_array.hc
+33
Näytä tiedosto
@@ -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);
}
};
}
samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
+48
Näytä tiedosto
@@ -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);
}
};
}
samples/0_Intro/hcc_dialects/vadd_hip.cpp
+51
Näytä tiedosto
@@ -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);
}
};
}