diff --git a/llvm/projects/hpvm-tensor-rt/CMakeLists.txt b/llvm/projects/hpvm-tensor-rt/CMakeLists.txt
index d86c8fa3cb7cb70bbe23bda5cb21b57637ab9544..76730ff31402aa90a589ee17f4b45b2ecc2ca7b0 100644
--- a/llvm/projects/hpvm-tensor-rt/CMakeLists.txt
+++ b/llvm/projects/hpvm-tensor-rt/CMakeLists.txt
@@ -1,49 +1,38 @@
 cmake_minimum_required (VERSION 2.6)
 project (hpvm-tensor-rt)
-
-
-
-SET(CUDA_SEPARABLE_COMPILATION ON)
 find_package(CUDA 6.5 REQUIRED)
+set(CUDA_SEPARABLE_COMPILATION ON CACHE BOOL "")
+set(CUDA_PROPAGATE_HOST_FLAGS OFF)
 
+# Addresses a bug where code is not compiled as C++11 in non-CUDA code and older g++ versions
+# Edit: using c++14 now
+set(CMAKE_CXX_STANDARD 14)
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -I/")
 set(
   CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};
   -gencode;arch=compute_60,code=sm_60;
   -gencode;arch=compute_60,code=compute_60;
-  -std=c++14 --expt-relaxed-constexpr -maxrregcount 32 # These are newly added
+  -std=c++14 --expt-relaxed-constexpr -maxrregcount 32 # These are for image ops
 )
-
-if (CMAKE_BUILD_TYPE STREQUAL "Debug")
+if(CMAKE_BUILD_TYPE STREQUAL "Debug")
   message("Debug mode")
-  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-g;-lineinfo;-Xcompiler;-ggdb;-lcurand)
+  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-g;-lineinfo;-Xcompiler;-ggdb)
 else()
-  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-DNDEBUG;-Xcompiler;-DNDEBUG;-lcurand)
+  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-DNDEBUG;-Xcompiler;-DNDEBUG)
 endif()
 
-set(CUDA_PROPAGATE_HOST_FLAGS OFF)
-
-# Addresses a bug where code is not compiled as C++11 in non-CUDA code and older g++ versions
-# Edit: using c++14 now
-set(CMAKE_CXX_STANDARD 14)
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14  -I/  " )
-
+# Default options
 if(USE_GFLAGS)
   add_definitions(-DUSE_GFLAGS)
 endif()
-
 if(USE_AUTOTUNER)
   remove_definitions(-DNO_INJECTION)
 endif()
-
-
 add_definitions(-DNO_INJECTION)
 add_definitions(-DPROMISE_TUNER_ENABLED)
 add_definitions(-DSIMULATION_MODE=true)
-add_definitions(-DONLINE_PROFILING=false)
-add_definitions(-DFP16_tuning=true)
-
-
 
+# Default include/link directories
 include_directories($ENV{CUDNN_PATH} $ENV{CUDNN_PATH}/include)
 include_directories(./tensor_runtime/include)
 include_directories(./dnn_sources/include)
@@ -51,6 +40,7 @@ include_directories(../gpu_profiler/include)
 include_directories(../soc_simulator/include)
 link_directories(${CUDA_TOOLKIT_ROOT_DIR}/lib64 $ENV{CUDNN_PATH} $ENV{CUDNN_PATH}/lib $ENV{CUDNN_PATH}/lib64)
 
+# Source files of runtime
 set(
   RUNTIME_SRCS_FILENAME
   approx_simulation.cu
@@ -79,44 +69,43 @@ foreach(FILE ${RUNTIME_SRCS_FILENAME})
   list(APPEND RUNTIME_SRCS "tensor_runtime/src/${FILE}")
 endforeach()
 
+# Default link libraries
 find_library(GPU_PROFILER_LIB
-    NAMES libgpu_profiler.a
-    HINTS ../gpu_profiler/lib
+  NAMES libgpu_profiler.a
+  HINTS ../gpu_profiler/lib
 )
 find_library(SOC_SIMULATOR_LIB
-    NAMES libpromise_profiler.a
-    HINTS ../soc_simulator/lib
+  NAMES libpromise_profiler.a
+  HINTS ../soc_simulator/lib
 )
-set(LINK_LIBS cudnn cufft stdc++fs -lcurand)
+set(LINK_LIBS cudnn cufft stdc++fs curand)
 if(USE_GFLAGS)
   list(APPEND LINK_LIBS gflags)
 endif()
 
-
 # Adding new rule for building a cuDNN runtime library
 # Offline version
 cuda_add_library(tensor_runtime ${RUNTIME_SRCS})
 cuda_add_cublas_to_target(tensor_runtime)
 target_link_libraries(tensor_runtime ${LINK_LIBS})
+target_compile_definitions(tensor_runtime PRIVATE -DONLINE_PROFILING=false -DFP16_tuning=true)
 
-
+# Install version (also offline)
 cuda_add_library(tensor_runtime_install ${RUNTIME_SRCS})
 cuda_add_cublas_to_target(tensor_runtime_install)
-target_link_libraries(tensor_runtime_install ${LINK_LIBS})
 # tensor_runtime_install is built AFTER tensor_runtime because of a nvcc bug (bug?)
 # that doesn't allow compiling the same file from multiple targets at once.
 # Same for tensor_runtime_online.
 add_dependencies(tensor_runtime_install tensor_runtime)
+target_link_libraries(tensor_runtime_install ${LINK_LIBS})
+target_compile_definitions(tensor_runtime_install PRIVATE -DONLINE_PROFILING=false -DFP16_tuning=true)
 
 # Online version
-remove_definitions(-DONLINE_PROFILING=false)
-add_definitions(-DONLINE_PROFILING=true)
-remove_definitions(-DFP16_tuning=true)
-add_definitions(-DFP16_tuning=false)
 cuda_add_library(tensor_runtime_online ${RUNTIME_SRCS})
 cuda_add_cublas_to_target(tensor_runtime_online)
-target_link_libraries(tensor_runtime_online  ${LINK_LIBS})
 add_dependencies(tensor_runtime_online tensor_runtime)
+target_link_libraries(tensor_runtime_online ${LINK_LIBS}) 
+target_compile_definitions(tensor_runtime_online PRIVATE -DONLINE_PROFILING=true -DFP16_tuning=false)
 
 # Adding new rule for building a cuDNN runtime library
 
diff --git a/llvm/projects/hpvm-tensor-rt/dnn_sources/include/utils.h b/llvm/projects/hpvm-tensor-rt/dnn_sources/include/utils.h
index ef0d9fe12fbf4438af167784800dfdee845769d8..500ff63bc86dce6cae0dee3f942639c07bf14ab3 100644
--- a/llvm/projects/hpvm-tensor-rt/dnn_sources/include/utils.h
+++ b/llvm/projects/hpvm-tensor-rt/dnn_sources/include/utils.h
@@ -35,7 +35,7 @@ void printTensorInfo(void* tensor_ptr){
 
 
 // FIXIT: Move this to debug.h and include in all files
-void dumpWeightsToFile(char* file_name, void* weights_ptr){
+void dumpWeightsToFile(const char* file_name, void* weights_ptr){
 
   struct Tensor* weights = (Tensor*) weights_ptr;
   // Move data back to host
diff --git a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/half/lenet_mnist_half.cc b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/half/lenet_mnist_half.cc
index 123d3a4429e6b358042d2c33703f678bf0dfc973..f04ec041644394e2258414575162b961f9849667 100644
--- a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/half/lenet_mnist_half.cc
+++ b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/half/lenet_mnist_half.cc
@@ -13,18 +13,9 @@
 #include "utils.h"
 
 
-bool Opentuner_run = false;
-
-
 /* NOTE: Reference Architecture to use for profiling */
 void testLenetTanh(){
-
   int total_runs = 1;
-  if(Opentuner_run){
-    total_runs = 1000000;
-  }
-
-  
   printf("********* Lenet-2 Architecture ********** \n");
   // FIXIT: Extend this to batch of images - currently 5 images
 
@@ -62,28 +53,6 @@ void testLenetTanh(){
   clearTensorMap();
   
   for(int i = 0; i < total_runs; i++){
-
-    if(Opentuner_run){
-
-      char* myfifo = "/tmp/myfifo";
-      int fd = open(myfifo, O_RDONLY);
-
-      int ret_val = fcntl(fd, F_GETFD);
-      if(ret_val == -1){
-	printf("Invalid descriptor \n");
-	abort();
-      }
-
-      char str[100];
-      read(fd, str, 80);
-      if(strcmp(str, "stop_run") == 0){
-	abort();
-      }
-
-      close(fd);
-    }
-
-    
     readOpenTunerFlags("opentuner_flags"); // Resets the OpenTuner counters
 
     // Start power and performnce profiling 
@@ -132,23 +101,7 @@ void testLenetTanh(){
     computeAccuracy2(labels, test_batch_size, result);
     
     dumpAccuracyNorms();
-    freeOutputTensors();  
-
-    if(Opentuner_run){
-
-      char* myfifo = "/tmp/myfifo";
-      int fd_out = open(myfifo, O_WRONLY);
-      int ret_val = fcntl(fd_out, F_GETFD);
-      if(ret_val == -1){
-	printf("Invalid descriptor \n");
-	abort();
-      }
-      
-      const char* str = "completed***!\n\0";
-      write(fd_out, str, 80);
-      close(fd_out);
-    }
-    
+    freeOutputTensors();
   }
 
 
@@ -157,10 +110,6 @@ void testLenetTanh(){
 
 
 int main(int argc, char* argv[]){
-
-  if(argc > 1)
-    Opentuner_run = true;
-
   llvm_hpvm_initTensorRt(0);
 
   testLenetTanh();
diff --git a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/lenet_mnist.cc b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/lenet_mnist.cc
index 1b73645354ae70a53bd63bc864cae80420fdff1d..c047ffe090a93711cb66973ef6622d46fccdcee3 100644
--- a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/lenet_mnist.cc
+++ b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/lenet_mnist.cc
@@ -12,20 +12,11 @@
 #include "tensor_runtime.h"
 #include "utils.h"
 
-
-bool Opentuner_run = false;
-
 int total_runs = 1;
 
 
 /* NOTE: Reference Architecture to use for profiling */
 void testLenetTanh(){
-
-  if(Opentuner_run){
-    total_runs = 1000000;
-  }
-
-  
   printf("********* Lenet-2 Architecture ********** \n");
   // FIXIT: Extend this to batch of images - currently 5 images
 
@@ -64,28 +55,6 @@ void testLenetTanh(){
   clearTensorMap();
   
   for(int i = 0; i < total_runs; i++){
-
-    if(Opentuner_run){
-
-      const char* myfifo = "/tmp/myfifo";
-      int fd = open(myfifo, O_RDONLY);
-
-      int ret_val = fcntl(fd, F_GETFD);
-      if(ret_val == -1){
-	printf("Invalid descriptor \n");
-	abort();
-      }
-
-      char str[100];
-      read(fd, str, 80);
-      if(strcmp(str, "stop_run") == 0){
-	abort();
-      }
-
-      close(fd);
-    }
-
-    
     readOpenTunerFlags("opentuner_flags"); // Resets the OpenTuner counters
 
     // Start power and performnce profiling 
@@ -138,22 +107,6 @@ void testLenetTanh(){
     //FIXME: remove the comment below to use piped autotuner
     //dumpAccuracyNorms();
     freeOutputTensors();  
-
-    if(Opentuner_run){
-
-      const char* myfifo = "/tmp/myfifo";
-      int fd_out = open(myfifo, O_WRONLY);
-      int ret_val = fcntl(fd_out, F_GETFD);
-      if(ret_val == -1){
-	printf("Invalid descriptor \n");
-	abort();
-      }
-      
-      const char* str = "completed***!\n\0";
-      write(fd_out, str, 80);
-      close(fd_out);
-    }
-    
   }
 
   dumpExecutionAccuracies();
diff --git a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
index 1ccc5bd1b4e48a0019a78c6550b7ea9d1f679b91..0809c3616b17ef807d6155edea1da674272bd113 100644
--- a/llvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
+++ b/llvm/projects/hpvm-tensor-rt/dnn_sources/src/unit_tests.cc
@@ -733,7 +733,7 @@ void testPerforationCalls(void* input, void* filter,
   float interpolation_rate = 1.0;
   for (int offset = 0; offset < 2; offset++){
   
-      printf("\n\n\n\**Test -- pad_h = %d pad_w = %d stride_h = %d stride_w = %d row = %d col = %d  offset= %d \n\n",
+      printf("\n\n\n**Test -- pad_h = %d pad_w = %d stride_h = %d stride_w = %d row = %d col = %d  offset= %d \n\n",
 	     pad_h, pad_w, stride_h, stride_w, row, col, offset);
 
     
@@ -800,7 +800,7 @@ void testPerforationCalls(void* input, void* filter,
     }
   
  
-  printf ("\n\n\n\ --- End of Test \n\n\n");
+  printf ("\n\n\n--- End of Test \n\n\n");
 }
 
 
@@ -1019,8 +1019,9 @@ void testSamplingCalls(void* input, void* filter,
 
   float interpolation_rate = 1.0;
   for (int offset = 0; offset < 2; offset++){
- 
-      printf("\n\n\n\**Test -- pad_h = %d pad_w = %d stride_h = %d stride_w = %d skip_every = %d offset= %d interpolation_rate = %f \n\n",
+
+  
+      printf("\n\n\n**Test -- pad_h = %d pad_w = %d stride_h = %d stride_w = %d skip_every = %d offset= %d interpolation_rate = %f \n\n",
 	     pad_h, pad_w, stride_h, stride_w, skip_every, offset, interpolation_rate);
 
     
@@ -1078,7 +1079,7 @@ void testSamplingCalls(void* input, void* filter,
     }
   
  
-  printf ("\n\n\n\ --- End of Test \n\n\n");
+  printf ("\n\n\n --- End of Test \n\n\n");
 }
 
 
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/device_math.h b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/device_math.h
index 0f3cd1970e2aec5dcc4cbebeec4785a31fa1086a..947556489de2fb94ef7d793dd46eb74db2240516 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/device_math.h
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/device_math.h
@@ -43,6 +43,7 @@ template <typename T> T reduceOpToIdentity(MathOp op) {
   default:
     ERROR("Operator does not have id value\n");
   }
+  return T(); // For some compilers
 }
 
 template <> half reduceOpToIdentity<half>(MathOp op);
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/functional/common.h b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/functional/common.h
index 6a20ff83c2ae2bc0e264a7cc7559b89524c6b8b7..76a26d249c2ef1dc9db196f80cd99c1054e12603 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/functional/common.h
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/functional/common.h
@@ -26,8 +26,8 @@ template <typename T> class HFProfileGuard {
   static const char *getEventName(bool end) {
     if (typeid(T) == typeid(half) || typeid(T) == typeid(half2))
       return end ? "F2H_end" : "F2H_start";
-    else
-      ERROR("Type not accepted\n");
+    ERROR("Type not accepted\n");
+    return ""; // For some compilers
   }
 
   static bool needProfiling() {
@@ -55,8 +55,10 @@ template <typename T> int getTensorType() {
     return (int)float2_type;
   else if (typeid(T) == typeid(half2))
     return (int)half2_type;
-  else
+  else {
     ERROR("Unsupported type!\n");
+    return 0; // For some compilers
+  }
 }
 
 template <typename T> T *convertAndGetGPUData(Tensor *t);
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/approx_simulation.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/approx_simulation.cu
index 45876d4b3542ed68970146ec4e21dd313b246567..5051b780894b6e758cf768e663739ea6b92c71e5 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/approx_simulation.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/approx_simulation.cu
@@ -26,7 +26,7 @@
 #include <string>
 #include <vector>
 #include <map>
-
+#include <cassert>
 
 
 
@@ -1422,7 +1422,6 @@ void* Autotuner_ConvLayer(void* input, float i_min, float i_max,
       ERROR("Online Profiling cannot be enabled with PROMISE Simulation \n");
     }
   }
-  
 
   swing = getSwing(swing);  
  
@@ -1438,20 +1437,17 @@ void* Autotuner_ConvLayer(void* input, float i_min, float i_max,
 			out_min, out_max, swing);
   }
 
-  if (isGPULayer(swing)){
+  assert(isGPULayer(swing));
 
-    return Autotuner_GPU_ConvLayer(input, i_min, i_max,
-				   filter, w_min, w_max,
-				   bias, b_min, b_max,
-				   conv_pad_h, conv_pad_w,
-				   conv_stride_h, conv_stride_w,
-				   pool_id, pool_size, pool_stride, 
-				   activation_id,
-				   out_min, out_max, swing);
-
-  }
+  return Autotuner_GPU_ConvLayer(input, i_min, i_max,
+			   filter, w_min, w_max,
+			   bias, b_min, b_max,
+			   conv_pad_h, conv_pad_w,
+			   conv_stride_h, conv_stride_w,
+			   pool_id, pool_size, pool_stride, 
+			   activation_id,
+			   out_min, out_max, swing);
 
-  
 }
 
 
@@ -1531,58 +1527,50 @@ void* FCLayer_PROMISE(void* input, float i_min, float i_max,
 		      out_min, out_max, swing);
   }
 
+  assert(isGPULayer(swing));
 
-  if(isGPULayer(swing)){
-  
-    void* gemm_out;
-    if(!isFullPrecision(swing)){
-      gemm_out = tensorHalfGemm(input, weights);
-    }
-    else{
-      gemm_out = tensorGemmGPU(input, weights);
-    }
+  void *gemm_out;
+  if (!isFullPrecision(swing)) {
+    gemm_out = tensorHalfGemm(input, weights);
+  } else {
+    gemm_out = tensorGemmGPU(input, weights);
+  }
 
-  
-    void* gemmbias_out;
-    if(bias != NULL){
-      // Swing 8 corresponds to FP32
-      if( isFullPrecision(swing) || (!FP16_tuning) ){
-	gemmbias_out = tensorAdd(gemm_out, bias);
-      }
-      else{
-	gemmbias_out = tensorHalfAdd(gemm_out, bias);
-      }
-    }
-    else{
-      gemmbias_out = gemm_out;
+  void *gemmbias_out;
+  if (bias != NULL) {
+    // Swing 8 corresponds to FP32
+    if (isFullPrecision(swing) || (!FP16_tuning)) {
+      gemmbias_out = tensorAdd(gemm_out, bias);
+    } else {
+      gemmbias_out = tensorHalfAdd(gemm_out, bias);
     }
- 
-    void* activation_out;
-    switch(activation_id){
+  } else {
+    gemmbias_out = gemm_out;
+  }
 
-    case -1:
-      activation_out = gemmbias_out;
-      INFO("No Activation Function \n");
-      break;
-    case 0:
-      activation_out = tensorTanh(gemmbias_out);
-      break;
-    case 1:
-      activation_out = tensorRelu(gemmbias_out);
-      break;
-    case 2:
-      activation_out = tensorRelu2(gemmbias_out, out_min, out_max);
-      break;
-    default:
-      ERROR("Activation id %d NOT supported \n", activation_out);
-      break;
-    }
-    
-    return activation_out;
+  void *activation_out;
+  switch (activation_id) {
 
+  case -1:
+    activation_out = gemmbias_out;
+    INFO("No Activation Function \n");
+    break;
+  case 0:
+    activation_out = tensorTanh(gemmbias_out);
+    break;
+  case 1:
+    activation_out = tensorRelu(gemmbias_out);
+    break;
+  case 2:
+    activation_out = tensorRelu2(gemmbias_out, out_min, out_max);
+    break;
+  default:
+    ERROR("Activation id %d NOT supported \n", activation_out);
+    break;
   }
 
-  
+  return activation_out;
+
 }
 
 #endif
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/device_math.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/device_math.cu
index 002c14b3c72a8f9947195f2dd5f930f4868a6708..163efd1a2dec403146ad83c1e00ef1cffa48222f 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/device_math.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/device_math.cu
@@ -119,6 +119,7 @@ template <> void *mathOpToFunc<float2>(MathOp op) {
     CASE_FUNC(Mul, f2mul)
   default:
     ERROR("Float2 function not found\n");
+    return nullptr;  // For some compilers
   }
 }
 
@@ -127,6 +128,7 @@ template <> void *mathOpToFunc<half2>(MathOp op) {
     CASE_FUNC(Mul, h2mul)
   default:
     ERROR("Half2 function not found\n");
+    return nullptr;  // For some compilers
   }
 }
 
@@ -148,6 +150,7 @@ template <> void *mathOpToFunc<float>(MathOp op) {
   default:
     ERROR("Float function not found\n");
   }
+  return nullptr;  // For some compilers
 }
 
 template <> void *mathOpToFunc<half>(MathOp op) {
@@ -165,6 +168,7 @@ template <> void *mathOpToFunc<half>(MathOp op) {
   default:
     ERROR("Half function not found\n");
   }
+  return nullptr;  // For some compilers
 }
 
 template <> half reduceOpToIdentity<half>(MathOp op) {
@@ -180,4 +184,5 @@ template <> half reduceOpToIdentity<half>(MathOp op) {
   default:
     ERROR("Operator does not have id value\n");
   }
+  return 0.0f;  // For some compilers
 }
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_runtime.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_runtime.cu
index c2cc1ef5f7a7e86d9265057fd8ff0f17071a705f..9814e2c6a371b423a725c533c0029425f209fe1d 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_runtime.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_runtime.cu
@@ -497,8 +497,10 @@ void *tensorMap2Half(MathOp f2, void *i1, void *i2) {
     profileEvent("#tensorMap2_end");
     return ret;
   }
-  else
+  else {
     ERROR("Type not recognized\n");
+    return nullptr; // For some compilers
+  }
 }
 
 void *tensorMap3(MathOp f3, void *i1, void *i2, void *i3) {
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_utils.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_utils.cu
index 74a87faa2d13839b75f5d57d564fd31d0d45e2f9..6bbccfabaf22a395e91748be22e1eaddcf32c0ba 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_utils.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_utils.cu
@@ -70,8 +70,9 @@ int getTypeSize(int data_type){
     case half2_type:
       return 4;
     default:
-      ERROR("Unknown type %s\n", std::to_string(data_type));
+      ERROR("Unknown type %d\n", data_type);
   }
+  return 0;
 }
 
 static int getFullPrecTypeSize(int data_type){
@@ -90,6 +91,7 @@ static int getFullPrecTypeSize(int data_type){
     default:
       ERROR("Unknown type %d\n", data_type);
   }
+  return 0;
 }
 
 static bool isFP16Compound(int data_type) {