Removed support for PROMISE configs and quant file

hpvm/lib/Transforms/DFG2LLVM_WrapperAPI/DFG2LLVM_WrapperAPI.cpp

@@ -49,11 +49,6 @@ using namespace inplacedfg;
 
 namespace {
 
-cl::opt<std::string> QuantizationInputsFilename(
-    "quantization-levels-filename",
-    cl::desc("<PROMISE quantization levels input file (path)>"),
-    cl::value_desc("filename"), cl::Required);
-
 cl::opt<std::string> ConfigurationInputsFilename(
     "configuration-inputs-filename",
     cl::desc("<Autotuner configurations input file (path)>"),
@@ -1241,8 +1236,6 @@ class CGT_WrapperAPI : public CodeGenTraversal {
 private:
   // Member variables
   unsigned nodeID; // Used as a node identifier
-
-  std::string QuantizationInputsFilenameStr;
   std::string ConfigurationInputsFilenameStr;
 
   InPlaceDFGAnalysis::InPlaceDFGParameter *IPP;
@@ -1267,10 +1260,8 @@ public:
   // Constructor
   CGT_WrapperAPI(Module &_M, BuildDFG &_DFG,
                  InPlaceDFGAnalysis::InPlaceDFGParameter &_IPP,
-                 std::string &_QuantizationInputsFilenameStr,
                  std::string &_ConfigurationInputsFilenameStr)
       : CodeGenTraversal(_M, _DFG), IPP(&_IPP),
-        QuantizationInputsFilenameStr(_QuantizationInputsFilenameStr),
         ConfigurationInputsFilenameStr(_ConfigurationInputsFilenameStr) {
     nodeID = 0;
     initRuntimeAPI();
@@ -1314,19 +1305,6 @@ void CGT_WrapperAPI::initRuntimeAPI() {
       ArrayRef<Value *>(ConstantInt::get(Type::getInt32Ty(M.getContext()), 0)),
       "", InitCall);
 
-  StringRef QRangesStrRef = StringRef(QuantizationInputsFilenameStr);
-  // Create string for node name, as first argument for wrapper API call
-  Constant *ConstArray1 =
-      ConstantDataArray::getString(M.getContext(), QRangesStrRef, true);
-  GlobalVariable *GV1 =
-      new GlobalVariable(M, ConstArray1->getType(), true,
-                         GlobalValue::ExternalLinkage, ConstArray1, "");
-  // Create GEP expression to access it
-  Constant *Int_0 = ConstantInt::get(Type::getInt32Ty(M.getContext()), 0);
-  Constant *GEPIndices[] = {Int_0, Int_0};
-  Constant *QRangesGEPConst = ConstantExpr::getGetElementPtr(
-      GV1->getType()->getPointerElementType(), GV1, GEPIndices);
-
   StringRef ConfsStrRef = StringRef(ConfigurationInputsFilenameStr);
   // Create string for node name, as first argument for wrapper API call
   Constant *ConstArray2 =
@@ -1334,10 +1312,12 @@ void CGT_WrapperAPI::initRuntimeAPI() {
   GlobalVariable *GV2 =
       new GlobalVariable(M, ConstArray2->getType(), true,
                          GlobalValue::ExternalLinkage, ConstArray2, "");
+  // Create GEP expression to access it
+  Constant *Int_0 = ConstantInt::get(Type::getInt32Ty(M.getContext()), 0);
+  Constant *GEPIndices[] = {Int_0, Int_0};
   Constant *ConfsGEPConst = ConstantExpr::getGetElementPtr(
       GV2->getType()->getPointerElementType(), GV2, GEPIndices);
-  Value *RTCInitArgs[] = {ConfsGEPConst, QRangesGEPConst};
-  CallInst::Create(llvm_hpvm_initializeRuntimeController, RTCInitArgs, "",
+  CallInst::Create(llvm_hpvm_initializeRuntimeController, {ConfsGEPConst}, "",
                    InitCall);
 
   Function *VC = M.getFunction("llvm.hpvm.cleanup");
@@ -1464,7 +1444,8 @@ bool DFG2LLVM_WrapperAPI::runOnModule(Module &M) {
 
   // Visitor for Code Generation Graph Traversal
   CGT_WrapperAPI *CGTVisitor = new CGT_WrapperAPI(
-      M, DFG, IPP, QuantizationInputsFilename, ConfigurationInputsFilename);
+    M, DFG, IPP, ConfigurationInputsFilename
+  );
 
   // Iterate over all the DFGs and produce code for each one of them
   for (auto rootNode : Roots) {

hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/configuration.h

@@ -11,14 +11,12 @@
 // Describes the internal choices made for an ApproxHPVM node
 class NodeConfiguration {
 public:
-  enum NODE_CONFIGURATION_TARGET { PROMISE, GPU, CPU, END };
+  enum NODE_CONFIGURATION_TARGET { GPU, CPU, END };
 
 protected:
   enum NODE_CONFIGURATION_TARGET NODE_CONFIGURATION_TARGET_ID;
 
 public:
-  bool isPROMISENodeConfiguration();
-
   bool isGPUNodeConfiguration();
 
   bool isCPUNodeConfiguration();
@@ -26,27 +24,6 @@ public:
   virtual void print() = 0;
 };
 
-class PROMISENodeConfiguration : public NodeConfiguration {
-public:
-  // Approximation methods available for this HW type
-  enum APPROX { SWING_LEVEL, END };
-
-private:
-  // A vector, containing pairs of approximation method and tunable parameter
-  // (expressed as int, or ignored when not applicable)
-  std::vector<std::pair<enum APPROX, int>> ApproxChoices;
-
-public:
-  void pushNewApproximationChoice(enum APPROX approx, int u);
-
-  std::vector<std::pair<enum APPROX, int>> &getApproxChoices();
-
-  PROMISENodeConfiguration();
-  ~PROMISENodeConfiguration();
-
-  void print() override;
-};
-
 class GPUNodeConfiguration : public NodeConfiguration {
 public:
   // Approximation methods available for this HW type

hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/hpvm-rt-controller.h

@@ -155,8 +155,6 @@ extern RuntimeController *RC;
 class RuntimeController {
 private:
   // Members
-  // Map from node names to quantization ranges
-  std::map<std::string, std::vector<float>> QuantizationMap;
 
   // Configurations.
   // Configurations initially read - all generated from autotuner
@@ -203,7 +201,6 @@ private:
   void stop_profiler();
 
   void setProfileInfoFilename(const char *);
-  void readQuantizationFile(const char *);
   void readConfigurationFile(const char *);
 
   void computeParetoConfigurationPoints();
@@ -224,8 +221,6 @@ public:
   double getCurrentConfigurationAccuracy();
   double getCurrentConfigurationAccuracyLoss();
 
-  std::vector<float> &getQuantizationRanges(const char *data);
-
   NodeConfiguration *getNodeConfiguration(const char *data);
 
   // Functions for runtime control
@@ -241,7 +236,7 @@ public:
   double getBaselineTime();
   Slowdowns *getSlowdowns();
 
-  void init(const char *Cstr, const char *Qstr);
+  void init(const char *Cstr);
 
   // Exposing functionality of ProfileInfo
   void end_iteration();
@@ -298,7 +293,6 @@ public:
   ~RuntimeController();
 
   // Helper Functions
-  void printQuantizationMap();
   void printConfigurations(std::vector<struct Configuration> &);
   void printConfigurations(std::vector<struct Configuration *> &);
 };

hpvm/projects/hpvm-tensor-rt/tensor_runtime/include/rt-controller-api.h

 extern "C" {
 // Functions to be inserted with initializeTensorRT and clearTensorRT
-void llvm_hpvm_initializeRuntimeController(const char *, const char *);
+void llvm_hpvm_initializeRuntimeController(const char *);
 void llvm_hpvm_clearRuntimeController();
 void llvm_hpvm_invokeRtControl(void *result, const char *str, int start,
                                int end);

hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/configuration.cpp

@@ -10,16 +10,11 @@
 
 #include "configuration.h"
 
-using P_APPROX = PROMISENodeConfiguration::APPROX;
 using G_APPROX = GPUNodeConfiguration::APPROX;
 using C_APPROX = CPUNodeConfiguration::APPROX;
 using G_TENSOR_OP = GPUNodeConfiguration::TENSOR_OP;
 using C_TENSOR_OP = CPUNodeConfiguration::TENSOR_OP;
 
-bool NodeConfiguration::isPROMISENodeConfiguration() {
-  return NODE_CONFIGURATION_TARGET_ID == PROMISE;
-}
-
 bool NodeConfiguration::isGPUNodeConfiguration() {
   return NODE_CONFIGURATION_TARGET_ID == GPU;
 }
@@ -28,22 +23,6 @@ bool NodeConfiguration::isCPUNodeConfiguration() {
   return NODE_CONFIGURATION_TARGET_ID == CPU;
 }
 
-void PROMISENodeConfiguration::pushNewApproximationChoice(
-    P_APPROX approx, int u) {
-  ApproxChoices.push_back(std::make_pair(approx, u));
-}
-
-std::vector<std::pair<P_APPROX, int>> &
-PROMISENodeConfiguration::getApproxChoices() {
-  return ApproxChoices;
-}
-
-PROMISENodeConfiguration::PROMISENodeConfiguration() {
-  NODE_CONFIGURATION_TARGET_ID = PROMISE;
-}
-
-PROMISENodeConfiguration::~PROMISENodeConfiguration() {}
-
 void GPUNodeConfiguration::pushNewTensorOperation(G_TENSOR_OP top) {
   std::vector<std::pair<G_APPROX, int>> emptyVec;
   ApproxChoices.push_back(std::make_pair(top, emptyVec));
@@ -124,26 +103,6 @@ operator()(const struct Configuration *a, const float &b) const {
 
 // Helper configuration print methods
 
-void PROMISENodeConfiguration::print() {
-
-  printf(" promise");
-  for (auto &it : ApproxChoices) {
-    printf(" ");
-    switch (it.first) {
-    case P_APPROX::SWING_LEVEL:
-      printf("swing_level");
-      break;
-      // TODO additional approx methods to be printed here
-    default:
-      ERROR("Unknown approximation option");
-      break;
-    }
-    printf(" %d", it.second);
-  }
-
-  printf("\n");
-}
-
 void GPUNodeConfiguration::print() {
 
   printf(" gpu");

hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/hpvm-rt-controller.cpp

@@ -420,12 +420,6 @@ double RuntimeController::getCurrentConfigurationAccuracyLoss() {
   return (double) (*Configurations)[configurationIdx]->accuracyLoss;
 }
 
-std::vector<float> &RuntimeController::getQuantizationRanges(const char *data) {
-  std::string s(data);
-  // All nodes are expected to have quantization ranges
-  return QuantizationMap.at(s);
-}
-
 NodeConfiguration *RuntimeController::getNodeConfiguration(const char *data) {
 
   // if visc.node.id Not specified for this HPVM Node
@@ -441,12 +435,10 @@ NodeConfiguration *RuntimeController::getNodeConfiguration(const char *data) {
   
 }
 
-void RuntimeController::init(const char *Cstr, const char *Qstr) {
+void RuntimeController::init(const char *Cstr) {
   // We initialize the path to the profile info output file,
   // based on the path given for the configuration file
   setProfileInfoFilename(Cstr);
-
-  readQuantizationFile(Qstr);
   readConfigurationFile(Cstr);
 
   // NOTE: Configurations is pareto-configs. InitialConfigurations is the full list (config file)
@@ -668,59 +660,6 @@ void RuntimeController::setProfileInfoFilename(const char *str) {
     PI->set_out_file_name(profile_filename);
   }
 }
-void RuntimeController::readQuantizationFile(const char *str) {
-
-  INFO("Reading Quantization Ranges File...\n");
-
-  if (std::string(str).empty()) {
-    INFO("Empty quantization file string.\n");
-    return;
-  }
-
-  std::ifstream qin(str);
-
-  if (!qin) {
-    ERROR("Failed to open PROMISE quantization file.");
-    abort();
-  }
-
-  while (!qin.eof()) {
-    char NodeName[NODE_NAME_BUFFER_SIZE];
-    std::vector<float> QuantRangeVector;
-
-    qin >> NodeName;
-
-    float qrange;
-    for (unsigned i = 0; i < 8; i++) {
-      qin >> qrange;
-      QuantRangeVector.push_back(qrange);
-    }
-    // See if we need to insert this in map instead - my lookup test seemed to
-    // work without it std::string s(NodeName);
-    QuantizationMap.insert(
-        std::pair<std::string, std::vector<float>>(NodeName, QuantRangeVector));
-  }
-
-  qin.close();
-  INFO("DONE.\n");
-}
-
-void RuntimeController::printQuantizationMap() {
-
-  DEBUG("Quantization Ranges Map:\n");
-
-  for (std::map<std::string, std::vector<float>>::const_iterator it =
-           QuantizationMap.begin();
-       it != QuantizationMap.end(); ++it) {
-    DEBUG("%s :", it->first.c_str());
-
-    for (unsigned i = 0; i < it->second.size(); i++) {
-      DEBUG(" %f", it->second[i]);
-    }
-
-    DEBUG("\n");
-  }
-}
 
 void RuntimeController::readConfigurationFile(const char *str) {
 
@@ -790,31 +729,7 @@ void RuntimeController::readConfigurationFile(const char *str) {
       continue;
     }
 
-    if (tokens[1] == "promise") {
-      DEBUG("Found promise configuration\n");
-
-      // There must be at least one approximation option
-      CUSTOM_ASSERT(
-          (tokens.size() >= 2) && "Not enough approximation options.");
-
-      PROMISENodeConfiguration *NodeConf = new PROMISENodeConfiguration();
-      InitialConfigurations.back().setup.insert(
-          std::make_pair(tokens[0], NodeConf));
-
-      // In increments of two, to handle pairs of approx option - tunable
-      // parameter
-      for (unsigned idx = 2; idx < tokens.size(); idx += 2) {
-        if (tokens[idx] == "swing_level") {
-          DEBUG("Found swing voltage option\n");
-          int vswing = std::stoi(tokens[idx + 1]);
-          DEBUG("vswing: %d\n", vswing);
-          NodeConf->pushNewApproximationChoice(
-              PROMISENodeConfiguration::APPROX::SWING_LEVEL, vswing);
-        }
-        // TODO: other approximation options handled here
-      }
-
-    } else if (tokens[1] == "gpu") {
+    if (tokens[1] == "gpu") {
       DEBUG("Found gpu configuration\n");
 
       // There must be at least one operation, with an approximation option
@@ -1483,10 +1398,9 @@ double RuntimeController::getBaselineTime() { return baseline_time; }
 Slowdowns *RuntimeController::getSlowdowns() { return slowdowns; }
 
 // Functions to be inserted with initializeTensorRT and clearTensorRT
-extern "C" void llvm_hpvm_initializeRuntimeController(
-    const char *ConfigFile, const char *QRangeFile) {
+extern "C" void llvm_hpvm_initializeRuntimeController(const char *ConfigFile) {
   RC = new RuntimeController();
-  RC->init(ConfigFile, QRangeFile);
+  RC->init(ConfigFile);
   return;
 }
 

hpvm/projects/hpvm-tensor-rt/tensor_runtime/src/wrapper_runtime.cu

@@ -59,68 +59,7 @@ extern "C"{
 
     NodeConfiguration *NodeConf = RC->getNodeConfiguration(hpvm_node_id);
 
-    if (NodeConf->isPROMISENodeConfiguration()) {
-      DEBUG("PROMISE Configuration for ConvLayer\n");
-      // Mapped to PROMISE - get a PROMISE node configuration
-      PROMISENodeConfiguration *PROMISEConf = (PROMISENodeConfiguration *)NodeConf;
-      std::vector<float> &QRanges = RC->getQuantizationRanges(hpvm_node_id);
-
-      std::vector<std::pair<PROMISENodeConfiguration::APPROX, int> > &approxTuples =
-	PROMISEConf->getApproxChoices();
-
-      if (approxTuples.size() == 1) {
-	enum PROMISENodeConfiguration::APPROX approx = approxTuples[0].first;
-	int param = approxTuples[0].second;
-	if (approx == PROMISENodeConfiguration::APPROX::SWING_LEVEL) {
-	  DEBUG("Approximation choice for ConvLayer: swing level %d\n", param);
-
-	  struct Tensor* input_tensor_cast = (struct Tensor*) input;
-	  struct Tensor* filter_tensor_cast = (struct Tensor*) filter;
-	  std::pair<double, double> pinfo =
-	    RC->conv_profile(input_tensor_cast->dims.dim_sizes[0], //n
-			     input_tensor_cast->dims.dim_sizes[1], //c
-			     input_tensor_cast->dims.dim_sizes[2], //h
-			     input_tensor_cast->dims.dim_sizes[3], //w
-			     filter_tensor_cast->dims.dim_sizes[0], //c_out
-			     filter_tensor_cast->dims.dim_sizes[1], //c_in
-			     filter_tensor_cast->dims.dim_sizes[2], //k_h
-			     filter_tensor_cast->dims.dim_sizes[3], //k_w
-			     conv_stride_h, //s_h
-			     conv_stride_w, //s_w
-			     param, //voltage_swing
-			     filter_tensor_cast->dims.dim_sizes[2] *
-                             filter_tensor_cast->dims.dim_sizes[3] /*patch_factor: k_h*k_w*/);
-	  RC->addToCurrentIterationComputeTime("ConvLayer_PROMISE", pinfo.first);
-	  RC->addToCurrentIterationComputeEnergy("ConvLayer_PROMISE", pinfo.second);
-	  void* t_out;
-	  // FIXIT: The pool_size in second param should be pool_stride
-	  t_out = PROMISE_Conv(input, QRanges[0], QRanges[1],
-			       filter, QRanges[2], QRanges[3],
-			       bias, QRanges[4], QRanges[5],
-			       conv_pad_h, conv_pad_w,
-			       conv_stride_h, conv_stride_w,
-			       pool_id, pool_size, pool_size,
-			       activation_id,
-			       QRanges[6], QRanges[7], param);
-
-	  return t_out;
-	} else {
-	  CUSTOM_ASSERT(false && "Unknown approximation type");
-	  ERROR("Unknown approximation type");
-	  abort();
-	}
-	// TODO additional approx methods implemented here
-
-      } else if (approxTuples.size() == 2) {
-	ERROR("Currently unsupported case");
-	abort();
-      } else {
-	ERROR("Unsupported case");
-	abort();
-      }
-    }
-    else
-      if (NodeConf->isGPUNodeConfiguration()) {
+    if (NodeConf->isGPUNodeConfiguration()) {
 	DEBUG("GPU Configuration for ConvLayer\n");
 	// Mapped to GPU - get a GPU node configuration
 	GPUNodeConfiguration *GPUConf = (GPUNodeConfiguration *)NodeConf;
@@ -271,69 +210,7 @@ extern "C"{
     INFO ("*** Conv Layer \n");
     
     NodeConfiguration *NodeConf = RC->getNodeConfiguration(hpvm_node_id);
-
-    if (NodeConf->isPROMISENodeConfiguration()) {
-      DEBUG("PROMISE Configuration for ConvLayer\n");
-      // Mapped to PROMISE - get a PROMISE node configuration
-      PROMISENodeConfiguration *PROMISEConf = (PROMISENodeConfiguration *)NodeConf;
-      std::vector<float> &QRanges = RC->getQuantizationRanges(hpvm_node_id);
-
-      std::vector<std::pair<PROMISENodeConfiguration::APPROX, int> > &approxTuples =
-	PROMISEConf->getApproxChoices();
-
-      if (approxTuples.size() == 1) {
-	enum PROMISENodeConfiguration::APPROX approx = approxTuples[0].first;
-	int param = approxTuples[0].second;
-	if (approx == PROMISENodeConfiguration::APPROX::SWING_LEVEL) {
-	  DEBUG("Approximation choice for ConvLayer: swing level %d\n", param);
-
-	  struct Tensor* input_tensor_cast = (struct Tensor*) input;
-	  struct Tensor* filter_tensor_cast = (struct Tensor*) filter;
-	  std::pair<double, double> pinfo =
-	    RC->conv_profile(input_tensor_cast->dims.dim_sizes[0], //n
-			     input_tensor_cast->dims.dim_sizes[1], //c
-			     input_tensor_cast->dims.dim_sizes[2], //h
-			     input_tensor_cast->dims.dim_sizes[3], //w
-			     filter_tensor_cast->dims.dim_sizes[0], //c_out
-			     filter_tensor_cast->dims.dim_sizes[1], //c_in
-			     filter_tensor_cast->dims.dim_sizes[2], //k_h
-			     filter_tensor_cast->dims.dim_sizes[3], //k_w
-			     conv_stride_h, //s_h
-			     conv_stride_w, //s_w
-			     param, //voltage_swing
-			     filter_tensor_cast->dims.dim_sizes[2] *
-                             filter_tensor_cast->dims.dim_sizes[3] /*patch_factor: k_h*k_w*/);
-	  RC->addToCurrentIterationComputeTime("ConvLayer_PROMISE", pinfo.first);
-	  RC->addToCurrentIterationComputeEnergy("ConvLayer_PROMISE", pinfo.second);
-	  void* t_out;
-	  // FIXIT: The pool_size in second param should be pool_stride
-	  t_out = PROMISE_Conv(input, QRanges[0], QRanges[1],
-			       filter, QRanges[2], QRanges[3],
-			       bias, QRanges[4], QRanges[5],
-			       conv_pad_h, conv_pad_w,
-			       conv_stride_h, conv_stride_w,
-			       pool_id, pool_size_v, pool_size_h,
-			       activation_id,
-			       QRanges[6], QRanges[7], param);
-
-	  return t_out;
-	} else {
-	  CUSTOM_ASSERT(false && "Unknown approximation type");
-	  ERROR("Unknown approximation type");
-	  abort();
-	}
-	// TODO additional approx methods implemented here
-
-      } else if (approxTuples.size() == 2) {
-	ERROR("Currently unsupported case");
-	abort();
-      } else {
-	ERROR("Unsupported case");
-	abort();
-      }
-    }
-    else
-      if (NodeConf->isGPUNodeConfiguration()) {
+	if (NodeConf->isGPUNodeConfiguration()) {
 	DEBUG("GPU Configuration for ConvLayer\n");
 	// Mapped to GPU - get a GPU node configuration
 	GPUNodeConfiguration *GPUConf = (GPUNodeConfiguration *)NodeConf;
@@ -495,64 +372,7 @@ extern "C"{
     INFO ("*** Dense Layer \n");
     
     NodeConfiguration *NodeConf = RC->getNodeConfiguration(hpvm_node_id);
-
-    if (NodeConf->isPROMISENodeConfiguration()) {
-      DEBUG("PROMISE Configuration for FCLayer\n");
-      // Mapped to PROMISE - get a PROMISE node configuration
-      PROMISENodeConfiguration *PROMISEConf = (PROMISENodeConfiguration *)NodeConf;
-      std::vector<float> &QRanges = RC->getQuantizationRanges(hpvm_node_id);
-
-      std::vector<std::pair<PROMISENodeConfiguration::APPROX, int> > &approxTuples =
-	PROMISEConf->getApproxChoices();
-
-      if (approxTuples.size() == 1) {
-	enum PROMISENodeConfiguration::APPROX approx = approxTuples[0].first;
-	int param = approxTuples[0].second;
-	if (approx == PROMISENodeConfiguration::APPROX::SWING_LEVEL) {
-	  DEBUG("Approximation choice for FCLayer: swing level %d\n", param);
-
-	  struct Tensor* input_tensor_cast = (struct Tensor*) input;
-	  struct Tensor* weights_tensor_cast = (struct Tensor*) weights;
-	  CUSTOM_ASSERT((input_tensor_cast->dims.dim_sizes[1] *
-			 input_tensor_cast->dims.dim_sizes[2] *
-			 input_tensor_cast->dims.dim_sizes[3] ==
-                         weights_tensor_cast->dims.dim_sizes[2]) &&
-			"Dimensions for matrix multiplication do not match.");
-	  std::pair<double, double> pinfo =
-	    RC->fc_profile(input_tensor_cast->dims.dim_sizes[0], //num_rows_a,
-			   input_tensor_cast->dims.dim_sizes[1] *
-                           input_tensor_cast->dims.dim_sizes[2] *
-                           input_tensor_cast->dims.dim_sizes[3], //num_cols_a,
-			   weights_tensor_cast->dims.dim_sizes[2], //num_rows_b,
-			   weights_tensor_cast->dims.dim_sizes[3], //num_cols_b,
-			   param, //voltage_swing,
-			   1 /*patch_factor*/);
-	  RC->addToCurrentIterationComputeTime("FCLayer_PROMISE", pinfo.first);
-	  RC->addToCurrentIterationComputeEnergy("FCLayer_PROMISE", pinfo.second);
-	  void* t_out;
-	  t_out = PROMISE_FC(input, QRanges[0], QRanges[1],
-			     weights, QRanges[2], QRanges[3],
-			     bias, QRanges[4], QRanges[5],
-			     activation_id,
-			     QRanges[6], QRanges[7], param);
-	  return t_out;
-	} else {
-	  CUSTOM_ASSERT(false && "Unknown approximation type");
-	  ERROR("Unknown approximation type");
-	  abort();
-	}
-	// TODO additional approx methods implemented here
-
-      } else if (approxTuples.size() == 2) {
-	ERROR("Currently unsupported case");
-	abort();
-      } else {
-	ERROR("Unsupported case");
-	abort();
-      }
-    }
-    else
-      if (NodeConf->isGPUNodeConfiguration()) {
+	if (NodeConf->isGPUNodeConfiguration()) {
 	DEBUG("GPU Configuration for FCLayer\n");
 	// Mapped to GPU - get a GPU node configuration
 	GPUNodeConfiguration *GPUConf = (GPUNodeConfiguration *)NodeConf;

hpvm/test/dnn_benchmarks/hpvm-c/CMakeLists.txt

@@ -43,7 +43,6 @@ foreach(dir ${entries})
   # Generate "tensor"-targeted code
   approxhpvm_py_codegen(
     ${dirname} ${dir}/${dirname}.cpp tensor 
-    --quant-file ${dir}/data/quant_ranges_rt.txt
     --conf-file ${dir}/data/tuner_confs.txt
   )
   # Run tensor binary

hpvm/tools/py-approxhpvm/main.py.in

@@ -34,13 +34,12 @@ def compile_hpvm_c(
     codegen_target: str = "tensor",
     include: List[PathLike] = None,
     working_dir: PathLike = None,
-    quant_file: PathLike = None,
     conf_file: PathLike = None,
 ):
     from subprocess import check_output
 
     codegen_functions = {
-        "tensor": lambda i, o: opt_codegen_tensor(i, o, quant_file, conf_file),
+        "tensor": lambda i, o: opt_codegen_tensor(i, o, conf_file),
         "cudnn": opt_codegen_cudnn
     }
     codegen_f = codegen_functions[codegen_target]
@@ -98,7 +97,7 @@ def opt_codegen_cudnn(src_file: PathLike, target_file: PathLike) -> List[str]:
 
 
 def opt_codegen_tensor(
-    src_file: PathLike, target_file: PathLike, quant_file: PathLike, conf_file: PathLike
+    src_file: PathLike, target_file: PathLike, conf_file: PathLike
 ):
     passes = [
         "LLVMBuildDFG", "LLVMInPlaceDFGAnalysis",
@@ -108,7 +107,6 @@ def opt_codegen_tensor(
     flags = [
         "buildDFG", "inplace", "hpvm-fuse",
         "dfg2llvm-wrapperapi",
-        f"quantization-levels-filename={quant_file}",
         f"configuration-inputs-filename={conf_file}",
         "dfg2llvm-cpu", "clearDFG",
     ]
@@ -159,10 +157,6 @@ def parse_args():
     parser.add_argument(
         "-d", "--working-dir", type=Path, help="Directory to generate temp files in"
     )
-    parser.add_argument(
-        "--quant-file", type=Path,
-        help="File to quantization levels of layers; required for 'tensor' target"
-    )
     parser.add_argument(
         "--conf-file", type=Path,
         help="File to approximation configurations; required for 'tensor' target"
@@ -174,10 +168,8 @@ def parse_args():
 
     args = parser.parse_args()
     if args.codegen_target == "tensor":
-        if args.quant_file is None:
-            parser.error('Codegen target "tensor" requires -quant-file argument')
         if args.conf_file is None:
-            parser.error('Codegen target "tensor" requires -conf-file argument')
+            parser.error('Codegen target "tensor" requires --conf-file argument')
     return args