diff --git a/llvm/lib/Transforms/GenVISC/GenVISC.cpp b/llvm/lib/Transforms/GenVISC/GenVISC.cpp index fef3e5e4c73b5d384a6547081249a21cdbc639ba..2e2a2ac093e0257e6836a92d0be1e63bc9a9a506 100644 --- a/llvm/lib/Transforms/GenVISC/GenVISC.cpp +++ b/llvm/lib/Transforms/GenVISC/GenVISC.cpp @@ -834,7 +834,8 @@ bool GenVISC::runOnModule(Module &M) { "Define LLVM_SRC_ROOT environment variable!"); Twine llvmSrcRoot = LLVM_SRC_ROOT; - Twine runtimeAPI = llvmSrcRoot+"/projects/visc-rt/visc-rt.ll"; + Twine runtimeAPI = llvmSrcRoot+"/../build/projects/visc-rt/visc-rt.ll"; + errs() << llvmSrcRoot << "\n"; std::unique_ptr<Module> runtimeModule = parseIRFile(runtimeAPI.str(), Err, M.getContext()); // Module* runtimeModule = @@ -862,14 +863,14 @@ bool GenVISC::runOnModule(Module &M) { // Insert init context in main Function* VI = M.getFunction("__visc__init"); assert(VI->getNumUses() == 1 && "__visc__init should only be used once"); - Instruction* I = cast<Instruction>(*VI->use_begin()); + Instruction* I = cast<Instruction>(*VI->user_begin()); initializeTimerSet(I); switchToTimer(visc_TimerID_NONE, I); // Insert print instruction at visc exit Function* VC = M.getFunction("__visc__cleanup"); assert(VC->getNumUses() == 1 && "__visc__cleanup should only be used once"); - I = cast<Instruction>(*VC->use_begin()); + I = cast<Instruction>(*VC->user_begin()); printTimerSet(I); diff --git a/llvm/test/VISC/parboil/benchmarks/sgemm/Makefile b/llvm/test/VISC/parboil/benchmarks/sgemm/Makefile index 3582119a311bb6922c26ea4bca094eec8f6ce0a2..6b52d4ffaed7dd84d734e2feedfafa834899398d 100644 --- a/llvm/test/VISC/parboil/benchmarks/sgemm/Makefile +++ b/llvm/test/VISC/parboil/benchmarks/sgemm/Makefile @@ -3,7 +3,7 @@ APP = sgemm # Default compile visc ifeq ($(VERSION),) - VERSION = visc_sh + VERSION = visc endif # Default use small test case diff --git a/llvm/test/VISC/parboil/benchmarks/sgemm/src/visc/main.cc b/llvm/test/VISC/parboil/benchmarks/sgemm/src/visc/main.cc index ff62514be606eedaf6e3d37bd886d01f95595ce5..b2f798a51490007331824b5a9e1237309f9579c2 100644 --- a/llvm/test/VISC/parboil/benchmarks/sgemm/src/visc/main.cc +++ b/llvm/test/VISC/parboil/benchmarks/sgemm/src/visc/main.cc @@ -39,7 +39,7 @@ extern char* readFile(const char*); void mysgemmNT( float* A, int lda, float* B, int ldb, float* C, int ldc, int k, float alpha, float beta ) { - __visc__hint(visc::SPIR_TARGET); + __visc__hint(visc::DEVICE); __visc__attributes(3, A, B, C, 1, C); float c = 0.0f; int m = get_global_id(0); diff --git a/llvm/test/VISC/parboil/common/platform/visc.default.mk b/llvm/test/VISC/parboil/common/platform/visc.default.mk index 8bbbabb4946823f56b2c9a2abe1abb0a2fe14ff5..03b4fd7ecd9e822195d15de4d227fc644e1ebb9f 100644 --- a/llvm/test/VISC/parboil/common/platform/visc.default.mk +++ b/llvm/test/VISC/parboil/common/platform/visc.default.mk @@ -12,19 +12,19 @@ #OPENCL_LIB_PATH=$(OPENCL_PATH)/lib/x86_64 #build -VISC_BUILD_DIR = Release+Asserts +VISC_BUILD_DIR = $(LLVM_SRC_ROOT)/../build # gcc (default) -CC = $(LLVM_SRC_ROOT)/$(VISC_BUILD_DIR)/bin/clang -PLATFORM_CFLAGS = -I$(LLVM_SRC_ROOT)/include +CC = $(VISC_BUILD_DIR)/bin/clang +PLATFORM_CFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(VISC_BUILD_DIR)/include -CXX = $(LLVM_SRC_ROOT)/$(VISC_BUILD_DIR)/bin/clang++ -PLATFORM_CXXFLAGS = -I$(LLVM_SRC_ROOT)/include +CXX = $(VISC_BUILD_DIR)/bin/clang++ +PLATFORM_CXXFLAGS = -I$(LLVM_SRC_ROOT)/include -I$(VISC_BUILD_DIR)/include -LINKER = $(LLVM_SRC_ROOT)/$(VISC_BUILD_DIR)/bin/clang++ +LINKER = $(VISC_BUILD_DIR)/bin/clang++ PLATFORM_LDFLAGS = -lm -lpthread -lOpenCL -LLVM_LIB_PATH = $(LLVM_SRC_ROOT)/$(VISC_BUILD_DIR)/lib -LLVM_BIN_PATH = $(LLVM_SRC_ROOT)/$(VISC_BUILD_DIR)/bin +LLVM_LIB_PATH = $(VISC_BUILD_DIR)/lib +LLVM_BIN_PATH = $(VISC_BUILD_DIR)/bin OPT = $(LLVM_BIN_PATH)/opt LLVM_LINK = $(LLVM_BIN_PATH)/llvm-link