diff --git a/hpvm/projects/hpvm-rt/CMakeLists.txt b/hpvm/projects/hpvm-rt/CMakeLists.txt
index 5d75c6e0f12622ac11957fdb29d5691eeb979017..b8a1716f2a7c18c7e1dcb8d907c3585707c22386 100644
--- a/hpvm/projects/hpvm-rt/CMakeLists.txt
+++ b/hpvm/projects/hpvm-rt/CMakeLists.txt
@@ -13,7 +13,7 @@ add_llvm_library(hpvm-rt.ll hpvm-rt.cpp
)
target_compile_options(hpvm-rt.ll PUBLIC -flto)
target_include_directories(hpvm-rt.ll PRIVATE ${OpenCL_INCLUDE_DIRS})
-target_link_directories(hpvm-rt.ll PRIVATE ${OpenCL_LIBRARY})
+link_directories(${OpenCL_LIBRARY})
add_custom_target(hpvm-rt.cpp.o ALL
COMMAND ar -x ${CMAKE_BINARY_DIR}/lib/libhpvm-rt.ll.a
diff --git a/hpvm/projects/hpvm-rt/hpvm-rt.cpp b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
index 20af30fdfa3d23401d05644a45c0caf124333524..e0e017c03e017edef7e6c1dfed17ceb8db9d2ba5 100644
--- a/hpvm/projects/hpvm-rt/hpvm-rt.cpp
+++ b/hpvm/projects/hpvm-rt/hpvm-rt.cpp
@@ -16,7 +16,7 @@
#endif
#include "hpvm-rt.h"
-#define DEBUG_BUILD
+//#define DEBUG_BUILD
#ifndef DEBUG_BUILD
#define DEBUG(s) \
{}
diff --git a/hpvm/test/CTestSuite/Makefile b/hpvm/test/CTestSuite/Makefile
deleted file mode 100644
index 1169e4e896a861975ac0562ebff8b208828bbf89..0000000000000000000000000000000000000000
--- a/hpvm/test/CTestSuite/Makefile
+++ /dev/null
@@ -1,26 +0,0 @@
-PASSES :=
-
-.PHONY: clean
-
-LLVM_INSTALL:=$(LLVM_SRC_ROOT)/Release+Asserts
-LIBCLC:=$(LLVM_SRC_ROOT)/../libclc
-HOST:=gemm gemm_2
-LLVM_CC:=$(LLVM_INSTALL)/bin/clang
-LLVM_OPT:=$(LLVM_INSTALL)/bin/opt
-BUILD_DIR:=build
-
-all: $(BUILD_DIR) $(HOST:%=$(BUILD_DIR)/%.ll) $(HOST:%=$(BUILD_DIR)/%.hpvm.ll)
-
-$(BUILD_DIR):
- mkdir -p $(BUILD_DIR)
-
-$(HOST:%=$(BUILD_DIR)/%.ll):$(BUILD_DIR)/%.ll:%.c
- $(LLVM_CC) -S -emit-llvm $< -O3 -o $@
-
-$(HOST:%=$(BUILD_DIR)/%.hpvm.ll):$(BUILD_DIR)/%.hpvm.ll:$(BUILD_DIR)/%.ll
- $(LLVM_OPT) -load $(LLVM_SRC_ROOT)/Release+Asserts/lib/LLVMGenHPVM.so -genhpvm -globaldce $< -S -o $@
- @cat RUN.script $@ > $@.tmp
- @mv $@.tmp $@
-
-clean :
- rm -f $(HOST:%=$(BUILD_DIR)/%.ll) $(HOST:%=$(BUILD_DIR)/%.hpvm.ll) $(HOST:%=$(BUILD_DIR)/%.hpvm.ll.kernels.ll) $(HOST:%=$(BUILD_DIR)/%.hpvm.ll.nvptx.s) $(BUILD_DIR)/DataflowGraph.dot*
diff --git a/hpvm/test/CTestSuite/RUN.script b/hpvm/test/CTestSuite/RUN.script
deleted file mode 100644
index 23fa1694ebf4b7448c731327b96b949c0509b62e..0000000000000000000000000000000000000000
--- a/hpvm/test/CTestSuite/RUN.script
+++ /dev/null
@@ -1,6 +0,0 @@
-; RUN: opt -load LLVMBuildDFG.so -load LLVMDFG2LLVM_NVPTX.so -load LLVMDFG2LLVM_X86.so -load LLVMClearDFG.so -dfg2llvm-nvptx -dfg2llvm-x86 -clearDFG -o %t.ll -S %s
-; RUN: llvm-link %llvm_src/../libclc/built_libs/nvptx--nvidiacl.bc %s.kernels.ll -o %t.ll.kernels.linked.bc
-; RUN: clang -O3 -target nvptx %t.ll.kernels.linked.bc -S -o %s.nvptx.s
-; RUN: llvm-link %t.ll %llvm_src/projects/hpvm-rt/hpvm-rt.ll -S -o %t.linked.ll
-; RUN: clang++ -O3 %t.linked.ll -lpthread -lOpenCL -o %t.bin
-; RUN: %t.bin
diff --git a/hpvm/test/CTestSuite/gemm.c b/hpvm/test/CTestSuite/gemm.c
deleted file mode 100644
index eb0a3c5e9204d9621c4a15ae7f07ef5158ac1d07..0000000000000000000000000000000000000000
--- a/hpvm/test/CTestSuite/gemm.c
+++ /dev/null
@@ -1,147 +0,0 @@
-#include <math.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-
-#define WA 1024
-#define HA 1024
-#define WB 1024
-#define HB WA
-#define WC WB
-#define HC HA
-
-// Thread block size
-#define BLOCK_SIZE 16
-
-// Allocates a matrix with random float entries.
-void randomInit(float *data, int size) {
- for (int i = 0; i < size; ++i)
- data[i] = rand() / (float)RAND_MAX;
-}
-
-//////////////////////////////////////////////////////////////////////////////
-//! Loads a Program file.
-//!
-//! @return the source string if succeeded, 0 otherwise
-//! @param cFilename program filename
-//! @param szFinalLength returned length of the code string
-//////////////////////////////////////////////////////////////////////////////
-
-// Check bool
-int isEqual(float a, float b) { return (fabs(a - b) < 0.001); }
-
-// Check Results
-
-__attribute__((noinline)) int checkResults(float *A, float *B, float *C) {
- unsigned int size_A = WA * HA;
- unsigned int size_B = WB * HB;
- unsigned int size_C = WC * HC;
- unsigned int bytesC = sizeof(float) * size_C;
- float *goldC = (float *)malloc(bytesC);
- for (int i = 0; i < HC; i++) {
- for (int j = 0; j < WC; j++) {
- goldC[i * WC + j] = 0;
- for (int k = 0; k < HB; k++) {
- goldC[i * WC + j] += A[i * WA + k] * B[k * WB + j];
- }
- if (!isEqual(goldC[i * WC + j], C[i * WC + j])) {
- printf("Mismatch at %d,%d --- C = %f and goldC = %f\n", i, j,
- C[i * WC + j], goldC[i * WC + j]);
- return 0;
- }
- }
- }
- return 1; // Success
-}
-
-// Dummy hpvm node execution call
-// void __hpvm__node(void kernel (float*, float*, float*, unsigned, unsigned),
-// int numDims, void* dims, int numInputs, void* inputs, int numOutputs, void*
-// outputs);
-
-void matrixMul(float *A, float *B, float *C, unsigned k, unsigned n) {
-
- __hpvm__attributes(2, A, B, 1, C);
- // printf("Entered function\n");
- int tx = get_local_id(0); // 2D Global Thread ID x
- int ty = get_local_id(1); // 2D Global Thread ID y
- // int tx = get_global_id(0); //2D Global Thread ID x
- // int ty = get_global_id(1); //2D Global Thread ID y
-
- // printf("Computing element (%d, %d)\n", tx, ty);
- // Initialize accumulator
- float res = 0.0f;
-
- // Perform dot-product of row-column
- for (int i = 0; i < k; i++) {
- // printf("Accessing k = %d, A[%d], B[%d]\n", k, ty*k+i, i*n+tx);
- res += A[ty * k + i] * B[i * n + tx];
- }
-
- // printf("Result computed\n");
- // Write in device memory
- C[ty * n + tx] = res;
-
- // printf("Result written to C\n");
-}
-
-// Main
-int main(int argc, char **argv) {
-
- // seed for rand()
- srand(2006);
-
- // Allocate host memory for matrices A and B
- unsigned int size_A = WA * HA;
- size_t bytes_A = sizeof(float) * size_A;
- float *h_A = (float *)malloc(bytes_A);
-
- unsigned int size_B = WB * HB;
- size_t bytes_B = sizeof(float) * size_B;
- float *h_B = (float *)malloc(bytes_B);
-
- // Initialize host memory
- randomInit(h_A, size_A);
- randomInit(h_B, size_B);
-
- /*
- // Print A and B
- printf("\n\nMatrix A\n");
- for(int i = 0; i < size_A; i++)
- {
- printf("%f ", h_A[i]);
- if(((i + 1) % WA) == 0)
- printf("\n");
- }
-
- printf("\n\nMatrix B\n");
- for(int i = 0; i < size_B; i++)
- {
- printf("%f ", h_B[i]);
- if(((i + 1) % WB) == 0)
- printf("\n");
- }
- */
-
- // Allocate host memory for the result matrix C
- unsigned int size_C = WC * HC;
- size_t bytes_C = sizeof(float) * size_C;
- float *h_C = (float *)malloc(bytes_C);
-
- // Compute using OpenCL
- // matrixMul(h_A, h_B, h_C, WA, WB);
- //__hpvm__node(matrixMul, 2, WB, HA, 3, h_A, h_B, h_C, 0);
- unsigned graphMM = __hpvm__node(matrixMul, 1, 2, WB, HA, 8, h_A, bytes_A, h_B,
- bytes_B, h_C, bytes_C, WA, WB, 0);
- __hpvm__wait(graphMM);
- if (checkResults(h_A, h_B, h_C))
- printf("\nPass!\n");
- else
- printf("\nFailed!\n");
- printf("\nDone!\n");
-
- // Deallocate memory
- free(h_A);
- free(h_B);
- free(h_C);
-}
diff --git a/hpvm/test/CTestSuite/gemm_2.c b/hpvm/test/CTestSuite/gemm_2.c
deleted file mode 100644
index df4555936316703cfccd4048f2ade4e28592e53a..0000000000000000000000000000000000000000
--- a/hpvm/test/CTestSuite/gemm_2.c
+++ /dev/null
@@ -1,148 +0,0 @@
-#include <math.h>
-#include <stdio.h>
-#include <stdlib.h>
-#include <string.h>
-
-#define WA 1024
-#define HA 1024
-#define WB 1024
-#define HB WA
-#define WC WB
-#define HC HA
-
-// Thread block size
-#define BLOCK_SIZE 16
-
-// Allocates a matrix with random float entries.
-void randomInit(float *data, int size) {
- for (int i = 0; i < size; ++i)
- data[i] = rand() / (float)RAND_MAX;
-}
-
-//////////////////////////////////////////////////////////////////////////////
-//! Loads a Program file.
-//!
-//! @return the source string if succeeded, 0 otherwise
-//! @param cFilename program filename
-//! @param szFinalLength returned length of the code string
-//////////////////////////////////////////////////////////////////////////////
-
-// Check bool
-int isEqual(float a, float b) { return (fabs(a - b) < 0.001); }
-
-// Check Results
-
-__attribute__((noinline)) int checkResults(float *A, float *B, float *C) {
- unsigned int size_A = WA * HA;
- unsigned int size_B = WB * HB;
- unsigned int size_C = WC * HC;
- unsigned int bytesC = sizeof(float) * size_C;
- float *goldC = (float *)malloc(bytesC);
- for (int i = 0; i < HC; i++) {
- for (int j = 0; j < WC; j++) {
- goldC[i * WC + j] = 0;
- for (int k = 0; k < HB; k++) {
- goldC[i * WC + j] += A[i * WA + k] * B[k * WB + j];
- }
- if (!isEqual(goldC[i * WC + j], C[i * WC + j])) {
- printf("Mismatch at %d,%d --- C = %f and goldC = %f\n", i, j,
- C[i * WC + j], goldC[i * WC + j]);
- return 0;
- }
- }
- }
- return 1; // Success
-}
-
-// Dummy hpvm node execution call
-// void __hpvm__node(void kernel (float*, float*, float*, unsigned, unsigned),
-// int numDims, void* dims, int numInputs, void* inputs, int numOutputs, void*
-// outputs);
-
-void matrixMul(float *A, float *B, float *C, unsigned k, unsigned n) {
- __hpvm__attributes(2, A, B, 1, C);
-
- // printf("Entered function\n");
- int tx = get_global_id(0); // 2D Global Thread ID x
- int ty = get_global_id(1); // 2D Global Thread ID y
- // int tx = get_global_id(0); //2D Global Thread ID x
- // int ty = get_global_id(1); //2D Global Thread ID y
-
- // printf("Computing element (%d, %d)\n", tx, ty);
- // Initialize accumulator
- float res = 0.0f;
-
- // Perform dot-product of row-column
- for (int i = 0; i < k; i++) {
- // printf("Accessing k = %d, A[%d], B[%d]\n", k, ty*k+i, i*n+tx);
- res += A[ty * k + i] * B[i * n + tx];
- }
-
- // printf("Result computed\n");
- // Write in device memory
- C[ty * n + tx] = res;
-
- // printf("Result written to C\n");
-}
-
-// Main
-int main(int argc, char **argv) {
-
- // seed for rand()
- srand(2006);
-
- // Allocate host memory for matrices A and B
- unsigned int size_A = WA * HA;
- size_t bytes_A = sizeof(float) * size_A;
- float *h_A = (float *)malloc(bytes_A);
-
- unsigned int size_B = WB * HB;
- size_t bytes_B = sizeof(float) * size_B;
- float *h_B = (float *)malloc(bytes_B);
-
- // Initialize host memory
- randomInit(h_A, size_A);
- randomInit(h_B, size_B);
-
- /*
- // Print A and B
- printf("\n\nMatrix A\n");
- for(int i = 0; i < size_A; i++)
- {
- printf("%f ", h_A[i]);
- if(((i + 1) % WA) == 0)
- printf("\n");
- }
-
- printf("\n\nMatrix B\n");
- for(int i = 0; i < size_B; i++)
- {
- printf("%f ", h_B[i]);
- if(((i + 1) % WB) == 0)
- printf("\n");
- }
- */
-
- // Allocate host memory for the result matrix C
- unsigned int size_C = WC * HC;
- size_t bytes_C = sizeof(float) * size_C;
- float *h_C = (float *)malloc(bytes_C);
-
- // Compute using OpenCL
- // matrixMul(h_A, h_B, h_C, WA, WB);
- //__hpvm__node(matrixMul, 2, WB, HA, 3, h_A, h_B, h_C, 0);
- unsigned graphMM =
- __hpvm__node(matrixMul, 2, 2, 16, 16, WB / 16, HA / 16, 8, h_A, bytes_A,
- h_B, bytes_B, h_C, bytes_C, WA, WB, 0);
- __hpvm__wait(graphMM);
- if (checkResults(h_A, h_B, h_C))
- printf("\nPass!\n");
- else
- printf("\nFailed!\n");
- printf("\nDone!\n");
-
- // Deallocate memory
- free(h_A);
- free(h_B);
- free(h_C);
-}
diff --git a/hpvm/test/include/Makefile.config.example b/hpvm/test/include/Makefile.config.example
index ae7fe03d9e6b42e51c9860fb3ae508fe90421f34..06267a75eebc094bb064b7162a24876c35d9fe12 100644
--- a/hpvm/test/include/Makefile.config.example
+++ b/hpvm/test/include/Makefile.config.example
@@ -1,12 +1,12 @@
# These paths should be set to your local directories of CUDA and OpenCL
-CUDA_PATH=/software/cuda-9.1
+CUDA_PATH=/software/cuda-9.1
CUDA_LIB_PATH=$(CUDA_PATH)/lib64
OPENCL_PATH=$(CUDA_PATH)
OPENCL_LIB_PATH=$(OPENCL_PATH)/lib64
# This path should be set to your HPVM build directory
-LLVM_BUILD_DIR =PATH/TO/LOCAL/HPVM/BUILD
-HPVM_TEST_DIR = PATH/TO/LOCAL/HPVM/BUILD
+LLVM_BUILD_DIR=PATH/TO/LOCAL/HPVM/BUILD
+HPVM_TEST_DIR=PATH/TO/HPVM/TEST/DIR
CC = $(LLVM_BUILD_DIR)/bin/clang
PLATFORM_CFLAGS = -I$(OPENCL_PATH)/include/CL/ -I$(HPVM_TEST_DIR)/include
diff --git a/hpvm/test/parboil/common/mk/hpvm.mk b/hpvm/test/parboil/common/mk/hpvm.mk
index e42e95f0d788a5af278ba46133d1fe7ce3eee44c..24778a990159365ab024f08c0784cf2eac27f32b 100755
--- a/hpvm/test/parboil/common/mk/hpvm.mk
+++ b/hpvm/test/parboil/common/mk/hpvm.mk
@@ -44,10 +44,6 @@ else
TESTGEN_OPTFLAGS += -hpvm-timers-gen
endif
-ifeq ($(DABSTRACTION),true)
- HPVM_OPTFLAGS += -hpvm-eda
-endif
-
# Rules common to all makefiles
########################################
diff --git a/hpvm/test/pipeline/Makefile b/hpvm/test/pipeline/Makefile
index f46b52391a33b846996dff6311981596bcc3bcad..0511f2b05a5de9e1b68896d1128f5d9913ffe329 100644
--- a/hpvm/test/pipeline/Makefile
+++ b/hpvm/test/pipeline/Makefile
@@ -22,7 +22,7 @@ CURRENT_DIR := $(dir $(abspath $(lastword $(MAKEFILE_LIST))))
EXE = pipeline-$(TARGET)
-INCLUDES += -I$(SRC_DIR) -I$(CAM_PIPE_SRC_DIR)
+INCLUDES += -I$(SRC_DIR)
## BEGIN HPVM MAKEFILE
SRCDIR_OBJS= io.ll
diff --git a/hpvm/test/pipeline/src/main.cc b/hpvm/test/pipeline/src/main.cc
index 2a69a75037f622d5fffcdaeda8dc67f5cbdaf81d..cda1d975a63fc07c174ed57ddef1e72f0973f033 100644
--- a/hpvm/test/pipeline/src/main.cc
+++ b/hpvm/test/pipeline/src/main.cc
@@ -834,7 +834,7 @@ int main(int argc, char *argv[]) {
resize(E, out, Size(HEIGHT, WIDTH));
imshow(input_window, in);
imshow(output_window, out);
- waitKey(0);
+// waitKey(0);
struct InStruct *args = (struct InStruct *)malloc(sizeof(InStruct));
packData(args, (float *)src.data, I_sz, (float *)Is.data, I_sz,
@@ -854,7 +854,6 @@ int main(int argc, char *argv[]) {
if (NUM_FRAMES >= 2) {
for (int i = 0; i < NUM_FRAMES; i++) {
- std::cout << "Frame " << i << "\n";
args->I = (float *)src.data;
*maxG = 0.0;
@@ -873,8 +872,8 @@ int main(int argc, char *argv[]) {
__hpvm__push(DFG, args);
void *ret = __hpvm__pop(DFG);
- std::cout << "Returned size: " << ((OutStruct *)ret)->ret
- << " expected " << I_sz << '\n';
+ // This is reading the result of the streaming graph
+ size_t framesize = ((OutStruct *)ret)->ret;
llvm_hpvm_request_mem(maxG, bytesMaxG);
llvm_hpvm_request_mem(E.data, I_sz);