From 1f56a6938d7e9be9def4242cb80bb3428dc30c6c Mon Sep 17 00:00:00 2001
From: Hashim Sharif <hsharif3@tyler.cs.illinois.edu>
Date: Sun, 5 Apr 2020 00:31:47 -0500
Subject: [PATCH] Creating Autotuner initializer for one-time initializations
---
.../include/approx_simulation.h | 3 +++
.../tensor_runtime/src/approx_simulation.cu | 27 +++++++++++++++----
.../tensor_runtime/src/tensor_runtime.cu | 8 ++++--
3 files changed, 31 insertions(+), 7 deletions(-)
create mode 100644 llvm/projects/hpvm-tensor-rt/tensor_runtime/include/approx_simulation.h
diff --git a/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/approx_simulation.h b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/approx_simulation.h
new file mode 100644
index 0000000000..25fc5dc7fa
--- /dev/null
+++ b/llvm/projects/hpvm-tensor-rt/tensor_runtime/include/approx_simulation.h
@@ -0,0 +1,3 @@
+
+
+void initializeAutotuner();
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 6a408ef059..1933480422 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
@@ -21,6 +21,15 @@
#include "device_math.h"
+
+
+
+
+
+
+
+
+
//N is new_data's size
//n, c, h, w are the dimensions of new_data
__global__
@@ -717,7 +726,6 @@ int getSwing(int swing){
-
class PerfParams{
public:
@@ -857,6 +865,7 @@ public:
};
+
class RedSampParams {
@@ -905,8 +914,18 @@ RedSampParams getRedSampParams(int swing) {
bool FP16_tuning = false;
+
/***** API for Autotuner Use - Not the ApproxHPVM Wrapper API */
+PerfParamSet* perfParamSet;
+SampParamSet* sampParamSet;
+
+
+void initializeAutotuner(){
+
+ sampParamSet = new SampParamSet;
+ perfParamSet = new PerfParamSet;
+}
void* Autotuner_SampConv(void* input, float i_min, float i_max,
@@ -919,8 +938,7 @@ void* Autotuner_SampConv(void* input, float i_min, float i_max,
float out_min, float out_max, int swing){
- SampParamSet* paramSet = new SampParamSet;
- SampParams params = paramSet->getSampParams(swing);
+ SampParams params = sampParamSet->getSampParams(swing);
DEBUG("params.skip_rate = %d, params.skip_offset = %d \n",
params.skip_rate, params.skip_offset);
@@ -969,8 +987,7 @@ void* Autotuner_PerforatedConv(void* input, float i_min, float i_max,
float out_min, float out_max, int swing){
- PerfParamSet* paramSet = new PerfParamSet;
- PerfParams params = paramSet->getPerfParams(swing);
+ PerfParams params = perfParamSet->getPerfParams(swing);
DEBUG("params.row = %d, params.col = %d, params.skip_offset = %d \n",
params.row, params.col, params.skip_offset);
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 442e0b802a..a80b2f726f 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
@@ -43,6 +43,7 @@
#include "tensor.h"
#include "op_overheads.h"
#include "half_precision_api.h"
+#include "approx_simulation.h"
//** Potential Improvements:
// 1) Add support for dataypes beyond floats and half
@@ -65,6 +66,8 @@ void llvm_hpvm_initTensorRt(int gpuid){
#ifdef PROMISE_TUNER_ENABLED
// readOpenTunerFlags("opentuner_flags");
readOpenTunerFlags("promise_flags");
+ initializeAutotuner();
+
#endif
@@ -571,7 +574,7 @@ void* tensorConvolution(void* input_ptr, void* filter_ptr,
op_counter++;
#endif
-
+
return output;
}
@@ -608,7 +611,8 @@ void* tensorPooling(void* input_ptr,
int w = (input->dims.dim_sizes[3] + (2 * horizontal_pad) - window_width) / horizontal_stride;
w = w + 1;
- DEBUG("n = %d, c = %d, h = %d, w = %d \n", n, c, h, w);
+ DEBUG("n = %d, c = %d, h = %d, w = %d , dim1 = %d , dim2 = %d \n",
+ n, c, h, w, input->dims.dim_sizes[2], input->dims.dim_sizes[3]);
// FIXIT: Don't be specific to floats
Tensor* output = (Tensor*) create4DTensor(CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, n, c, h, w);
--
GitLab