Skip to content
Snippets Groups Projects
Commit 1f56a693 authored by Hashim Sharif's avatar Hashim Sharif
Browse files

Creating Autotuner initializer for one-time initializations

parent 899bdc22
No related branches found
No related tags found
No related merge requests found
void initializeAutotuner();
...@@ -21,6 +21,15 @@ ...@@ -21,6 +21,15 @@
#include "device_math.h" #include "device_math.h"
//N is new_data's size //N is new_data's size
//n, c, h, w are the dimensions of new_data //n, c, h, w are the dimensions of new_data
__global__ __global__
...@@ -717,7 +726,6 @@ int getSwing(int swing){ ...@@ -717,7 +726,6 @@ int getSwing(int swing){
class PerfParams{ class PerfParams{
public: public:
...@@ -857,6 +865,7 @@ public: ...@@ -857,6 +865,7 @@ public:
}; };
class RedSampParams { class RedSampParams {
...@@ -905,8 +914,18 @@ RedSampParams getRedSampParams(int swing) { ...@@ -905,8 +914,18 @@ RedSampParams getRedSampParams(int swing) {
bool FP16_tuning = false; bool FP16_tuning = false;
/***** API for Autotuner Use - Not the ApproxHPVM Wrapper API */ /***** 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, 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, ...@@ -919,8 +938,7 @@ void* Autotuner_SampConv(void* input, float i_min, float i_max,
float out_min, float out_max, int swing){ float out_min, float out_max, int swing){
SampParamSet* paramSet = new SampParamSet; SampParams params = sampParamSet->getSampParams(swing);
SampParams params = paramSet->getSampParams(swing);
DEBUG("params.skip_rate = %d, params.skip_offset = %d \n", DEBUG("params.skip_rate = %d, params.skip_offset = %d \n",
params.skip_rate, params.skip_offset); params.skip_rate, params.skip_offset);
...@@ -969,8 +987,7 @@ void* Autotuner_PerforatedConv(void* input, float i_min, float i_max, ...@@ -969,8 +987,7 @@ void* Autotuner_PerforatedConv(void* input, float i_min, float i_max,
float out_min, float out_max, int swing){ float out_min, float out_max, int swing){
PerfParamSet* paramSet = new PerfParamSet; PerfParams params = perfParamSet->getPerfParams(swing);
PerfParams params = paramSet->getPerfParams(swing);
DEBUG("params.row = %d, params.col = %d, params.skip_offset = %d \n", DEBUG("params.row = %d, params.col = %d, params.skip_offset = %d \n",
params.row, params.col, params.skip_offset); params.row, params.col, params.skip_offset);
......
...@@ -43,6 +43,7 @@ ...@@ -43,6 +43,7 @@
#include "tensor.h" #include "tensor.h"
#include "op_overheads.h" #include "op_overheads.h"
#include "half_precision_api.h" #include "half_precision_api.h"
#include "approx_simulation.h"
//** Potential Improvements: //** Potential Improvements:
// 1) Add support for dataypes beyond floats and half // 1) Add support for dataypes beyond floats and half
...@@ -65,6 +66,8 @@ void llvm_hpvm_initTensorRt(int gpuid){ ...@@ -65,6 +66,8 @@ void llvm_hpvm_initTensorRt(int gpuid){
#ifdef PROMISE_TUNER_ENABLED #ifdef PROMISE_TUNER_ENABLED
// readOpenTunerFlags("opentuner_flags"); // readOpenTunerFlags("opentuner_flags");
readOpenTunerFlags("promise_flags"); readOpenTunerFlags("promise_flags");
initializeAutotuner();
#endif #endif
...@@ -571,7 +574,7 @@ void* tensorConvolution(void* input_ptr, void* filter_ptr, ...@@ -571,7 +574,7 @@ void* tensorConvolution(void* input_ptr, void* filter_ptr,
op_counter++; op_counter++;
#endif #endif
return output; return output;
} }
...@@ -608,7 +611,8 @@ void* tensorPooling(void* input_ptr, ...@@ -608,7 +611,8 @@ void* tensorPooling(void* input_ptr,
int w = (input->dims.dim_sizes[3] + (2 * horizontal_pad) - window_width) / horizontal_stride; int w = (input->dims.dim_sizes[3] + (2 * horizontal_pad) - window_width) / horizontal_stride;
w = w + 1; 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 // FIXIT: Don't be specific to floats
Tensor* output = (Tensor*) create4DTensor(CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, n, c, h, w); Tensor* output = (Tensor*) create4DTensor(CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, n, c, h, w);
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment