Tc check

tc/aten/aten_autotuner.h

@@ -35,7 +35,7 @@ namespace aten {
  *    auto best = tuner.tune("tc_function_name", inputs, baseOption, cacheFn)
  *
  * The best options may then be used to compile an executor and run.
- *    CHECK_GT(best.size(), 0);
+ *    TC_CHECK_GT(best.size(), 0);
  *    auto pExecutor = compile(tc, "tc_function_name", inputs, best[0]);
  *    auto outputs = prepareOutputs(tc, "tc_function_name", inputs);
  *    // memoize the executor and outputs if needed

tc/aten/aten_compiler-inl.h

@@ -20,6 +20,7 @@
 #include <vector>
 
 #include "tc/aten/aten.h"
+#include "tc/core/check.h"
 #include "tc/core/compiler.h"
 #include "tc/core/tc_executor.h"
 #include "tc/core/tensor.h"
@@ -71,7 +72,7 @@ void uncheckedRun(
     const Executor& executor,
     const std::vector<at::Tensor>& inputs,
     std::vector<at::Tensor>& outputs) {
-  CHECK_GE(outputs.size(), 1u);
+  TC_CHECK_GE(outputs.size(), 1u);
   std::vector<const void*> rawInputs(inputs.size(), nullptr);
   std::vector<void*> rawOutputs(outputs.size(), nullptr);
   for (size_t i = 0; i < inputs.size(); ++i) {

tc/aten/aten_compiler.cc

@@ -20,6 +20,7 @@
 #include <vector>
 
 #include "tc/aten/aten.h"
+#include "tc/core/check.h"
 #include "tc/core/compiler.h"
 #include "tc/core/tc_executor.h"
 #include "tc/core/tensor.h"
@@ -32,7 +33,7 @@ std::vector<tc::DLTensorUPtr> inferOutputTensorInfo(
     const std::vector<at::Tensor>& inputs) {
   auto parsedTcs = tc::detail::parse(tc);
   if (parsedTcs.count(entryPoint) != 1u) {
-    CHECK_GE(parsedTcs.size(), 1u)
+    TC_CHECK_GE(parsedTcs.size(), 1u)
         << "No TC was parsed, should have thrown earlier";
     throw lang::ErrorReport(parsedTcs.begin()->second)
         << "\nattempting to access undefined entryPoint: " << entryPoint;
@@ -51,7 +52,7 @@ std::vector<at::Tensor> prepareOutputs(
   if (outTensorInfo.size() == 0) {
     return outputs;
   }
-  CHECK_GE(inputs.size(), 1u)
+  TC_CHECK_GE(inputs.size(), 1u)
       << "NYI: Need >= 1 input tensors to determine "
       << "backend and prepare ATen outputs. Add an overload with just an ATen "
       << "backend";

tc/autotuner/autotuner-inl.h

@@ -21,6 +21,7 @@
 #include <glog/stl_logging.h>
 
 #include "tc/autotuner/utils.h"
+#include "tc/core/check.h"
 #include "tc/core/compiler.h"
 #include "tc/core/flags.h"
 #include "tc/core/scope_guard.h"
@@ -127,9 +128,9 @@ void TuningHarness<Backend>::doEvaluate(
     size_t populationSize,
     Printer& printer) {
   typename Backend::WithDevice wd(device);
-  CHECK_EQ(inputs_.count(device), 1u);
+  TC_CHECK_EQ(inputs_.count(device), 1u);
   auto& inputs = inputs_.at(device);
-  CHECK_EQ(outputs_.count(device), 1u);
+  TC_CHECK_EQ(outputs_.count(device), 1u);
   auto& outputs = outputs_.at(device);
 
   while (true) {
@@ -158,7 +159,7 @@ void TuningHarness<Backend>::doEvaluate(
     if (!pExecutor.get()) {
       // If I popped an empty executor then compilation didn't go as
       // planned, skip it.
-      CHECK(pConf->invalid);
+      TC_CHECK(pConf->invalid);
       continue;
     }
 
@@ -252,8 +253,8 @@ void TuningHarness<Backend>::runOneIteration(
     size_t iteration) {
   // Define tensors per device once globally
   auto devices = detail::parseDevices<Backend>(FLAGS_tuner_devices);
-  CHECK(executors_.empty());
-  CHECK(configurations_.empty());
+  TC_CHECK(executors_.empty());
+  TC_CHECK(configurations_.empty());
 
   {
     // Initialize for this round
@@ -384,7 +385,7 @@ std::vector<size_t> inputDivisorsAndPowers2(
 }
 
 size_t largestDim(const std::vector<const DLConstTensor*>& inputs) {
-  CHECK_GE(inputs.size(), 1u);
+  TC_CHECK_GE(inputs.size(), 1u);
   auto maxElement = std::max_element(
       inputs.begin(),
       inputs.end(),
@@ -398,7 +399,7 @@ size_t largestDim(const std::vector<const DLConstTensor*>& inputs) {
 void setupTuningParameters(
     const std::vector<const DLConstTensor*>& inputs,
     TuningConfiguration& configuration) {
-  CHECK_GE(inputs.size(), 1u);
+  TC_CHECK_GE(inputs.size(), 1u);
   auto range = inputDivisorsAndPowers2(inputs);
   // 0 is a valid tiling annotation and signals no tiling of that dimension
   // 0 is not a valid block / grid annotation
@@ -428,12 +429,12 @@ Autotuner<Backend, SearchStrategy>::tune(
     const std::string& cacheFileName,
     const TuningParameterFixer& fixedParams) {
   std::map<std::string, lang::TreeRef> tcEntryPointMap(tc::detail::parse(tc));
-  CHECK_EQ(tcEntryPointMap.count(tcEntryPoint), 1u)
+  TC_CHECK_EQ(tcEntryPointMap.count(tcEntryPoint), 1u)
       << "Error looking up " << tcEntryPoint;
 
   // Initialize a model configuration
   TuningConfiguration modelConfiguration;
-  CHECK_GE(inputs.size(), 1u);
+  TC_CHECK_GE(inputs.size(), 1u);
   setupTuningParameters(inputs.begin()->second, modelConfiguration);
   modelConfiguration.fixParameters(fixedParams);
 

tc/autotuner/genetic_search.cc

@@ -19,6 +19,8 @@
 #include <random>
 #include <sstream>
 
+#include "tc/core/check.h"
+
 namespace tc {
 namespace autotune {
 
@@ -130,13 +132,13 @@ void dropInvalidConfigurations(GeneticSearch::Population& population) {
 
 } // namespace
 
-#define VALIDATE()                                   \
-  CHECK_LT(numberElites, maxPopulationSize);         \
-  CHECK(mutationRate >= 0 and mutationRate <= 100)   \
-      << "the mutation rate (" << mutationRate       \
-      << ") should be in the [0,100] interval";      \
-  CHECK(crossOverRate >= 0 and crossOverRate <= 100) \
-      << "the crossover (" << crossOverRate          \
+#define VALIDATE()                                      \
+  TC_CHECK_LT(numberElites, maxPopulationSize);         \
+  TC_CHECK(mutationRate >= 0 and mutationRate <= 100)   \
+      << "the mutation rate (" << mutationRate          \
+      << ") should be in the [0,100] interval";         \
+  TC_CHECK(crossOverRate >= 0 and crossOverRate <= 100) \
+      << "the crossover (" << crossOverRate             \
       << ") rate should be in the [0,100] interval";
 
 namespace {
@@ -170,7 +172,7 @@ GeneticSearch::GeneticSearch(
       rng{std::random_device{}()} {
   restoreRngState(rng);
   VALIDATE();
-  CHECK(not confs.empty()) << "empty set of predefined configurations";
+  TC_CHECK(not confs.empty()) << "empty set of predefined configurations";
 
   population.reserve(populationSize);
   size_t size = 0;
@@ -296,7 +298,7 @@ void GeneticSearch::updateParameters() {
           make_unique<CandidateConfiguration>(lastBestConf));
     }
     // Don't lose the first one which was the best from before
-    CHECK_LT(0u, population.size());
+    TC_CHECK_LT(0u, population.size());
     randomizePopulation(population.begin() + 1, population.end(), rng);
     return;
   }

tc/autotuner/options_cache-inl.h

@@ -26,6 +26,7 @@
 
 #include <llvm/ADT/Optional.h>
 
+#include "tc/core/check.h"
 #include "tc/core/tensor.h"
 #include "tc/core/utils/math.h"
 #include "tc/core/utils/time.h"
@@ -307,7 +308,7 @@ template <typename Backend>
 void OptionsCache<Backend>::fromProtobuf(
     const typename Backend::OptionsCacheProtoType& proto) {
   std::lock_guard<std::mutex> lock(mutex);
-  CHECK_EQ(proto.keys().size(), proto.values().size());
+  TC_CHECK_EQ(proto.keys().size(), proto.values().size());
   for (int i = 0; i < proto.keys().size(); ++i) {
     OptionsCacheKey key(OptionsCacheKey::fromProtobuf(proto.keys().Get(i)));
     OptionsCacheValue<Backend> value(

tc/autotuner/parameters.cc

@@ -23,6 +23,8 @@
 #include <sstream>
 #include <typeinfo>
 
+#include "tc/core/check.h"
+
 namespace tc {
 namespace autotune {
 
@@ -97,7 +99,7 @@ RangeParameter& RangeParameter::operator=(const RangeParameter& other) {
 }
 
 void BoolParameter::selectOption(size_t idx) {
-  CHECK_LE(idx, 1u);
+  TC_CHECK_LE(idx, 1u);
   selectValue(idx);
 }
 
@@ -106,7 +108,7 @@ void BoolParameter::selectValue(bool val) {
 }
 
 void RangeParameter::selectOption(size_t idx) {
-  CHECK_LE(idx, values_.size());
+  TC_CHECK_LE(idx, values_.size());
   selected_ = idx;
 }
 
@@ -124,8 +126,8 @@ void RangeParameter::selectFromValue(size_t value) {
 }
 
 void ParameterView::overwrite(const ParameterView& pv) {
-  CHECK_EQ(rangePtr == nullptr, pv.rangePtr == nullptr);
-  CHECK_EQ(boolPtr == nullptr, pv.boolPtr == nullptr);
+  TC_CHECK_EQ(rangePtr == nullptr, pv.rangePtr == nullptr);
+  TC_CHECK_EQ(boolPtr == nullptr, pv.boolPtr == nullptr);
   if (rangePtr) {
     *rangePtr = *pv.rangePtr;
   } else {
@@ -134,7 +136,7 @@ void ParameterView::overwrite(const ParameterView& pv) {
 }
 
 bool ParameterView::isForced() const {
-  CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
+  TC_CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
   if (rangePtr) {
     return rangePtr->fixedValue_.hasValue();
   } else {
@@ -143,7 +145,7 @@ bool ParameterView::isForced() const {
 }
 
 size_t ParameterView::numberOptions() const {
-  CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
+  TC_CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
   if (rangePtr) {
     return rangePtr->numberOptions();
   } else {
@@ -152,7 +154,7 @@ size_t ParameterView::numberOptions() const {
 }
 
 void ParameterView::selectOption(size_t idx) {
-  CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
+  TC_CHECK((rangePtr == nullptr) xor (boolPtr == nullptr));
   if (rangePtr) {
     return rangePtr->selectOption(idx);
   } else {
@@ -361,8 +363,8 @@ TuningConfiguration::TuningConfiguration()
         case 1:
           return b0v;
         default:
-          CHECK(false) << "Must have (1-3) block dims, got: "
-                       << conf.blockParams.numberDims.value();
+          TC_CHECK(false) << "Must have (1-3) block dims, got: "
+                          << conf.blockParams.numberDims.value();
       }
       return b0v;
     }();

tc/benchmarks/MLP_model.cc

@@ -26,6 +26,7 @@
 #include "tc/aten/aten.h"
 
 #include "tc/aten/aten_compiler.h"
+#include "tc/core/check.h"
 #include "tc/core/cuda/cuda_mapping_options.h"
 
 #include "../test/caffe2/cuda/test_harness.h"
@@ -200,8 +201,6 @@ class ProductionModel : public Benchmark {
 };
 
 void ProductionModel::run1LUT(const tc::CudaMappingOptions& options) {
-  CHECK_LT(0, E1);
-
   // This test uses an c2 OpTester because we need to run the C2 reference
   // implementation for TcLUTOp.
   auto ws_init_func = [=](Workspace& w) {
@@ -266,7 +265,7 @@ def _1LUT(float(E1, D) LUT1, int32(B, L1) I1) -> (O1) {
           inputs,
           options,
           check_fun);
-      CHECK_GE(bestOptions.size(), 1u);
+      TC_CHECK_GE(bestOptions.size(), 1u);
     }
     Check(tc, "_1LUT", options, inputs, check_fun);
   }
@@ -294,9 +293,8 @@ void ProductionModel::runATen1LUT() {
 }
 
 void ProductionModel::run2LUT(const tc::CudaMappingOptions& options) {
-  CHECK_LT(0, E1);
-  CHECK_LT(0, E2);
-
+  TC_CHECK_LT(0, E1);
+  TC_CHECK_LT(0, E2);
   auto ws_init_func = [=](Workspace& w) {
     AddDeterministicallyRandomInput<caffe2::CUDABackend, float>(
         w, {E1, D}, "LUT1");
@@ -385,7 +383,7 @@ def _2LUT(float(E1, D) LUT1, int32(B, L1) I1, float(E2, D) LUT2, int32(B, L2) I2
           inputs,
           options,
           check_fun);
-      CHECK_GE(bestOptions.size(), 1u);
+      TC_CHECK_GE(bestOptions.size(), 1u);
     }
     Check(tc, "_2LUT", bestOptions[0], inputs, check_fun);
   }
@@ -454,7 +452,7 @@ def _C3(float(B,WX) I, float(WY, WX) W) -> (C3) {
         inputs,
         options,
         check_fun);
-    CHECK_GE(bestOptions.size(), 1u);
+    TC_CHECK_GE(bestOptions.size(), 1u);
   }
   Check(tc, "_C3", bestOptions[0], inputs, check_fun);
 }
@@ -519,7 +517,7 @@ def mlp1(float(B,M) I, float(M, N) W1, float(N) B1) -> (O1) {
         inputs,
         options,
         check_fun);
-    CHECK_GE(bestOptions.size(), 1u);
+    TC_CHECK_GE(bestOptions.size(), 1u);
   }
   Check(tc, "mlp1", bestOptions[0], inputs, check_fun);
 }
@@ -602,7 +600,7 @@ def mlp3(float(B,N) I, float(O,N) W2, float(O) B2, float(P,O) W3, float(P) B3,
         inputs,
         options,
         check_fun);
-    CHECK_GE(bestOptions.size(), 1u);
+    TC_CHECK_GE(bestOptions.size(), 1u);
   }
   Check(tc, "mlp3", bestOptions[0], inputs, check_fun);
 }

tc/benchmarks/benchmark_fixture.h

@@ -29,6 +29,7 @@
 #include "tc/aten/aten_compiler.h"
 #include "tc/autotuner/genetic_search.h"
 #include "tc/autotuner/utils.h"
+#include "tc/core/check.h"
 #include "tc/core/cuda/cuda.h"
 #include "tc/core/cuda/cuda_mapping_options.h"
 #include "tc/core/cuda/cuda_rtc.h"
@@ -65,33 +66,33 @@ struct Benchmark : public ::testing::Test {
   void SetUp() {
     if (!FLAGS_disable_version_checks) {
       auto cudnnVersion = cudnnGetVersion();
-      CHECK_LE(6021, cudnnVersion)
+      TC_CHECK_LE(6021, cudnnVersion)
           << "[CUDNN][VERSION] Enforce version compatibility check";
 
       auto cudaRtVersion = cudnnGetCudartVersion();
-      CHECK_LE(8000, cudaRtVersion)
+      TC_CHECK_LE(8000, cudaRtVersion)
           << "[CUDART][VERSION] Enforce version compatibility check";
 
       int cublasVersion;
       cublasHandle_t handle;
       TC_CUDA_CUBLAS_ENFORCE(cublasCreate_v2(&handle));
       TC_CUDA_CUBLAS_ENFORCE(cublasGetVersion_v2(handle, &cublasVersion));
-      CHECK_LE(8000, cublasVersion)
+      TC_CHECK_LE(8000, cublasVersion)
           << "[CUBLAS][VERSION] Enforce version compatibility check";
       tc::ScopeGuard sg(
           [&handle]() { TC_CUDA_CUBLAS_ENFORCE(cublasDestroy_v2(handle)); });
 
       int cudaRuntimeVersion;
       TC_CUDA_RUNTIMEAPI_ENFORCE(cudaRuntimeGetVersion(&cudaRuntimeVersion));
-      CHECK_LE(8000, cudaRuntimeVersion)
+      TC_CHECK_LE(8000, cudaRuntimeVersion)
           << "[CUDA RUNTIME][VERSION] Enforce version compatibility check";
 
       int nvrtcVersionMajor;
       int nvrtcVersionMinor;
       TC_NVRTC_CHECK(nvrtcVersion(&nvrtcVersionMajor, &nvrtcVersionMinor));
-      CHECK_LE(8, nvrtcVersionMajor)
+      TC_CHECK_LE(8, nvrtcVersionMajor)
           << "[NVRTC][MAJOR][VERSION] Enforce version compatibility check";
-      CHECK_LE(0, nvrtcVersionMinor)
+      TC_CHECK_LE(0, nvrtcVersionMinor)
           << "[NVRTC][MINOR][VERSION] Enforce version compatibility check";
     }
   }
@@ -191,8 +192,8 @@ struct Benchmark : public ::testing::Test {
     auto bestOptions = [&]() {
       auto options = geneticAutotuneATen.tune(
           kernelName, inputs, baseMapping, cacheFilename, fixedParams);
-      CHECK_GE(options.size(), 1u) << "Benchmark mode: at least one "
-                                   << "options expected";
+      TC_CHECK_GE(options.size(), 1u) << "Benchmark mode: at least one "
+                                      << "options expected";
       return options[0];
     }();
     Check(tc, kernelName, bestOptions, inputs, check_fun);