From e862382c38562a47f09e98a7c93642409bdece09 Mon Sep 17 00:00:00 2001
From: Yifan Zhao <yifanz16@illinois.edu>
Date: Thu, 13 Feb 2020 15:56:17 -0600
Subject: [PATCH] Fixed some warnings in tensor runtime

---
 .../tensor_runtime/include/image/stb_image.h         |  6 +++---
 .../tensor_runtime/include/op_overheads.h            |  2 +-
 .../tensor_runtime/src/approx_simulation.cu          | 12 ++++++------
 .../hpvm-tensor-rt/tensor_runtime/src/error.cu       |  2 +-
 .../tensor_runtime/src/img_tensor_utils.cpp          |  2 +-
 .../tensor_runtime/src/op_overheads.cc               |  2 +-
 .../tensor_runtime/src/tensor_runtime.cu             |  2 +-
 7 files changed, 14 insertions(+), 14 deletions(-)

diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/image/stb_image.h b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/image/stb_image.h
index 196dfd5cc6..f4b2b17f4b 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/image/stb_image.h
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/image/stb_image.h
@@ -4012,7 +4012,7 @@ stbi_inline static int stbi__zhuffman_decode(stbi__zbuf *a, stbi__zhuffman *z)
 static int stbi__zexpand(stbi__zbuf *z, char *zout, int n)  // need to make room for n bytes
 {
    char *q;
-   int cur, limit, old_limit;
+   int cur, limit, old_limit __attribute__((unused));
    z->zout = zout;
    if (!z->z_expandable) return stbi__err("output buffer limit","Corrupt PNG");
    cur   = (int) (z->zout     - z->zout_start);
@@ -4891,7 +4891,7 @@ static int stbi__parse_png_file(stbi__png *z, int scan, int req_comp)
             if (scan == STBI__SCAN_header) { s->img_n = pal_img_n; return 1; }
             if ((int)(ioff + c.length) < (int)ioff) return 0;
             if (ioff + c.length > idata_limit) {
-               stbi__uint32 idata_limit_old = idata_limit;
+               stbi__uint32 idata_limit_old __attribute__((unused)) = idata_limit;
                stbi_uc *p;
                if (idata_limit == 0) idata_limit = c.length > 4096 ? c.length : 4096;
                while (ioff + c.length > idata_limit)
@@ -5111,7 +5111,7 @@ static int stbi__shiftsigned(unsigned int v, int shift, int bits)
       v <<= -shift;
    else
       v >>= shift;
-   STBI_ASSERT(v >= 0 && v < 256);
+   STBI_ASSERT(v < 256);
    v >>= (8-bits);
    STBI_ASSERT(bits >= 0 && bits <= 8);
    return (int) ((unsigned) v * mul_table[bits]) >> shift_table[bits];
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/op_overheads.h b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/op_overheads.h
index d2b8f89d30..efa5c5e92e 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/op_overheads.h
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/op_overheads.h
@@ -49,7 +49,7 @@ extern "C"{
 
   void add_norms(void* norms_ptr, char* op_name, int error_value);
 
-  void dump_result(char* file_name);
+  void dump_result(const char* file_name);
 
 }
 
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 c5f1d1af70..95c1604fa2 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
@@ -1089,7 +1089,7 @@ void *autotuner_tensorFft(void *input, bool inverse) {
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
@@ -1110,7 +1110,7 @@ void *autotuner_tensorReduce(void *input, size_t axis, MathOp func) {
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
@@ -1141,7 +1141,7 @@ void *autotuner_tensorProjectiveT(void *input, void *transformation) {
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
@@ -1159,7 +1159,7 @@ void *autotuner_tensorMap1(MathOp func, void *input) {
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
@@ -1181,7 +1181,7 @@ void *autotuner_tensorMap2(MathOp func, void *input1, void *input2) {
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
@@ -1204,7 +1204,7 @@ void *autotuner_tensorMap3(MathOp func, void *input1, void *input2,
     abort();
   }
 
-  int swing;
+  int swing = 0;
   swing = getSwing(swing);
 
   if (isFullPrecision(swing)) {
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
index 2d370eb4d0..6b8ee15a42 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/error.cu
@@ -293,7 +293,7 @@ __global__ void normComputeKernel(float* A, float * B, double* l1_A, double* l2_
 __inline__ __device__ double warpReduceSum(double val) {
 
   for (int offset = warpSize/2; offset > 0; offset /= 2)
-    val += __shfl_down(val, offset);
+    val += __shfl_down_sync(0xFFFFFFFF, val, offset);
 
   return val;
 }
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_utils.cpp b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_utils.cpp
index de254b2188..07f2a5e692 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_utils.cpp
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/img_tensor_utils.cpp
@@ -395,7 +395,7 @@ float violationRate(
   for (float v : values) {
     if (std::isnan(v))
       ++violation;
-    if (higher_better && v < threshold || !higher_better && v > threshold)
+    if ((higher_better && v < threshold) || (!higher_better && v > threshold))
       ++violation;
   }
   return (float)violation / values.size();
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/op_overheads.cc b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/op_overheads.cc
index e14920868a..8f673f07ac 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/op_overheads.cc
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/op_overheads.cc
@@ -213,7 +213,7 @@ void add_norms(void* norms_ptr, char* op_name, int error_value){
 }
 
 
-void dump_result(char* file_name){
+void dump_result(const char* file_name){
 
   FILE* fp = fopen(file_name, "w+");
   if(fp != NULL){
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
index a5d14909cf..6e97043591 100644
--- a/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/src/tensor_runtime.cu
@@ -85,7 +85,7 @@ void llvm_hpvm_initTensorRt(int gpuid){
 
 
 void llvm_hpvm_cleanupTensorRt(){
-  DEBUG("\**** llvm_hpvm_cleanupTensorRt ***\n");
+  DEBUG("**** llvm_hpvm_cleanupTensorRt ***\n");
   dumpAccuracyNorms();
 }
 
-- 
GitLab