diff --git a/tc/aten/aten_autotuner.h b/tc/aten/aten_autotuner.h index f11442169..e59f29618 100644 --- a/tc/aten/aten_autotuner.h +++ b/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 diff --git a/tc/aten/aten_compiler-inl.h b/tc/aten/aten_compiler-inl.h index 453f55598..8f64dec12 100644 --- a/tc/aten/aten_compiler-inl.h +++ b/tc/aten/aten_compiler-inl.h @@ -20,6 +20,7 @@ #include #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& inputs, std::vector& outputs) { - CHECK_GE(outputs.size(), 1u); + TC_CHECK_GE(outputs.size(), 1u); std::vector rawInputs(inputs.size(), nullptr); std::vector rawOutputs(outputs.size(), nullptr); for (size_t i = 0; i < inputs.size(); ++i) { diff --git a/tc/aten/aten_compiler.cc b/tc/aten/aten_compiler.cc index 5d6a6881e..b34fbd316 100644 --- a/tc/aten/aten_compiler.cc +++ b/tc/aten/aten_compiler.cc @@ -20,6 +20,7 @@ #include #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 inferOutputTensorInfo( const std::vector& 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 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"; diff --git a/tc/autotuner/autotuner-inl.h b/tc/autotuner/autotuner-inl.h index f9da58aed..2b5a035c6 100644 --- a/tc/autotuner/autotuner-inl.h +++ b/tc/autotuner/autotuner-inl.h @@ -21,6 +21,7 @@ #include #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::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::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::runOneIteration( size_t iteration) { // Define tensors per device once globally auto devices = detail::parseDevices(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 inputDivisorsAndPowers2( } size_t largestDim(const std::vector& 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& inputs) { void setupTuningParameters( const std::vector& 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::tune( const std::string& cacheFileName, const TuningParameterFixer& fixedParams) { std::map 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); diff --git a/tc/autotuner/genetic_search.cc b/tc/autotuner/genetic_search.cc index 416528e88..b004b4916 100644 --- a/tc/autotuner/genetic_search.cc +++ b/tc/autotuner/genetic_search.cc @@ -19,6 +19,8 @@ #include #include +#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(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; } diff --git a/tc/autotuner/options_cache-inl.h b/tc/autotuner/options_cache-inl.h index d4b03bf47..9f8d5836e 100644 --- a/tc/autotuner/options_cache-inl.h +++ b/tc/autotuner/options_cache-inl.h @@ -26,6 +26,7 @@ #include +#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 void OptionsCache::fromProtobuf( const typename Backend::OptionsCacheProtoType& proto) { std::lock_guard 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 value( diff --git a/tc/autotuner/parameters.cc b/tc/autotuner/parameters.cc index b0ec123b4..1d421e495 100644 --- a/tc/autotuner/parameters.cc +++ b/tc/autotuner/parameters.cc @@ -23,6 +23,8 @@ #include #include +#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; }(); diff --git a/tc/benchmarks/MLP_model.cc b/tc/benchmarks/MLP_model.cc index bcbc44607..f30867ebf 100644 --- a/tc/benchmarks/MLP_model.cc +++ b/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( 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); } diff --git a/tc/benchmarks/benchmark_fixture.h b/tc/benchmarks/benchmark_fixture.h index f208d8ef4..98da1b5e1 100644 --- a/tc/benchmarks/benchmark_fixture.h +++ b/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); diff --git a/tc/benchmarks/group_normalization.cc b/tc/benchmarks/group_normalization.cc index e11895f7a..cf2b683fe 100644 --- a/tc/benchmarks/group_normalization.cc +++ b/tc/benchmarks/group_normalization.cc @@ -106,7 +106,7 @@ std::vector GroupNormalization::runGroupNormalization( tc::TC_GroupNormalization_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } auto pExecutorMoments = tc::aten::compile( @@ -168,7 +168,7 @@ std::vector GroupNormalization::runGroupNormalizationSingleKernel( tc::TC_GroupNormalizationSingleKernel_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } return Check( tc::TC_GroupNormalization, diff --git a/tc/benchmarks/kronecker.cc b/tc/benchmarks/kronecker.cc index 53ec62127..65f58bada 100644 --- a/tc/benchmarks/kronecker.cc +++ b/tc/benchmarks/kronecker.cc @@ -112,7 +112,7 @@ std::vector Kronecker::runKronecker3_1( tc::TC_Kronecker3_1_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } return Check( tc::TC_Kronecker3_1, tc::TC_Kronecker3_1_NAME, bestOptions[0], inputs); @@ -142,7 +142,7 @@ std::vector Kronecker::runKronecker3_2( tc::TC_Kronecker3_2_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } return Check( tc::TC_Kronecker3_2, tc::TC_Kronecker3_2_NAME, bestOptions[0], inputs); @@ -173,7 +173,7 @@ std::vector Kronecker::runKronecker3_3( tc::TC_Kronecker3_3_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } return Check( tc::TC_Kronecker3_3, tc::TC_Kronecker3_3_NAME, bestOptions[0], inputs); @@ -297,7 +297,7 @@ void Kronecker::checkKronecker3Full( auto r3 = runKronecker3_3(options3, &W0, &r2[0]); auto checkFun = makeKronecker3CheckFunction(M, D0, D1, D2, N0, N1, N2); - CHECK(checkFun({W0, W1, W2, X}, r3)); + TC_CHECK(checkFun({W0, W1, W2, X}, r3)); } // Generic diff --git a/tc/benchmarks/kronecker.h b/tc/benchmarks/kronecker.h index a1ee1823e..d17d27e74 100644 --- a/tc/benchmarks/kronecker.h +++ b/tc/benchmarks/kronecker.h @@ -78,7 +78,7 @@ void cpu_kronecker_real_forward( const float* W_k = Ws[k]; if (k > 0) { // assert(Ysize[k-1] == M * N); - CHECK_EQ(M * N, Ysize[k - 1]) + TC_CHECK_EQ(M * N, Ysize[k - 1]) << "@k=" << k - 1 << ": " << M * N << " vs " << Ysize[k - 1]; } cpu_kronecker_real_forward_kernel(M, N, rowk, colk, stride, W_k, X_k, Y_k); @@ -88,7 +88,7 @@ void cpu_kronecker_real_forward( offset += rowk * colk; X_k = Y_k; // assert(Ysize[k] == M * N); - CHECK_EQ(M * N, Ysize[k]) + TC_CHECK_EQ(M * N, Ysize[k]) << "@k=" << k << ": " << M * N << " vs " << Ysize[k]; Y_k += Ysize[k]; } diff --git a/tc/benchmarks/moments.cc b/tc/benchmarks/moments.cc index df60d3b6d..32d352ebd 100644 --- a/tc/benchmarks/moments.cc +++ b/tc/benchmarks/moments.cc @@ -93,7 +93,7 @@ void Moments2_2D_1D::autotuneAndCheck( entryPoint, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } Check(tc::TC_Moments, entryPoint, bestOptions[0], inputs, checkFun); } diff --git a/tc/benchmarks/tmm.cc b/tc/benchmarks/tmm.cc index 5007e355a..8f8d8bc1c 100644 --- a/tc/benchmarks/tmm.cc +++ b/tc/benchmarks/tmm.cc @@ -95,7 +95,7 @@ def tmm(float(M,K) A, float(N,K) B) -> (C) { inputs, options, check_fun); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } Check(tc, "tmm", bestOptions[0], inputs, check_fun); } diff --git a/tc/benchmarks/wavenet.cc b/tc/benchmarks/wavenet.cc index de6a91614..c48c185c2 100644 --- a/tc/benchmarks/wavenet.cc +++ b/tc/benchmarks/wavenet.cc @@ -122,7 +122,7 @@ void WaveNet::runWaveNet1(const tc::CudaMappingOptions& options) { tc::TC_WAVENET1_NAME, inputs, options); - CHECK_GE(bestOptions.size(), 1u); + TC_CHECK_GE(bestOptions.size(), 1u); } Check(tc::TC_WAVENET, tc::TC_WAVENET1_NAME, bestOptions[0], inputs); } diff --git a/tc/c2/convolution_op.h b/tc/c2/convolution_op.h index 5a1b0b8bb..2a5f1cffb 100644 --- a/tc/c2/convolution_op.h +++ b/tc/c2/convolution_op.h @@ -19,6 +19,7 @@ #include #include "tc/c2/tc_op.h" +#include "tc/core/check.h" #include "tc/library/convolution.h" namespace caffe2 { @@ -58,7 +59,7 @@ class TcConvolutionOp : public TcOp { padR = OperatorBase::GetSingleArgument("pad_r", 0); } - CHECK(padT == 0 && padL == 0 && padB == 0 && padR == 0) + TC_CHECK(padT == 0 && padL == 0 && padB == 0 && padR == 0) << "NYI: padding larger than 0"; this->tc_ = tc::makeConvolution2DTc(strideH, strideW); diff --git a/tc/c2/dlpack_c2.h b/tc/c2/dlpack_c2.h index a3c3cc81a..5bc8f83eb 100644 --- a/tc/c2/dlpack_c2.h +++ b/tc/c2/dlpack_c2.h @@ -24,6 +24,7 @@ #include "tc/core/tensor.h" #include "caffe2/core/common.h" +#include "tc/core/check.h" namespace caffe2 { namespace dlpack { @@ -49,8 +50,8 @@ inline DLDataType getDLDataType(const TypeMeta& meta) { } else if (meta.Match()) { res.code = DLDataTypeCode::kDLInt; } else { - CHECK(false) << "NYI: getDLDataType(caffe2::Meta::Make<" << meta.name() - << ">))"; + TC_CHECK(false) << "NYI: getDLDataType(caffe2::Meta::Make<" << meta.name() + << ">))"; } res.bits = 32; res.lanes = 1; diff --git a/tc/c2/group_convolution_op.h b/tc/c2/group_convolution_op.h index 1fa5cf4ac..2b1359ff0 100644 --- a/tc/c2/group_convolution_op.h +++ b/tc/c2/group_convolution_op.h @@ -19,6 +19,7 @@ #include #include "tc/c2/tc_op.h" +#include "tc/core/check.h" #include "tc/library/group_convolution.h" namespace caffe2 { @@ -35,7 +36,7 @@ class TcGroupConvolutionOp : public TcOp { caffe2::Workspace* ws) : TcOp(operator_def, ws), group_(OperatorBase::GetSingleArgument("group", -1)) { - CHECK_EQ(-1, group_) + TC_CHECK_EQ(-1, group_) << "Caffe2 implements group convolution as a dilated convolution. " << "Someone (not us) needs to reshape."; @@ -65,7 +66,7 @@ class TcGroupConvolutionOp : public TcOp { padR = OperatorBase::GetSingleArgument("pad_r", 0); } - CHECK(padT == 0 && padL == 0 && padB == 0 && padR == 0) + TC_CHECK(padT == 0 && padL == 0 && padB == 0 && padR == 0) << "NYI: padding larger than 0"; this->tc_ = tc::makeGroupConvolution2DTc(strideH, strideW); diff --git a/tc/core/check.h b/tc/core/check.h new file mode 100644 index 000000000..e9f54ae19 --- /dev/null +++ b/tc/core/check.h @@ -0,0 +1,263 @@ +#pragma once + +#include +#include +#include +#include +#include + +#include "tc/core/utils/type_traits.h" + +/* + * Each TC_CHECK(_*) macro checks for a condition and throws an exception if + * the condition does not hold + * + * + *Additional information can be passed through operator<< and is + *included in the exception's error message, e.g.: + *TC_CHECK_EQ(x, 42) << "x is not the answer"; + * + * + * The message in the throw exception is: + * Check failed [filename:line_number] error(: info) + * + * filename: the name of the file where TC_CHECK(_*) was used + * lineno: the number of the line in which TC_CHECK(_*) was used + * error: this shows what failed, (e.g, "1 is not equal to 42") + * info: if operator<< was called then the information passed to it is info + * + * + * + * WARNING/CORNER CASE: + * + * Checker's destructor throws. This means that if another exception is thrown + * before a fully constructed Checker object is destroyed then the program will + * std::terminate. here is one unavoidable corner case: + * + * TC_CHECK(foo) << bar; + * + * If bar is a function/constructor call and throws then the program will + * std::terminate (because when Checker's destructor runs it will throw a + * second exception). + * + * + * Exception type: + * The default exception type is std::runtime_error, a different type can be + * specified by passing an extra argument, e.g.: + * TC_CHECK(i, whatever.size(), std::out_of_range); + * + * + * List of check macros: + * TC_CHECK(condition) //checks if condition is true + * TC_CHECK_EQ(x,y) //checks if x == y + * TC_CHECK_NE(x,y) //checks if x != y + * TC_CHECK_LT(x,y) //checks if x < y + * TC_CHECK_GT(x,y) //checks if x > y + * TC_CHECK_LE(x,y) //checks if x <= y + * TC_CHECK_GE(x,y) //checks if x >= y + */ + +// condition should either be a bool or convertible to bool +#define TC_CHECK_IMPL(condition, exception_type) \ + tc::detail::tc_check( \ + static_cast(condition), __FILE__, __LINE__) +// checks if x == y +#define TC_CHECK_EQ_IMPL(x, y, exception_type) \ + tc::detail::tc_check_eq(x, y, __FILE__, __LINE__) +// checks if x != y +#define TC_CHECK_NE_IMPL(x, y, exception_type) \ + tc::detail::tc_check_ne(x, y, __FILE__, __LINE__) +// checks if x < y +#define TC_CHECK_LT_IMPL(x, y, exception_type) \ + tc::detail::tc_check_lt(x, y, __FILE__, __LINE__) +// checks if x > y +#define TC_CHECK_GT_IMPL(x, y, exception_type) \ + tc::detail::tc_check_gt(x, y, __FILE__, __LINE__) +// checks if x <= y +#define TC_CHECK_LE_IMPL(x, y, exception_type) \ + tc::detail::tc_check_le(x, y, __FILE__, __LINE__) +// checks if x >= y +#define TC_CHECK_GE_IMPL(x, y, exception_type) \ + tc::detail::tc_check_ge(x, y, __FILE__, __LINE__) + +#define TC_CHECK_DEFAULT(condition) TC_CHECK_IMPL(condition, std::runtime_error) +#define TC_CHECK_EQ_DEFAULT(x, y, ...) \ + TC_CHECK_EQ_IMPL(x, y, std::runtime_error) +#define TC_CHECK_NE_DEFAULT(x, y, ...) \ + TC_CHECK_NE_IMPL(x, y, std::runtime_error) +#define TC_CHECK_LT_DEFAULT(x, y, ...) \ + TC_CHECK_LT_IMPL(x, y, std::runtime_error) +#define TC_CHECK_GT_DEFAULT(x, y, ...) \ + TC_CHECK_GT_IMPL(x, y, std::runtime_error) +#define TC_CHECK_LE_DEFAULT(x, y, ...) \ + TC_CHECK_LE_IMPL(x, y, std::runtime_error) +#define TC_CHECK_GE_DEFAULT(x, y, ...) \ + TC_CHECK_GE_IMPL(x, y, std::runtime_error) + +#define TC_GET_MACRO12(_1, _2, NAME, ...) NAME +#define TC_GET_MACRO23(_1, _2, _3, NAME, ...) NAME + +#define TC_CHECK(...) \ + TC_GET_MACRO12(__VA_ARGS__, TC_CHECK_IMPL, TC_CHECK_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_EQ(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_EQ_IMPL, TC_CHECK_EQ_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_NE(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_NE_IMPL, TC_CHECK_NE_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_LT(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_LT_IMPL, TC_CHECK_LT_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_GT(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_GT_IMPL, TC_CHECK_GT_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_LE(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_LE_IMPL, TC_CHECK_LE_DEFAULT) \ + (__VA_ARGS__) + +#define TC_CHECK_GE(...) \ + TC_GET_MACRO23(__VA_ARGS__, TC_CHECK_GE_IMPL, TC_CHECK_GE_DEFAULT) \ + (__VA_ARGS__) + +namespace tc { + +namespace detail { +template +class Checker { + public: + Checker(bool condition, std::string location, std::string baseErrorMsg) + : condition_(condition), + location_(location), + baseErrorMsg_(baseErrorMsg){}; + ~Checker() noexcept(false) { + if (condition_) { + return; + } + std::stringstream ss; + ss << "Check failed [" << location_ << ']'; + + if (not baseErrorMsg_.empty()) { + ss << ' ' << baseErrorMsg_; + } + + if (not additionalMsg_.empty()) { + ss << ": " << additionalMsg_; + } + throw ExceptionType(ss.str()); + } + + template + typename std::enable_if::value, Checker&>::type + operator<<(const T& msg) { + try { + std::stringstream ss; + ss << additionalMsg_ << msg; + additionalMsg_ = ss.str(); + } catch (...) { + // If the above throws and we don't catch the exception then the + // destructor will throw a second one and the program will terminate. + } + return *this; + } + + template + typename std::enable_if::value, Checker&>::type + operator<<(const C& msg) { + try { + std::stringstream ss; + ss << additionalMsg_; + for (const auto& x : msg) { + ss << x << ','; + } + additionalMsg_ = ss.str(); + if (msg.begin() != msg.end()) { + additionalMsg_.pop_back(); + } + } catch (...) { + // If the above throws and we don't catch the exception then the + // destructor will throw a second one and the program will terminate. + } + return *this; + } + + private: + bool condition_; + std::string location_; + std::string baseErrorMsg_; + std::string additionalMsg_; +}; // namespace detail + +inline std::string makeLocation(const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << filename << ':' << lineno; + return ss.str(); +} + +template +Checker +tc_check(bool condition, const char* filename, uint64_t lineno) { + return Checker(condition, makeLocation(filename, lineno), {}); +} + +template +Checker +tc_check_eq(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " not equal to " << y; + return Checker( + x == y, makeLocation(filename, lineno), ss.str()); +} + +template +Checker +tc_check_ne(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " equal to " << y; + return Checker( + x != y, makeLocation(filename, lineno), ss.str()); +} + +template +Checker +tc_check_lt(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " not less than " << y; + return Checker( + x < y, makeLocation(filename, lineno), ss.str()); +} + +template +Checker +tc_check_gt(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " not greater than " << y; + return Checker( + x > y, makeLocation(filename, lineno), ss.str()); +} + +template +Checker +tc_check_le(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " not less than or equal to " << y; + return Checker( + x <= y, makeLocation(filename, lineno), ss.str()); +} + +template +Checker +tc_check_ge(const X& x, const Y& y, const char* filename, uint64_t lineno) { + std::stringstream ss; + ss << x << " not greater than or equal to " << y; + return Checker( + x >= y, makeLocation(filename, lineno), ss.str()); +} + +} // namespace detail +} // namespace tc diff --git a/tc/core/compiler-inl.h b/tc/core/compiler-inl.h index e72ddfadb..40dd53e2f 100644 --- a/tc/core/compiler-inl.h +++ b/tc/core/compiler-inl.h @@ -18,6 +18,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/halide_utils.h" #include "tc/core/tensor.h" @@ -38,7 +39,7 @@ std::unique_ptr compile( /* TODO: in the future also pass outputs for stride and alignment info */ const typename Backend::MappingOptionsType& options) { auto parsedTcs = detail::parse(tc); - CHECK_EQ(parsedTcs.count(entryPoint), 1u) + TC_CHECK_EQ(parsedTcs.count(entryPoint), 1u) << "attempting to access undefined function " << entryPoint; return compile(parsedTcs[entryPoint], inputs, options); } diff --git a/tc/core/compiler.cc b/tc/core/compiler.cc index 83cc3c1da..ef57b7756 100644 --- a/tc/core/compiler.cc +++ b/tc/core/compiler.cc @@ -18,6 +18,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/exceptions.h" #include "tc/core/flags.h" #include "tc/core/halide_utils.h" @@ -30,7 +31,7 @@ std::vector inferOutputTensorInfo( const std::string& entryPoint, const std::vector inputs) { auto parsedTcs = detail::parse(tc); - CHECK_EQ(parsedTcs.count(entryPoint), 1u) + TC_CHECK_EQ(parsedTcs.count(entryPoint), 1u) << "attempting to access undefined function " << entryPoint; return tc::detail::inferOutputTensorInfo(parsedTcs[entryPoint], inputs); } diff --git a/tc/core/cpu/cpu_mapping_options.cc b/tc/core/cpu/cpu_mapping_options.cc index 663e273f2..fa00a935a 100644 --- a/tc/core/cpu/cpu_mapping_options.cc +++ b/tc/core/cpu/cpu_mapping_options.cc @@ -23,6 +23,7 @@ #include "tc/proto/mapping_options.pb.h" +#include "tc/core/check.h" #include "tc/core/cpu/cpu_mapping_options_cpp_printer.h" #include "tc/core/flags.h" #include "tc/core/utils/string.h" @@ -37,7 +38,7 @@ CpuMappingOptions::CpuMappingOptions() CpuMappingOptions::CpuMappingOptions(const std::string& str) : CpuMappingOptions() { bool parsed = ownedProto_.ParseFromString(str); - CHECK(parsed) << "could not parse protobuf string"; + TC_CHECK(parsed) << "could not parse protobuf string"; } CpuMappingOptions::CpuMappingOptions(const CpuMappingOptions& options) diff --git a/tc/core/cuda/cuda_mapping_options.cc b/tc/core/cuda/cuda_mapping_options.cc index 826939488..f4b4c743c 100644 --- a/tc/core/cuda/cuda_mapping_options.cc +++ b/tc/core/cuda/cuda_mapping_options.cc @@ -26,6 +26,7 @@ #include "tc/core/cuda/cuda_mapping_options_cpp_printer.h" #include "tc/core/flags.h" #include "tc/core/utils/string.h" + #include "tc/external/isl.h" namespace tc { @@ -50,8 +51,9 @@ std::string CudaDimView::toCommaSeparatedString() const { // CudaDimView & CudaDim // CudaDim::CudaDim(std::vector il) : ownedProto_(), view(ownedProto_) { - CHECK_GT(il.size(), 0u) << "list of values in CudaDimView must be non-empty"; - CHECK_LE(il.size(), 3u) << "at most 3 values allowed in CudaDimView"; + TC_CHECK_GT(il.size(), 0u) + << "list of values in CudaDimView must be non-empty"; + TC_CHECK_LE(il.size(), 3u) << "at most 3 values allowed in CudaDimView"; switch (il.size()) { case 3: @@ -62,7 +64,7 @@ CudaDim::CudaDim(std::vector il) : ownedProto_(), view(ownedProto_) { ownedProto_.set_x(*il.begin()); break; default: - CHECK(false) << "unreachable"; + TC_CHECK(false) << "unreachable"; } } @@ -81,7 +83,7 @@ CudaDim::CudaDim(uint64_t x, uint64_t y, uint64_t z) } size_t CudaDimView::size() const { - CHECK(!(!proto.has_y() && proto.has_z())) << "CudaDimView has z but not y"; + TC_CHECK(!(!proto.has_y() && proto.has_z())) << "CudaDimView has z but not y"; if (proto.has_z() && proto.has_y()) { return 3; @@ -92,7 +94,7 @@ size_t CudaDimView::size() const { } std::vector CudaDimView::extractVector() const { - CHECK(!(!proto.has_y() && proto.has_z())) << "CudaDimView has z but not y"; + TC_CHECK(!(!proto.has_y() && proto.has_z())) << "CudaDimView has z but not y"; std::vector result; result.push_back(proto.x()); @@ -110,13 +112,13 @@ std::array CudaDimView::extractDefaultedArray() const { CudaDimView::defaultDim, CudaDimView::defaultDim}; auto v = extractVector(); - CHECK_LE(v.size(), 3u); + TC_CHECK_LE(v.size(), 3u); std::copy(v.begin(), v.end(), arr.begin()); return arr; } ValueAccessor CudaDimView::operator[](size_t i) { - CHECK_LT(i, 3u) << "index overflow"; + TC_CHECK_LT(i, 3u) << "index overflow"; if (i == 0) { return ValueAccessor( [this](uint64_t u) { this->proto.set_x(u); }, @@ -139,7 +141,7 @@ ValueAccessor CudaDimView::operator[](size_t i) { } uint64_t CudaDimView::operator[](size_t i) const { - CHECK_LT(i, 3u) << "index overflow"; + TC_CHECK_LT(i, 3u) << "index overflow"; if (i == 0) { return proto.x(); } else if (i == 1) { @@ -189,7 +191,7 @@ CudaMappingOptions::CudaMappingOptions(const std::string& str) block = CudaDimView(*ownedProto_.mutable_block()); grid = CudaDimView(*ownedProto_.mutable_grid()); bool parsed = ownedProto_.ParseFromString(str); - CHECK(parsed) << "could not parse protobuf string"; + TC_CHECK(parsed) << "could not parse protobuf string"; } CudaMappingOptions& CudaMappingOptions::operator=( @@ -222,8 +224,8 @@ CudaMappingOptions::mapToThreads(uint64_t x, uint64_t y, uint64_t z) { CudaMappingOptions& CudaMappingOptions::mapToThreads( const std::vector& threads) { - CHECK_GT(threads.size(), 0u) << "expected at least one thread size"; - CHECK_LE(threads.size(), 3u) << "expected at most three thread sizes"; + TC_CHECK_GT(threads.size(), 0u) << "expected at least one thread size"; + TC_CHECK_LE(threads.size(), 3u) << "expected at most three thread sizes"; uint64_t x = threads[0]; uint64_t y = threads.size() > 1 ? threads[1] : CudaDimView::defaultDim; @@ -246,8 +248,8 @@ CudaMappingOptions::mapToBlocks(uint64_t x, uint64_t y, uint64_t z) { CudaMappingOptions& CudaMappingOptions::mapToBlocks( const std::vector& blocks) { - CHECK_GT(blocks.size(), 0u) << "expected at least one thread size"; - CHECK_LE(blocks.size(), 3u) << "expected at most three thread sizes"; + TC_CHECK_GT(blocks.size(), 0u) << "expected at least one thread size"; + TC_CHECK_LE(blocks.size(), 3u) << "expected at most three thread sizes"; uint64_t x = blocks[0]; uint64_t y = blocks.size() > 1 ? blocks[1] : CudaDimView::defaultDim; @@ -290,9 +292,9 @@ CudaMappingOptions& CudaMappingOptions::useReadOnlyCache(bool b) { CudaMappingOptions& CudaMappingOptions::mapToThreads( const std::string& commaSeparatedSizes) { auto sizes = parseCommaSeparatedIntegers(commaSeparatedSizes); - CHECK_GT(sizes.size(), 0u) + TC_CHECK_GT(sizes.size(), 0u) << "expected at least one block size in " << commaSeparatedSizes; - CHECK_LE(sizes.size(), 3u) + TC_CHECK_LE(sizes.size(), 3u) << "expected at most three block sizes in " << commaSeparatedSizes; sizes.resize(3, CudaDimView::defaultDim); return mapToThreads(sizes[0], sizes[1], sizes[2]); @@ -301,9 +303,9 @@ CudaMappingOptions& CudaMappingOptions::mapToThreads( CudaMappingOptions& CudaMappingOptions::mapToBlocks( const std::string& commaSeparatedSizes) { auto sizes = parseCommaSeparatedIntegers(commaSeparatedSizes); - CHECK_GT(sizes.size(), 0u) + TC_CHECK_GT(sizes.size(), 0u) << "expected at least one grid size in " << commaSeparatedSizes; - CHECK_LE(sizes.size(), 3u) + TC_CHECK_LE(sizes.size(), 3u) << "expected at most three grid sizes in " << commaSeparatedSizes; sizes.resize(3, CudaDimView::defaultDim); return mapToBlocks(sizes[0], sizes[1], sizes[2]); diff --git a/tc/core/cuda/cuda_rtc.cc b/tc/core/cuda/cuda_rtc.cc index b25c968a9..412e397dc 100644 --- a/tc/core/cuda/cuda_rtc.cc +++ b/tc/core/cuda/cuda_rtc.cc @@ -20,6 +20,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/cuda/cuda.h" #include "tc/core/cuda/cuda_rtc.h" #include "tc/core/flags.h" @@ -152,7 +153,8 @@ Duration CudaRTCFunction::Launch( constexpr size_t kNumMaxParameters = 100; std::array args_voidp{0}; - CHECK_GE(kNumMaxParameters, params.size() + outputs.size() + inputs.size()); + TC_CHECK_GE( + kNumMaxParameters, params.size() + outputs.size() + inputs.size()); int ind = 0; for (auto& p : params) { args_voidp[ind++] = &p; diff --git a/tc/core/cuda/cuda_tc_executor.cc b/tc/core/cuda/cuda_tc_executor.cc index 72a1350ad..cf27d5f29 100644 --- a/tc/core/cuda/cuda_tc_executor.cc +++ b/tc/core/cuda/cuda_tc_executor.cc @@ -15,6 +15,7 @@ */ #include "tc/core/cuda/cuda_tc_executor.h" +#include "tc/core/check.h" #include "tc/core/cuda/cuda_mapping_options_cpp_printer.h" #include "tc/core/halide_utils.h" #include "tc/core/polyhedral/cuda/mapped_scop.h" @@ -111,9 +112,9 @@ void CudaTcExecutor::uncheckedRun( const std::vector& inputs, const std::vector& outputs, typename CudaBackend::RuntimeInformation info) const { - CHECK(rtcFun_) << "No rtcFun_ attached, cannot launch"; - CHECK_NE(grid_.view[0], 0u) << "Grid dims are not set up"; - CHECK_NE(block_.view[0], 0u) << "Block dims are not set up"; + TC_CHECK(rtcFun_) << "No rtcFun_ attached, cannot launch"; + TC_CHECK_NE(grid_.view[0], 0u) << "Grid dims are not set up"; + TC_CHECK_NE(block_.view[0], 0u) << "Block dims are not set up"; rtcFun_->Launch( grid_.view.extractDefaultedArray(), block_.view.extractDefaultedArray(), @@ -128,10 +129,10 @@ ProfilingInfo CudaTcExecutor::profileUnchecked( const std::vector& inputs, const std::vector& outputs) const { auto start = std::chrono::system_clock::now(); - CHECK(rtcFun_) << "No rtcFun_ attached, cannot launch"; + TC_CHECK(rtcFun_) << "No rtcFun_ attached, cannot launch"; cudaStream_t stream = 0; - CHECK_NE(grid_.view[0], 0u) << "Grid dims are not set up"; - CHECK_NE(block_.view[0], 0u) << "Block dims are not set up"; + TC_CHECK_NE(grid_.view[0], 0u) << "Grid dims are not set up"; + TC_CHECK_NE(block_.view[0], 0u) << "Block dims are not set up"; Duration kernelRuntime(rtcFun_->Launch( grid_.view.extractDefaultedArray(), block_.view.extractDefaultedArray(), diff --git a/tc/core/halide2isl.cc b/tc/core/halide2isl.cc index c99b221e6..08b35453e 100644 --- a/tc/core/halide2isl.cc +++ b/tc/core/halide2isl.cc @@ -18,6 +18,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/polyhedral/schedule_isl_conversion.h" #include "tc/core/polyhedral/schedule_transforms.h" @@ -132,8 +133,8 @@ inline std::vector combineSingleAffs( isl::aff (isl::aff::*combine)(isl::aff) const) { auto left = makeIslAffBoundsFromExpr(space, op->a, false, false); auto right = makeIslAffBoundsFromExpr(space, op->b, false, false); - CHECK_LE(left.size(), 1u); - CHECK_LE(right.size(), 1u); + TC_CHECK_LE(left.size(), 1u); + TC_CHECK_LE(right.size(), 1u); if (left.size() == 0 || right.size() == 0) { return {}; @@ -164,7 +165,7 @@ std::vector makeIslAffBoundsFromExpr( const Expr& e, bool allowMin, bool allowMax) { - CHECK(!(allowMin && allowMax)); + TC_CHECK(!(allowMin && allowMax)); using Halide::Internal::Max; using Halide::Internal::Min; @@ -204,7 +205,7 @@ std::vector makeIslAffBoundsFromExpr( // We cannot span multiple constraints if a modulo operation is involved. // x > max(a,b) % C is not equivalent to (x > a % C && x > b % C). auto lhs = makeIslAffBoundsFromExpr(space, op->a, false, false); - CHECK_EQ(lhs.size(), 1u); + TC_CHECK_EQ(lhs.size(), 1u); if (const int64_t* b = as_const_int(op->b)) { return {lhs[0].mod(isl::val(space.get_ctx(), *b))}; } @@ -217,7 +218,7 @@ std::vector makeIslAffBoundsFromExpr( isl::aff makeIslAffFromExpr(isl::space space, const Expr& e) { auto list = makeIslAffBoundsFromExpr(space, e, false, false); - CHECK_LE(list.size(), 1u) + TC_CHECK_LE(list.size(), 1u) << "Halide expr " << e << " unrolled into more than 1 isl aff" << " but min/max operations were disabled"; @@ -373,7 +374,7 @@ isl::schedule makeScheduleTreeHelper( // Then we add our new loop bound constraints. auto lbs = halide2isl::makeIslAffBoundsFromExpr( set.get_space(), op->min, false, true); - CHECK_GT(lbs.size(), 0u) + TC_CHECK_GT(lbs.size(), 0u) << "could not obtain polyhedral lower bounds from " << op->min; for (auto lb : lbs) { set = set.intersect(loopVar.ge_set(lb)); @@ -382,7 +383,7 @@ isl::schedule makeScheduleTreeHelper( Expr max = simplify(op->min + op->extent - 1); auto ubs = halide2isl::makeIslAffBoundsFromExpr(set.get_space(), max, true, false); - CHECK_GT(ubs.size(), 0u) + TC_CHECK_GT(ubs.size(), 0u) << "could not obtain polyhedral upper bounds from " << max; for (auto ub : ubs) { set = set.intersect(ub.ge_set(loopVar)); diff --git a/tc/core/halide_utils.cc b/tc/core/halide_utils.cc index 647af7df4..3258ffe12 100644 --- a/tc/core/halide_utils.cc +++ b/tc/core/halide_utils.cc @@ -19,6 +19,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/tc2halide.h" #include "tc/core/tensor.h" @@ -72,7 +73,7 @@ std::unordered_map computeParamValueMap( } } else { // it was a constant const int64_t* c = as_const_int(extent); - CHECK(c != NULL); + TC_CHECK(c != NULL); if (*c != tensor->shape[d]) { throw lang::ErrorReport(dimExpTree) << "Constant dimension expected size " << *c << " but found " @@ -149,7 +150,7 @@ std::string halideCodegenC(const Stmt& stmt) { stream << "]"; } stream << " = "; - CHECK_EQ(1u, op->values.size()) + TC_CHECK_EQ(1u, op->values.size()) << "Cannot generate C for provide with != 1 values"; op->values[0].accept(this); stream << ";\n"; diff --git a/tc/core/mapping_options-inl.h b/tc/core/mapping_options-inl.h index 970ac2917..53dfe229f 100644 --- a/tc/core/mapping_options-inl.h +++ b/tc/core/mapping_options-inl.h @@ -15,6 +15,7 @@ */ #pragma once +#include "tc/core/check.h" #include "tc/core/utils/vararg.h" namespace tc { @@ -41,14 +42,14 @@ size_t TilingView::size() const { } ValueAccessor TilingView::operator[](size_t i) { - CHECK_LT(i, static_cast(proto.sizes_size())) << "index overflow"; + TC_CHECK_LT(i, static_cast(proto.sizes_size())) << "index overflow"; return ValueAccessor( [this, i](uint64_t u) { this->proto.set_sizes(i, u); }, [this, i]() { return this->proto.sizes(i); }); } uint64_t TilingView::operator[](size_t i) const { - CHECK_LT(i, static_cast(proto.sizes_size())) << "index overflow"; + TC_CHECK_LT(i, static_cast(proto.sizes_size())) << "index overflow"; return proto.sizes(i); } @@ -161,9 +162,9 @@ MappingOptionsView& MappingOptionsView::scheduleFusionStrategy( MappingOptionsView& MappingOptionsView::scheduleFusionStrategy( const std::string& str) { - FusionStrategy fs; + FusionStrategy fs(FusionStrategy::Max); bool couldParse = FusionStrategy_Parse(str, &fs); - CHECK(couldParse) << "unknown FusionStrategy " << str; + TC_CHECK(couldParse) << "unknown FusionStrategy " << str; return scheduleFusionStrategy(fs); } @@ -175,9 +176,9 @@ MappingOptionsView& MappingOptionsView::outerScheduleFusionStrategy( MappingOptionsView& MappingOptionsView::outerScheduleFusionStrategy( const std::string& str) { - FusionStrategy fs; + FusionStrategy fs(FusionStrategy::Max); bool couldParse = FusionStrategy_Parse(str, &fs); - CHECK(couldParse) << "unknown FusionStrategy " << str; + TC_CHECK(couldParse) << "unknown FusionStrategy " << str; return outerScheduleFusionStrategy(fs); } @@ -199,9 +200,9 @@ MappingOptionsView& MappingOptionsView::intraTileScheduleFusionStrategy( MappingOptionsView& MappingOptionsView::intraTileScheduleFusionStrategy( const std::string& str) { - FusionStrategy fs; + FusionStrategy fs(FusionStrategy::Max); bool couldParse = FusionStrategy_Parse(str, &fs); - CHECK(couldParse) << "unknown FusionStrategy " << str; + TC_CHECK(couldParse) << "unknown FusionStrategy " << str; return intraTileScheduleFusionStrategy(fs); } diff --git a/tc/core/polyhedral/codegen_llvm.cc b/tc/core/polyhedral/codegen_llvm.cc index f59f21e89..e418f8488 100644 --- a/tc/core/polyhedral/codegen_llvm.cc +++ b/tc/core/polyhedral/codegen_llvm.cc @@ -33,6 +33,7 @@ #include "Halide.h" +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/flags.h" #include "tc/core/halide2isl.h" @@ -71,7 +72,7 @@ namespace { thread_local llvm::LLVMContext llvmCtx; int64_t toSInt(isl::val v) { - CHECK(v.is_int()); + TC_CHECK(v.is_int()); static_assert(sizeof(long) <= 8, "long is assumed to fit into 64bits"); return v.get_num_si(); } @@ -82,7 +83,7 @@ llvm::Value* getLLVMConstantSignedInt64(int64_t v) { int64_t IslExprToSInt(isl::ast_expr e) { auto intExpr = e.as(); - CHECK(intExpr); + TC_CHECK(intExpr); return toSInt(intExpr.get_val()); } @@ -90,7 +91,7 @@ int64_t islIdToInt(isl::ast_expr_id e, isl::set context) { auto space = context.get_space(); isl::aff param(isl::aff::param_on_domain_space(space, e.get_id())); auto p = context.sample_point(); - CHECK(context.is_equal(p)); + TC_CHECK(context.is_equal(p)); return toSInt(param.eval(p)); } @@ -99,7 +100,7 @@ int64_t getTensorSize(isl::set context, const Halide::Expr& e) { // simplifying the expression. auto aff = halide2isl::makeIslAffFromExpr(context.get_space(), e); auto p = context.sample_point(); - CHECK(context.is_equal(p)); + TC_CHECK(context.is_equal(p)); return toSInt(aff.eval(p)); } @@ -111,7 +112,7 @@ std::vector getTensorSizesWithoutLeadingDim( sizes.reserve(dims); for (int d = 1; d < dims; ++d) { Halide::Expr extent = t.parameter().extent_constraint(d); - CHECK(extent.defined()) + TC_CHECK(extent.defined()) << "Undefined extent on input/output tensor. Forward bounds inference should have set these\n"; sizes.push_back(getTensorSize(context, extent)); } @@ -132,7 +133,7 @@ class IslAstExprInterpeter { } else if (auto opExpr = e.as()) { return interpretOp(opExpr); } else { - CHECK(false) << "NYI"; + TC_CHECK(false) << "NYI"; return 0; // avoid warning } }; @@ -145,7 +146,7 @@ class IslAstExprInterpeter { case 2: return interpretBinaryOp(e); default: - CHECK(false) << "NYI: " << e; + TC_CHECK(false) << "NYI: " << e; return 0; // avoid warning } } @@ -158,7 +159,7 @@ class IslAstExprInterpeter { } else if (e.as()) { return left - right; } else { - CHECK(false) << "NYI: " << e; + TC_CHECK(false) << "NYI: " << e; return 0; // avoid warning } } @@ -168,7 +169,7 @@ class IslAstExprInterpeter { if (e.as()) { return -val; } else { - CHECK(false) << "NYI"; + TC_CHECK(false) << "NYI"; return 0; // avoid warning } } @@ -255,10 +256,10 @@ class CodeGen_TC : public Halide::Internal::CodeGen_X86 { } else if (op->type.is_float()) { value = builder->CreateFPCast(value, ty); } else { - CHECK(false) << "Type inconsistency not handled. " - << "Variable " << op->name << " is " << op->type - << ", but its corresponding llvm::Value is " - << toString(value->getType()) << "."; + TC_CHECK(false) << "Type inconsistency not handled. " + << "Variable " << op->name << " is " << op->type + << ", but its corresponding llvm::Value is " + << toString(value->getType()) << "."; } } } @@ -454,8 +455,8 @@ class LLVMCodegen { llvm::Type* makePtrToArrayType( llvm::Type* baseTy, const std::vector& sizes) { - CHECK_GE(sizes.size(), 1u); - CHECK(baseTy); + TC_CHECK_GE(sizes.size(), 1u); + TC_CHECK(baseTy); llvm::Type* arrTy = llvm::ArrayType::get(baseTy, sizes.back()); for (auto s = sizes.rbegin() + 1; s != sizes.rend(); ++s) { arrTy = llvm::ArrayType::get(arrTy, *s); @@ -502,11 +503,11 @@ class LLVMCodegen { phi->addIncoming(getLLVMConstantSignedInt64(initVal), incoming); auto cond_expr = node.get_cond().as(); - CHECK(cond_expr.as() or cond_expr.as()) + TC_CHECK(cond_expr.as() or cond_expr.as()) << "I only know how to codegen lt and le"; auto condLHS = cond_expr.get_arg(0).as(); - CHECK(condLHS); - CHECK_EQ(condLHS.get_id(), iterator); + TC_CHECK(condLHS); + TC_CHECK_EQ(condLHS.get_id(), iterator); IslAstExprInterpeter i(scop_.context()); auto condRHSVal = i.interpret(cond_expr.get_arg(1)); @@ -518,7 +519,7 @@ class LLVMCodegen { } else if (cond_expr.as()) { return halide_cg.get_builder().CreateICmpSLE(phi, constant); } else { - CHECK(false) << "NYI"; + TC_CHECK(false) << "NYI"; return static_cast(nullptr); // avoid warning } }(); @@ -578,8 +579,8 @@ class LLVMCodegen { auto id = usrExp.get_arg(0).as().get_id(); auto provide = scop_.halide.statements.at(id); auto op = provide.as(); - CHECK(op) << "Expected a Provide node: " << provide << '\n'; - CHECK(op->values.size() == 1) + TC_CHECK(op) << "Expected a Provide node: " << provide << '\n'; + TC_CHECK(op->values.size() == 1) << "Multi-valued Provide: " << Halide::Internal::Stmt(provide) << "\n"; auto arrayName = op->name; const auto& subscripts = stmtSubscripts_.at(id); @@ -628,13 +629,13 @@ isl::ast_node collectIteratorMaps( const Scop& scop, StmtSubscriptExprMapType& stmtSubscripts) { auto user = node.as(); - CHECK(user); + TC_CHECK(user); auto expr = user.get_expr().as(); auto schedule = build.get_schedule(); auto scheduleMap = isl::map::from_union_map(schedule); auto stmtId = expr.get_arg(0).as().get_id(); - CHECK_EQ(0u, iteratorMaps.count(stmtId)) << "entry exists: " << stmtId; + TC_CHECK_EQ(0u, iteratorMaps.count(stmtId)) << "entry exists: " << stmtId; auto iteratorMap = isl::pw_multi_aff(scheduleMap.reverse()); auto iterators = scop.halide.iterators.at(stmtId); auto& stmtIteratorMap = iteratorMaps[stmtId]; @@ -650,7 +651,7 @@ isl::ast_node collectIteratorMaps( auto space = map.get_space().params(); auto aff = scop.makeIslAffFromStmtExpr(stmtId, space, e); auto pulled = isl::pw_aff(aff).pullback(map); - CHECK_EQ(pulled.n_piece(), 1); + TC_CHECK_EQ(pulled.n_piece(), 1); subscripts.push_back(build.expr_from(pulled)); } return node.set_annotation(stmtId); @@ -661,7 +662,6 @@ IslCodegenRes codegenISL(const Scop& scop) { StmtSubscriptExprMapType stmtSubscripts; auto collect = [&iteratorMaps, &scop, &stmtSubscripts]( isl::ast_node n, isl::ast_build b) -> isl::ast_node { - auto& uv = iteratorMaps; return collectIteratorMaps(n, b, uv, scop, stmtSubscripts); }; diff --git a/tc/core/polyhedral/cuda/codegen.cc b/tc/core/polyhedral/cuda/codegen.cc index ee1643984..c3425ca1c 100644 --- a/tc/core/polyhedral/cuda/codegen.cc +++ b/tc/core/polyhedral/cuda/codegen.cc @@ -20,6 +20,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/islpp_wrap.h" #include "tc/core/libraries.h" @@ -153,7 +154,7 @@ void emitKernelSignature( stringstream& ss, const std::string& specializedName, const Scop& scop) { - CHECK_NE(specializedName, "") << "name not provided"; + TC_CHECK_NE(specializedName, "") << "name not provided"; ss << "__global__ void " << specializedName << "("; emitArgs(ss, scop); ss << ") {" << endl; @@ -189,7 +190,7 @@ void emitTensorView( for (int i = 1; i < p.dimensions(); ++i) { // Skip the outermost dimension Halide::Expr extent = p.parameter().extent_constraint(i); extent = Halide::Internal::substitute(paramValues, extent); - CHECK(extent.defined()) + TC_CHECK(extent.defined()) << "Undefined extent on input/output tensor. Forward bounds inference should have set these\n"; ssViewType << "[" << extent << "]"; } @@ -249,8 +250,8 @@ void AstPrinter::emitIf(isl::ast_node_if node) { void emitReductionOpName(const Halide::Expr& e, const CodegenContext& context) { auto call = e.as(); - CHECK(call); - CHECK(call->is_intrinsic(tc2halide::kReductionUpdate)); + TC_CHECK(call); + TC_CHECK(call->is_intrinsic(tc2halide::kReductionUpdate)); context.ss << "__tc::ReductionOp::"; if (call->args[0].as()) { context.ss << "Sum"; @@ -261,7 +262,7 @@ void emitReductionOpName(const Halide::Expr& e, const CodegenContext& context) { } else if (call->args[0].as()) { context.ss << "Max"; } else { - CHECK(false) << "unsupported reduction type: " << e << "\n"; + TC_CHECK(false) << "unsupported reduction type: " << e << "\n"; } } @@ -271,7 +272,7 @@ void emitTreeSyncCall( isl::id id, isl::id reductionUpdateNodeId, const CodegenStatementContext& context) { - CHECK_EQ(1u, context.scop().treeSyncUpdateMap.count(id)); + TC_CHECK_EQ(1u, context.scop().treeSyncUpdateMap.count(id)); isl::id updateId = context.scop().treeSyncUpdateMap.at(id); // Halide reduction. @@ -312,14 +313,14 @@ void emitTreeSyncCall( } void emitUserStmt(isl::id stmtId, const CodegenStatementContext& context) { - CHECK(context.scop().halide.statements.count(stmtId)) + TC_CHECK(context.scop().halide.statements.count(stmtId)) << "No stmt with id " << stmtId << "\n"; auto provide = context.scop().halide.statements.at(stmtId); auto op = provide.as(); - CHECK(op) << "Expected a Provide node: " << provide << '\n'; + TC_CHECK(op) << "Expected a Provide node: " << provide << '\n'; detail::emitMappedTensorAccess(op->name, op, op->args, context); context.ss << " = "; - CHECK(op->values.size() == 1) + TC_CHECK(op->values.size() == 1) << "Multi-valued Provide: " << Halide::Internal::Stmt(provide) << "\n"; detail::emitHalideExpr(op->values[0], context); context.ss << ";" << endl; @@ -356,7 +357,7 @@ void emitReductionInit( .as(); context.ss << makeReductionTmpName(updateId, context.scop()) << " = "; auto call = provide->values[0].as(); - CHECK(call && call->is_intrinsic(tc2halide::kReductionUpdate)); + TC_CHECK(call && call->is_intrinsic(tc2halide::kReductionUpdate)); auto assoc = prove_associativity(provide->name, provide->args, call->args); if (!assoc.associative()) { std::stringstream ss; @@ -367,7 +368,7 @@ void emitReductionInit( throw codegen::NotAssociativeError(ss.str()); } auto statementContext = CodegenStatementContext(context, stmtId); - CHECK_EQ(assoc.pattern.identities.size(), 1u); + TC_CHECK_EQ(assoc.pattern.identities.size(), 1u); detail::emitHalideExpr(assoc.pattern.identities[0], statementContext); context.ss << ";" << endl; } @@ -428,7 +429,7 @@ void AstPrinter::emitStmt(isl::ast_node_user node) { auto stmtId = usrExp.get_arg(0).as().get_id(); auto nodeId = node.get_annotation(); auto statementContext = CodegenStatementContext(context_, nodeId); - CHECK_EQ(context_.nodeInfoMap.count(nodeId), 1u) + TC_CHECK_EQ(context_.nodeInfoMap.count(nodeId), 1u) << "no info for node " << nodeId; WS ws; @@ -454,7 +455,7 @@ void AstPrinter::emitStmt(isl::ast_node_user node) { emitCopyStmt(statementContext); } else { // regular statement auto mappedStmtId = statementContext.statementId(); - CHECK_EQ(stmtId, mappedStmtId) + TC_CHECK_EQ(stmtId, mappedStmtId) << "statement ids in expr (" << stmtId << ") and in iteratorMaps (" << mappedStmtId << ") do not match"; emitUserStmt(stmtId, statementContext); @@ -471,7 +472,7 @@ void AstPrinter::emitAst(isl::ast_node node) { emitAst(child); } } else if (node.as()) { - CHECK(false) << "mark"; + TC_CHECK(false) << "mark"; // emitAst(node.mark_get_node()); } else if (auto userNode = node.as()) { emitStmt(userNode); @@ -489,7 +490,7 @@ isl::pw_aff makeAffFromMappedExpr( const CodegenStatementContext& context) { // We only expect this to be called on encountering a free // variable. Compound expressions should be emitted as Halide. - CHECK(expr.as()); + TC_CHECK(expr.as()); auto aff = context.makeIslAffFromExpr(expr); auto pwaff = isl::pw_aff(aff).pullback(context.iteratorMap()); return pwaff; @@ -501,8 +502,8 @@ isl::space findDomainSpaceById(const CodegenStatementContext& context) { return d.get_space(); } } - CHECK(false) << "could not find domain for " << context.statementId() - << " in " << context.scop().domain(); + TC_CHECK(false) << "could not find domain for " << context.statementId() + << " in " << context.scop().domain(); return isl::space(); } @@ -510,7 +511,8 @@ isl::multi_aff makeMultiAffAccess( isl::id tensorId, const std::vector& subscripts, const CodegenStatementContext& context) { - CHECK_NE(subscripts.size(), 0u) << "cannot build subscript aff for a scalar"; + TC_CHECK_NE(subscripts.size(), 0u) + << "cannot build subscript aff for a scalar"; auto domainSpace = findDomainSpaceById(context); auto tensorSpace = domainSpace.params().named_set_from_params_id( @@ -596,7 +598,7 @@ void emitMappedTensorAccess( return; } - CHECK_EQ(context.scop().halide.accesses.count(node), 1u) + TC_CHECK_EQ(context.scop().halide.accesses.count(node), 1u) << "attempting to generate code for tensor " << name << " reference not present in Scop" << node; auto refId = context.scop().halide.accesses.at(node); @@ -604,7 +606,7 @@ void emitMappedTensorAccess( Scop::PromotionInfo promotionInfo; for (auto pi : context.activePromotions()) { if (pi.group->referenceIds().count(refId)) { - CHECK(!promotionInfo.groupId) + TC_CHECK(!promotionInfo.groupId) << "reference " << refId << " belongs to two groups: " << promotionInfo.groupId << " and " << pi.groupId; @@ -639,7 +641,7 @@ void emitMappedTensorAccess( isl::map::from_union_map(promotionInfo.outerSchedule.intersect_domain( context.domain())); // map :: D -> S - CHECK(schedule.is_single_valued()) + TC_CHECK(schedule.is_single_valued()) << "expected single-valued schedule, got " << schedule; // PMA :: A -> S auto astToSchedule = isl::pw_multi_aff(schedule).pullback(iteratorMap); @@ -740,8 +742,8 @@ string emitCudaKernel( const std::string& specializedName, const MappedScop& mscop) { // Expecting a schedule with domain root and context first child. - CHECK(mscop.schedule()->elemAs()); - CHECK( + TC_CHECK(mscop.schedule()->elemAs()); + TC_CHECK( mscop.schedule()->child({0})->elemAs()); const auto& scop = mscop.scop(); @@ -766,7 +768,7 @@ string emitCudaKernel( isl::ast_build build, NodeInfoMapType* nodeInfoMap) -> isl::ast_node { auto user = node.as(); - CHECK(user); + TC_CHECK(user); auto expr = user.get_expr().as(); auto stmtId = expr.get_arg(0).as().get_id(); auto schedule = build.get_schedule(); @@ -775,7 +777,7 @@ string emitCudaKernel( auto nodeId = isl::id( node.get_ctx(), std::string(kAstNodeIdPrefix) + std::to_string(nAstNodes()++)); - CHECK_EQ(0u, nodeInfoMap->count(nodeId)) << "entry exists: " << nodeId; + TC_CHECK_EQ(0u, nodeInfoMap->count(nodeId)) << "entry exists: " << nodeId; auto& nodeInfo = (*nodeInfoMap)[nodeId]; nodeInfo.iteratorMap = isl::pw_multi_aff(scheduleMap.reverse()); diff --git a/tc/core/polyhedral/cuda/cuda_mapping_types-inl.h b/tc/core/polyhedral/cuda/cuda_mapping_types-inl.h index 99aa16547..aeb90fea0 100644 --- a/tc/core/polyhedral/cuda/cuda_mapping_types-inl.h +++ b/tc/core/polyhedral/cuda/cuda_mapping_types-inl.h @@ -15,11 +15,13 @@ */ #pragma once +#include "tc/core/check.h" + namespace tc { namespace polyhedral { namespace mapping { ThreadId ThreadId::makeId(size_t dim) { - CHECK(dim < 3); + TC_CHECK_LT(dim, 3ul); if (dim == 0) { return ThreadId::x(); } @@ -50,7 +52,7 @@ ThreadId ThreadId::z() { } BlockId BlockId::makeId(size_t dim) { - CHECK(dim < 3); + TC_CHECK_LT(dim, 3ul); if (dim == 0) { return BlockId::x(); } diff --git a/tc/core/polyhedral/cuda/mapped_scop.cc b/tc/core/polyhedral/cuda/mapped_scop.cc index e0dc474ae..fdad777f6 100644 --- a/tc/core/polyhedral/cuda/mapped_scop.cc +++ b/tc/core/polyhedral/cuda/mapped_scop.cc @@ -23,6 +23,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/gpu.h" #include "tc/core/libraries.h" @@ -106,7 +107,7 @@ detail::ScheduleTree* MappedScop::map( isl::union_pw_aff_list list) { size_t nToMap = list.n(); const auto& extent = mappingSize(this).view; - CHECK_LE(nToMap, extent.size()) << "dimension overflow"; + TC_CHECK_LE(nToMap, extent.size()) << "dimension overflow"; auto root = scop_->scheduleRoot(); auto domain = activeDomainPoints(root, tree).universe(); @@ -116,7 +117,7 @@ detail::ScheduleTree* MappedScop::map( for (size_t i = 0; i < nToMap; ++i) { auto id = MappingTypeId::makeId(i); auto upa = list.get(i); - CHECK_NE(extent[i], 0u) << "NYI: mapping to 0"; + TC_CHECK_NE(extent[i], 0u) << "NYI: mapping to 0"; upa = upa.mod_val(isl::val(tree->ctx_, extent[i])); affList = affList.add(upa); idList.emplace_back(id); @@ -139,7 +140,7 @@ detail::ScheduleTree* MappedScop::mapBlocksForward( detail::ScheduleTree* band, size_t nToMap) { auto bandNode = band->elemAs(); - CHECK(bandNode) << "expected a band, got " << *band; + TC_CHECK(bandNode) << "expected a band, got " << *band; auto list = bandNode->mupa_.get_union_pw_aff_list(); list = list.drop(nToMap, list.n() - nToMap); @@ -154,7 +155,7 @@ void MappedScop::mapToBlocksAndScaleBand( using namespace tc::polyhedral::detail; auto bandNode = band->elemAs(); - CHECK(bandNode->permutable_) << "cannot map non-permutable band to blocks"; + TC_CHECK(bandNode->permutable_) << "cannot map non-permutable band to blocks"; auto nBlocksToMap = bandNode->nOuterCoincident(); // Can map at most 3 dimensions @@ -288,9 +289,9 @@ bool MappedScop::needReductionSeparation(const detail::ScheduleTree* st) { isl::multi_union_pw_aff MappedScop::reductionMapSchedule( const detail::ScheduleTree* st) { - CHECK(reductionBandUpdates_.count(st) == 1); + TC_CHECK(reductionBandUpdates_.count(st) == 1); auto reductionBand = st->elemAs(); - CHECK(reductionBand); + TC_CHECK(reductionBand); // Drop band members following the reduction dimension and preceding those // mapped to threads. @@ -298,7 +299,7 @@ isl::multi_union_pw_aff MappedScop::reductionMapSchedule( auto nMember = reductionBand->nMember(); auto reductionDim = reductionBand->nOuterCoincident(); auto nMappedThreads = std::min(numThreads.view.size(), reductionDim + 1); - CHECK_GE(nMember, reductionDim); + TC_CHECK_GE(nMember, reductionDim); reductionSchedule = reductionSchedule.drop_dims( isl::dim_type::set, reductionDim + 1, nMember - (reductionDim + 1)); reductionSchedule = reductionSchedule.drop_dims( @@ -359,10 +360,10 @@ detail::ScheduleTree* MappedScop::separateReduction(detail::ScheduleTree* st) { detail::ScheduleTree* MappedScop::mapThreadsBackward( detail::ScheduleTree* band) { auto bandNode = band->elemAs(); - CHECK(bandNode); + TC_CHECK(bandNode); auto nMember = bandNode->nMember(); auto nToMap = std::min(nMember, numThreads.view.size()); - CHECK_LE(nToMap, 3u) << "mapping to too many threads"; + TC_CHECK_LE(nToMap, 3u) << "mapping to too many threads"; auto ctx = band->ctx_; insertNodeBelow(band, detail::ScheduleTree::makeThreadSpecificMarker(ctx)); @@ -393,7 +394,7 @@ size_t MappedScop::mapToThreads(detail::ScheduleTree* band) { // this member has to be mapped as well. // In particular, it will get mapped to threadIdx.x if (isReduction) { - CHECK(reductionBandUpdates_.at(band).separated); + TC_CHECK(reductionBandUpdates_.at(band).separated); nCanMap++; } @@ -423,7 +424,7 @@ size_t MappedScop::mapToThreads(detail::ScheduleTree* band) { bandSplit(scop_->scheduleRoot(), band, nMappedThreads); } - CHECK_GT(nMappedThreads, 0u) << "not mapping to threads"; + TC_CHECK_GT(nMappedThreads, 0u) << "not mapping to threads"; if (isReduction) { band = splitOutReductionTileAndInsertSyncs(band); @@ -551,10 +552,10 @@ Scop::SyncLevel MappedScop::findBestSync( return Scop::SyncLevel::None; } - CHECK_LE(1u, scop_->scheduleRoot()->children().size()); + TC_CHECK_LE(1u, scop_->scheduleRoot()->children().size()); auto contextSt = scop_->scheduleRoot()->children()[0]; auto contextElem = contextSt->elemAs(); - CHECK(nullptr != contextElem); + TC_CHECK(nullptr != contextElem); dependences = dependences.intersect_params(contextElem->context_); if (dependences.is_subset(dependences.eq_at(domainToThread))) { @@ -716,7 +717,7 @@ std::vector> MappedScop::findBestSyncConfigInSeq( } void MappedScop::insertBestSyncInSeq(detail::ScheduleTree* seq) { - CHECK(seq->elemAs()); + TC_CHECK(seq->elemAs()); auto children = seq->children(); auto nChildren = children.size(); @@ -817,9 +818,9 @@ size_t MappedScop::mapInnermostBandsToThreads(detail::ScheduleTree* st) { // member, insert a synchronization after its last child. // The node must have children if some of them were mapped to threads, // double-check. Note that a band node has at most one child. - CHECK_EQ(st->numChildren(), 1u); + TC_CHECK_EQ(st->numChildren(), 1u); // The mapping should be always complete, double-check. - CHECK_EQ(n, numThreads.view.size()); + TC_CHECK_EQ(n, numThreads.view.size()); scop_->insertSyncAfter(st->child({0})); } } @@ -987,7 +988,7 @@ std::unique_ptr MappedScop::makeWithOuterBlockInnerThreadStrategy( scop = Scop::makeScheduled(*scop, generic.outerScheduleOptions); // 3. Tile - CHECK_LT(0u, generic.tiling.size()) + TC_CHECK_LT(0u, generic.tiling.size()) << "Must pass tile vector with >= 1 tile sizes"; auto outerBand = scop->tileOuterBand(generic.tiling); @@ -1010,7 +1011,7 @@ std::unique_ptr MappedScop::makeWithOuterBlockInnerThreadStrategy( // 6. Map to threads if (outerBand->numChildren() > 0) { - CHECK_EQ(1u, outerBand->numChildren()); + TC_CHECK_EQ(1u, outerBand->numChildren()); // 6.1. Optionally detect reductions while mapping to threads if (generic.proto.match_library_calls()) { diff --git a/tc/core/polyhedral/cuda/tighten_launch_bounds.cc b/tc/core/polyhedral/cuda/tighten_launch_bounds.cc index a5f3f4248..65ef48f47 100644 --- a/tc/core/polyhedral/cuda/tighten_launch_bounds.cc +++ b/tc/core/polyhedral/cuda/tighten_launch_bounds.cc @@ -16,6 +16,7 @@ #include "tc/core/polyhedral/cuda/tighten_launch_bounds.h" +#include "tc/core/check.h" #include "tc/core/polyhedral/cuda/mapping_types.h" #include "tc/core/polyhedral/exceptions.h" #include "tc/core/polyhedral/functional.h" @@ -56,11 +57,11 @@ std::pair rangeOfMappingParameter( if (max.is_nan() || max.is_infty()) { return std::make_pair(0, std::numeric_limits::max()); } - CHECK(max.is_int()) << max.to_str(); - CHECK(max.is_nonneg()) << max.to_str(); + TC_CHECK(max.is_int()) << max.to_str(); + TC_CHECK(max.is_nonneg()) << max.to_str(); auto min = active.min_val(a); - CHECK(min.is_int()) << max.to_str(); - CHECK(min.is_nonneg()) << max.to_str(); + TC_CHECK(min.is_int()) << max.to_str(); + TC_CHECK(min.is_nonneg()) << max.to_str(); return std::make_pair( static_cast(min.get_num_si()), diff --git a/tc/core/polyhedral/llvm_jit.cc b/tc/core/polyhedral/llvm_jit.cc index b3e495e56..98937b0d4 100644 --- a/tc/core/polyhedral/llvm_jit.cc +++ b/tc/core/polyhedral/llvm_jit.cc @@ -15,6 +15,7 @@ */ #include +#include "tc/core/check.h" #include "tc/core/polyhedral/llvm_jit.h" #include "llvm/ExecutionEngine/ExecutionEngine.h" @@ -27,6 +28,7 @@ #include "llvm/Support/DynamicLibrary.h" #include "llvm/Transforms/Utils/Cloning.h" +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/polyhedral/codegen_llvm.h" @@ -99,7 +101,7 @@ void Jit::addModule(std::shared_ptr M) { }); auto res = compileLayer_.addModule(M, std::move(Resolver)); - CHECK(res) << "Failed to jit compile."; + TC_CHECK(res) << "Failed to jit compile."; } #else Jit::Jit() @@ -142,7 +144,7 @@ void Jit::addModule(std::shared_ptr M) { M->setTargetTriple(TM_->getTargetTriple().str()); auto K = ES.allocateVModule(); llvm::Error res = compileLayer_.addModule(K, CloneModule(*M)); - CHECK(!res) << "Failed to jit compile."; + TC_CHECK(!res) << "Failed to jit compile."; } #endif @@ -168,7 +170,7 @@ JITSymbol Jit::findSymbol(const std::string Name) { JITTargetAddress Jit::getSymbolAddress(const std::string Name) { auto res = findSymbol(Name).getAddress(); - CHECK(res) << "Could not find jit-ed symbol"; + TC_CHECK(res) << "Could not find jit-ed symbol"; return *res; } diff --git a/tc/core/polyhedral/memory_promotion.cc b/tc/core/polyhedral/memory_promotion.cc index f26ea03e2..3c6eccd1f 100644 --- a/tc/core/polyhedral/memory_promotion.cc +++ b/tc/core/polyhedral/memory_promotion.cc @@ -20,6 +20,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/polyhedral/exceptions.h" #include "tc/core/polyhedral/schedule_tree.h" #include "tc/core/polyhedral/scop.h" @@ -45,7 +46,7 @@ isl::map removeRangeStrides( isl::map relation, isl::multi_val strides, isl::multi_aff offsets) { - CHECK_EQ(strides.size(), offsets.size()); + TC_CHECK_EQ(strides.size(), offsets.size()); auto space = relation.get_space(); auto stridesMA = isl::multi_aff::identity(space.range().map_from_set()); diff --git a/tc/core/polyhedral/reduction_matcher.cc b/tc/core/polyhedral/reduction_matcher.cc index ff86fee9e..99865ac6d 100644 --- a/tc/core/polyhedral/reduction_matcher.cc +++ b/tc/core/polyhedral/reduction_matcher.cc @@ -17,6 +17,7 @@ #include +#include "tc/core/check.h" #include "tc/core/polyhedral/schedule_tree.h" #include "tc/core/polyhedral/scop.h" #include "tc/external/isl.h" @@ -53,7 +54,7 @@ bool isReductionUpdateId( isl::id id, const Scop& scop, std::vector& reductionDims) { - CHECK_EQ(scop.halide.statements.count(id), 1u) + TC_CHECK_EQ(scop.halide.statements.count(id), 1u) << "id is not a statement in scop" << id; auto provideNode = scop.halide.statements.at(id); if (!isSupportedReduction(provideNode)) { diff --git a/tc/core/polyhedral/schedule_isl_conversion.cc b/tc/core/polyhedral/schedule_isl_conversion.cc index 0f4b2b43a..8fa4ecb74 100644 --- a/tc/core/polyhedral/schedule_isl_conversion.cc +++ b/tc/core/polyhedral/schedule_isl_conversion.cc @@ -21,6 +21,7 @@ #include "tc/external/isl.h" +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/polyhedral/schedule_transforms.h" #include "tc/external/isl.h" @@ -66,7 +67,7 @@ isl::schedule_node insertBranch( auto filters = isl::union_set_list(node.get_ctx(), st->numChildren()); for (size_t i = 0; i < pos.size(); ++i) { auto filter = st->child({pos[i]})->elemAsBase(); - CHECK(filter); + TC_CHECK(filter); filters = filters.add(filter->filter_); } if (st->elemAs()) { @@ -100,10 +101,10 @@ std::vector findCorePositions( const ScheduleTree* st, isl::union_set domain) { std::vector positions; - CHECK(st->elemAs()); + TC_CHECK(st->elemAs()); for (size_t i = 0; i < st->numChildren(); ++i) { auto filter = st->child({i})->elemAsBase(); - CHECK(filter); + TC_CHECK(filter); if (!filter->filter_.intersect(domain).is_empty()) { positions.emplace_back(i); } @@ -121,7 +122,7 @@ isl::schedule_node graftFromFilterSubtree( const ScheduleTree* st, isl::union_map extension) { auto filter = st->elemAsBase(); - CHECK(filter); + TC_CHECK(filter); auto filterExtension = extension.intersect_range(filter->filter_); auto extensionNode = isl::schedule_node::from_extension(filterExtension); return extendChild(extensionNode, st); @@ -145,7 +146,7 @@ isl::schedule_node insertExtension( auto domain = node.get_universe_domain(); auto child = st->child({0}); auto corePos = findCorePositions(child, domain); - CHECK(!corePos.empty()); + TC_CHECK(!corePos.empty()); node = insertBranch(node, child, corePos); auto extension = st->elemAs()->extension_; @@ -244,7 +245,7 @@ isl::schedule_node extendChild( */ isl::schedule toIslSchedule(const ScheduleTree* root) { auto domain = root->elemAs(); - CHECK(domain) << "Root node should be domain node" << *root; + TC_CHECK(domain) << "Root node should be domain node" << *root; auto node = isl::schedule_node::from_domain(domain->domain_); node = extendChild(node, root); return node.get_schedule(); diff --git a/tc/core/polyhedral/schedule_print.cc b/tc/core/polyhedral/schedule_print.cc index d12b43541..0f5ea94cd 100644 --- a/tc/core/polyhedral/schedule_print.cc +++ b/tc/core/polyhedral/schedule_print.cc @@ -19,6 +19,7 @@ #include #include "tc/external/isl.h" +#include "tc/core/check.h" #include "tc/core/polyhedral/schedule_tree.h" #include "tc/core/polyhedral/schedule_tree_elem.h" #include "tc/external/isl.h" @@ -247,7 +248,7 @@ std::ostream& operator<<( } std::ostream& operator<<(std::ostream& os, const ScheduleTree& st) { - CHECK(st.elem_.get()); + TC_CHECK(st.elem_.get()); os << *st.elem_ << "\n"; os << st.children_; diff --git a/tc/core/polyhedral/schedule_transforms-inl.h b/tc/core/polyhedral/schedule_transforms-inl.h index 22309f59d..6b1f8893c 100644 --- a/tc/core/polyhedral/schedule_transforms-inl.h +++ b/tc/core/polyhedral/schedule_transforms-inl.h @@ -15,13 +15,15 @@ */ #pragma once +#include "tc/core/check.h" + namespace tc { namespace polyhedral { inline detail::ScheduleTree* insertNodeAbove( detail::ScheduleTree* root, detail::ScheduleTree* tree, ScheduleTreeUPtr&& node) { - CHECK_EQ(node->numChildren(), 0u); + TC_CHECK_EQ(node->numChildren(), 0u); auto parent = tree->ancestor(root, 1); auto childPos = tree->positionInParent(parent); node->appendChild(parent->detachChild(childPos)); @@ -32,9 +34,9 @@ inline detail::ScheduleTree* insertNodeAbove( inline detail::ScheduleTree* insertNodeBelow( detail::ScheduleTree* tree, ScheduleTreeUPtr&& node) { - CHECK_EQ(node->numChildren(), 0u); + TC_CHECK_EQ(node->numChildren(), 0u); auto numChildren = tree->numChildren(); - CHECK_LE(numChildren, 1u); + TC_CHECK_LE(numChildren, 1u); node->appendChildren(tree->detachChildren()); tree->appendChild(std::move(node)); return tree->child({0}); diff --git a/tc/core/polyhedral/schedule_transforms.cc b/tc/core/polyhedral/schedule_transforms.cc index d43363f03..31ed9a59f 100644 --- a/tc/core/polyhedral/schedule_transforms.cc +++ b/tc/core/polyhedral/schedule_transforms.cc @@ -28,6 +28,7 @@ #include "tc/external/isl.h" +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/polyhedral/functional.h" #include "tc/core/polyhedral/schedule_tree_elem.h" @@ -110,7 +111,7 @@ isl::union_set activeDomainPointsHelper( const ScheduleTree* root, const vector& nodes) { auto domainElem = root->elemAs(); - CHECK(domainElem) << "root must be a Domain node" << *root; + TC_CHECK(domainElem) << "root must be a Domain node" << *root; auto domain = domainElem->domain_; @@ -120,7 +121,7 @@ isl::union_set activeDomainPointsHelper( } else if (auto extensionElem = anc->elemAs()) { auto parentSchedule = prefixSchedule(root, anc); auto extension = extensionElem->extension_; - CHECK(parentSchedule) << "missing root domain node"; + TC_CHECK(parentSchedule) << "missing root domain node"; parentSchedule = parentSchedule.intersect_domain(domain); domain = domain.unite(parentSchedule.range().apply(extension)); } @@ -171,7 +172,7 @@ ScheduleTree* swapSubtree( ScheduleTree* relativeRoot, ScheduleTree* tree, ScheduleTreeUPtr& newTree) { - CHECK(relativeRoot != tree) << "Need a strict relative root to graft"; + TC_CHECK(relativeRoot != tree) << "Need a strict relative root to graft"; auto cpos = tree->positionRelativeTo(relativeRoot).back(); auto parent = tree->ancestor(relativeRoot, 1); auto rawPtr = newTree.get(); @@ -192,7 +193,7 @@ namespace { */ ScheduleTree* joinBandsHelper(ScheduleTree* st, bool& moveChildren) { moveChildren = false; - CHECK(st->elemAs()); + TC_CHECK(st->elemAs()); if (st->numChildren() != 1) { return st; } @@ -224,7 +225,7 @@ ScheduleTree* joinBands(ScheduleTree* st, bool permutable) { if (moveChildren) { // Just overwrite children and let shared pointers go out of scope auto children = st->detachChildren(); - CHECK_EQ(1u, children.size()) << "expected a sequence of bands"; + TC_CHECK_EQ(1u, children.size()) << "expected a sequence of bands"; st->appendChildren(children[0]->detachChildren()); } st->elemAs()->permutable_ = permutable; @@ -238,7 +239,7 @@ ScheduleTree* joinBandsIterative(ScheduleTree* st, bool permutable) { // Stupid private access hack, remove when moving to unique_ptr if (moveChildren) { auto children = st->detachChildren(); - CHECK_EQ(1u, children.size()) << "expected a sequence of bands"; + TC_CHECK_EQ(1u, children.size()) << "expected a sequence of bands"; st->appendChildren(children[0]->detachChildren()); } } @@ -270,12 +271,12 @@ void applyTileOptions(isl::ctx& ctx, TileOptions tileOptions) { ScheduleTree* bandSplit(ScheduleTree* relativeRoot, ScheduleTree* tree, size_t pos) { - CHECK(tree->elemAs()) << "Not a band:\n" << *tree; + TC_CHECK(tree->elemAs()) << "Not a band:\n" << *tree; auto band = tree->elemAs(); size_t n = band->nMember(); - CHECK_LT(0u, n) << "no bands to split"; - CHECK_LE(0u, pos) << "position out of bounds"; - CHECK_GE(n, pos) << "position out of bounds"; + TC_CHECK_LT(0u, n) << "no bands to split"; + TC_CHECK_LE(0u, pos) << "position out of bounds"; + TC_CHECK_GE(n, pos) << "position out of bounds"; // Detach and reattach children to avoid making copies. auto children = tree->detachChildren(); @@ -292,7 +293,7 @@ bandSplit(ScheduleTree* relativeRoot, ScheduleTree* tree, size_t pos) { ScheduleTree* bandSplitOut(ScheduleTree* relativeRoot, ScheduleTree* tree, size_t pos) { auto band = tree->elemAs(); - CHECK(band); + TC_CHECK(band); auto size = band->nMember(); if (pos != size - 1) { tree = bandSplit(relativeRoot, tree, pos + 1); @@ -320,13 +321,13 @@ ScheduleTree* bandTile( const vector& tileSizes, TileOptions tileOptions) { auto eb = st->elemAs(); - CHECK(eb) << "Not a band: " << *st; + TC_CHECK(eb) << "Not a band: " << *st; if (tileSizes.size() == 0) { return st; } auto& band = *eb; - CHECK(band.permutable_) << "Can't tile an non-permutable band" << band; + TC_CHECK(band.permutable_) << "Can't tile an non-permutable band" << band; auto ts = tileSizes; if (band.nMember() > ts.size()) { @@ -337,7 +338,7 @@ ScheduleTree* bandTile( << " entries: " << ts; ts.resize(band.nMember()); } - CHECK_EQ(band.nMember(), ts.size()) << "NYI: incorrect sizes: " << ts; + TC_CHECK_EQ(band.nMember(), ts.size()) << "NYI: incorrect sizes: " << ts; // TODO: adapt size // TODO: imperfectly nested loop tiling @@ -360,7 +361,7 @@ ScheduleTree* bandTile( } auto ebChild = childUPtr->elemAs(); - CHECK(ebChild) << "Not a band: " << *childUPtr; + TC_CHECK(ebChild) << "Not a band: " << *childUPtr; auto& childBand = *ebChild; // No need for isl_schedule_band_point, it's almost done if (tileOptions & TileOptions::ShiftPointLoops) { @@ -379,7 +380,7 @@ ScheduleTree* bandTile( ScheduleTree* bandScale(ScheduleTree* tree, const vector& scales) { auto eb = tree->elemAs(); - CHECK(eb) << "Not a band: " << *tree; + TC_CHECK(eb) << "Not a band: " << *tree; auto& band = *eb; // This mimics the behavior of bandTile... @@ -449,7 +450,7 @@ isl::multi_union_pw_aff infixScheduleMupa( const ScheduleTree* relativeRoot, const ScheduleTree* tree) { auto domainElem = root->elemAs(); - CHECK(domainElem); + TC_CHECK(domainElem); auto domain = domainElem->domain_.universe(); auto zero = isl::multi_val::zero(domain.get_space().set_from_params()); auto prefix = isl::multi_union_pw_aff(domain, zero); @@ -473,7 +474,7 @@ isl::multi_union_pw_aff partialScheduleMupa( const detail::ScheduleTree* root, const detail::ScheduleTree* tree) { auto band = tree->elemAs(); - CHECK(band); + TC_CHECK(band); return prefixScheduleMupa(root, tree).flat_range_product(band->mupa_); } @@ -484,7 +485,7 @@ void updateTopLevelContext(detail::ScheduleTree* root, isl::set context) { } auto contextElem = const_cast( root->child({0})->elemAs()); - CHECK(contextElem) << "Expected domain(context(any()))"; + TC_CHECK(contextElem) << "Expected domain(context(any()))"; contextElem->context_ = contextElem->context_ & context; } @@ -517,7 +518,7 @@ void insertSequenceBelow( const detail::ScheduleTree* root, detail::ScheduleTree* tree) { auto numChildren = tree->numChildren(); - CHECK_LE(numChildren, 1u); + TC_CHECK_LE(numChildren, 1u); auto filter = activeDomainPointsBelow(root, tree).universe(); auto node = ScheduleTree::makeFilter(filter, tree->detachChildren()); tree->appendChild(ScheduleTree::makeSequence(std::move(node))); @@ -546,7 +547,7 @@ detail::ScheduleTree* insertEmptyExtensionAbove( ScheduleTree* relativeRoot, ScheduleTree* st) { auto domain = root->elemAs(); - CHECK(domain); + TC_CHECK(domain); auto space = domain->domain_.get_space(); auto extension = isl::union_map::empty(space); return insertExtensionAbove(relativeRoot, st, extension); @@ -594,8 +595,8 @@ void insertExtensionAt( extensionTree = insertEmptyExtensionAbove(root, relativeRoot, seqNode); extensionNode = extensionTree->elemAs(); } - CHECK(extensionNode); - CHECK(seqNode->elemAs()); + TC_CHECK(extensionNode); + TC_CHECK(seqNode->elemAs()); extensionNode->extension_ = extensionNode->extension_.unite(extension); seqNode->insertChild(pos, std::move(filterNode)); } diff --git a/tc/core/polyhedral/schedule_tree-inl.h b/tc/core/polyhedral/schedule_tree-inl.h index f3e4e614d..bb7b79bed 100644 --- a/tc/core/polyhedral/schedule_tree-inl.h +++ b/tc/core/polyhedral/schedule_tree-inl.h @@ -15,6 +15,8 @@ */ #pragma once +#include "tc/core/check.h" + namespace tc { namespace polyhedral { namespace detail { @@ -23,14 +25,14 @@ inline ScheduleTreeUPtr ScheduleTree::makeMappingFilter( const std::vector& mappedIds, isl::union_pw_aff_list mappedAffs, std::vector&& children) { - CHECK_EQ(mappedIds.size(), static_cast(mappedAffs.n())) + TC_CHECK_EQ(mappedIds.size(), static_cast(mappedAffs.n())) << "expected as many mapped ids as affs"; ScheduleTreeElemMappingFilter::Mapping mapping; for (size_t i = 0, n = mappedAffs.n(); i < n; ++i) { mapping.emplace(mappedIds.at(i), mappedAffs.get(i)); } - CHECK_GE(mapping.size(), 1u) << "empty mapping"; - CHECK_EQ(mappedIds.size(), mapping.size()) + TC_CHECK_GE(mapping.size(), 1u) << "empty mapping"; + TC_CHECK_EQ(mappedIds.size(), mapping.size()) << "some id is used more than once in the mapping"; auto ctx = mappedIds[0].get_ctx(); ScheduleTreeUPtr res(new ScheduleTree(ctx)); diff --git a/tc/core/polyhedral/schedule_tree.cc b/tc/core/polyhedral/schedule_tree.cc index c4b916e8f..c011ead5e 100644 --- a/tc/core/polyhedral/schedule_tree.cc +++ b/tc/core/polyhedral/schedule_tree.cc @@ -27,6 +27,7 @@ #include "tc/external/isl.h" +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/polyhedral/functional.h" #include "tc/core/polyhedral/schedule_tree_elem.h" @@ -75,7 +76,7 @@ deque findDescendant( vector positionRelativeToSubtree( const ScheduleTree* relativeRoot, const ScheduleTree* target) { - CHECK(relativeRoot != target) + TC_CHECK(relativeRoot != target) << "Need a strict relative root to find position"; auto res = findDescendant(relativeRoot, target); return vector{res.begin(), res.end()}; @@ -90,7 +91,7 @@ vector constAncestorsInSubTree( vector cp(positionRelativeToSubtree(relativeRoot, target)); if (cp.size() == 0) { // Special case, this must be the root - CHECK_EQ(relativeRoot, target); + TC_CHECK_EQ(relativeRoot, target); return {}; } vector res(cp.size() + 1, nullptr); @@ -100,11 +101,11 @@ vector constAncestorsInSubTree( res[i + 1] = res[i]->child({cp[i]}); } // Check last element is self for consistency - CHECK_EQ(res.back(), target) + TC_CHECK_EQ(res.back(), target) << "Could not find " << *target << " under " << *relativeRoot << "\n"; // Resize to drop self, and check again for consistency res.resize(cp.size()); - CHECK_NE(res.back(), target); + TC_CHECK_NE(res.back(), target); return res; } @@ -148,8 +149,8 @@ ScheduleTree* ScheduleTree::child(const vector& positions) { const ScheduleTree* ScheduleTree::child(const vector& positions) const { auto st = this; for (auto pos : positions) { - CHECK_LE(0u, pos) << "Reached a leaf"; - CHECK_GT(st->children_.size(), pos) << "Out of children bounds"; + TC_CHECK_LE(0u, pos) << "Reached a leaf"; + TC_CHECK_GT(st->children_.size(), pos) << "Out of children bounds"; st = st->children_[pos].get(); } return st; @@ -165,9 +166,9 @@ ScheduleTree* ScheduleTree::ancestor( const ScheduleTree* ScheduleTree::ancestor( const ScheduleTree* relativeRoot, size_t generations) const { - CHECK_LT(0u, generations) << "Nonpositive ancestor generation"; + TC_CHECK_LT(0u, generations) << "Nonpositive ancestor generation"; auto as = constAncestorsInSubTree(relativeRoot, this); - CHECK_GE(as.size(), generations) << "Out of ancestors bounds"; + TC_CHECK_GE(as.size(), generations) << "Out of ancestors bounds"; return as[as.size() - generations]; } @@ -210,7 +211,7 @@ std::unique_ptr ScheduleTree::makeBand( ScheduleTreeUPtr ScheduleTree::makeEmptyBand(const ScheduleTree* root) { auto domain = root->elemAs(); - CHECK(domain); + TC_CHECK(domain); auto space = domain->domain_.get_space().set_from_params(); auto mv = isl::multi_val::zero(space); auto zero = isl::multi_union_pw_aff(domain->domain_, mv); @@ -359,7 +360,7 @@ bool ScheduleTree::operator==(const ScheduleTree& other) const { if (!elemEquals(elem_.get(), other.elem_.get(), type_)) { return false; } - CHECK(!other.elemAs()) + TC_CHECK(!other.elemAs()) << "NYI: isl_node_type::set comparison"; for (size_t i = 0; i < children_.size(); ++i) { if (*children_[i] != *other.children_[i]) { diff --git a/tc/core/polyhedral/schedule_tree.h b/tc/core/polyhedral/schedule_tree.h index da3211985..981ad9c19 100644 --- a/tc/core/polyhedral/schedule_tree.h +++ b/tc/core/polyhedral/schedule_tree.h @@ -20,6 +20,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/polyhedral/options.h" #include "tc/core/polyhedral/schedule_tree_elem.h" #include "tc/core/utils/vararg.h" @@ -135,9 +136,9 @@ struct ScheduleTree { // Swap a tree with with the given tree. void swapChild(size_t pos, ScheduleTreeUPtr& swappee) { - CHECK_GE(pos, 0u) << "position out of children bounds"; - CHECK_LE(pos, children_.size()) << "position out of children bounds"; - CHECK(swappee.get()) << "Cannot swap in a null tree"; + TC_CHECK_GE(pos, 0u) << "position out of children bounds"; + TC_CHECK_LE(pos, children_.size()) << "position out of children bounds"; + TC_CHECK(swappee.get()) << "Cannot swap in a null tree"; std::swap(children_[pos], swappee); } @@ -150,10 +151,10 @@ struct ScheduleTree { // Manipulators for the list of children. void insertChildren(size_t pos, std::vector&& children) { - CHECK_GE(pos, 0u) << "position out of children bounds"; - CHECK_LE(pos, children_.size()) << "position out of children bounds"; + TC_CHECK_GE(pos, 0u) << "position out of children bounds"; + TC_CHECK_LE(pos, children_.size()) << "position out of children bounds"; for (const auto& c : children) { - CHECK(c.get()) << "inserting null or moved-from child"; + TC_CHECK(c.get()) << "inserting null or moved-from child"; } children_.insert( @@ -177,8 +178,8 @@ struct ScheduleTree { } ScheduleTreeUPtr detachChild(size_t pos) { - CHECK_GE(pos, 0u) << "position out of children bounds"; - CHECK_LT(pos, children_.size()) << "position out of children bounds"; + TC_CHECK_GE(pos, 0u) << "position out of children bounds"; + TC_CHECK_LT(pos, children_.size()) << "position out of children bounds"; ScheduleTreeUPtr child = std::move(children_[pos]); children_.erase(children_.begin() + pos); @@ -199,8 +200,8 @@ struct ScheduleTree { } ScheduleTreeUPtr replaceChild(size_t pos, ScheduleTreeUPtr&& child) { - CHECK_GE(pos, 0u) << "position out of children bounds"; - CHECK_LT(pos, children_.size()) << "position out of children bounds"; + TC_CHECK_GE(pos, 0u) << "position out of children bounds"; + TC_CHECK_LT(pos, children_.size()) << "position out of children bounds"; ScheduleTreeUPtr oldChild = std::move(children_[pos]); children_[pos] = std::move(child); @@ -242,7 +243,7 @@ struct ScheduleTree { inline size_t positionInParent(const ScheduleTree* parent) const { auto p = positionRelativeTo(parent); - CHECK_EQ(1u, p.size()) << *parent << " is not the parent of " << *this; + TC_CHECK_EQ(1u, p.size()) << *parent << " is not the parent of " << *this; return p[0]; } @@ -346,7 +347,7 @@ struct ScheduleTree { // Flatten nested nodes of the same type. void flattenSequenceOrSet() { // This should be enforced by the type system... - CHECK( + TC_CHECK( type_ == detail::ScheduleTreeType::Sequence || type_ == detail::ScheduleTreeType::Set); diff --git a/tc/core/polyhedral/schedule_tree_elem.cc b/tc/core/polyhedral/schedule_tree_elem.cc index e418b8a31..9cb6cd5df 100644 --- a/tc/core/polyhedral/schedule_tree_elem.cc +++ b/tc/core/polyhedral/schedule_tree_elem.cc @@ -23,6 +23,7 @@ #include +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/flags.h" #include "tc/core/polyhedral/schedule_isl_conversion.h" @@ -136,13 +137,13 @@ std::unique_ptr ScheduleTreeElemBand::fromMultiUnionPwAff( // Return the number of scheduling dimensions in the band size_t ScheduleTreeElemBand::nMember() const { size_t res = mupa_.size(); - CHECK_EQ(res, coincident_.size()); - CHECK_EQ(res, unroll_.size()); + TC_CHECK_EQ(res, coincident_.size()); + TC_CHECK_EQ(res, unroll_.size()); return res; } size_t ScheduleTreeElemBand::nOuterCoincident() const { - CHECK_EQ(nMember(), coincident_.size()); + TC_CHECK_EQ(nMember(), coincident_.size()); size_t i; for (i = 0; i < nMember(); ++i) { if (!coincident_[i]) { @@ -153,9 +154,9 @@ size_t ScheduleTreeElemBand::nOuterCoincident() const { } void ScheduleTreeElemBand::drop(size_t pos, size_t n) { - CHECK_LE(0u, n) << "range out of bounds"; - CHECK_LE(0u, pos) << "range out of bounds"; - CHECK_GE(nMember(), pos + n) << "range out of bounds"; + TC_CHECK_LE(0u, n) << "range out of bounds"; + TC_CHECK_LE(0u, pos) << "range out of bounds"; + TC_CHECK_GE(nMember(), pos + n) << "range out of bounds"; auto nBegin = nMember(); mupa_ = mupa_.drop_dims(isl::dim_type::set, pos, n); @@ -167,7 +168,7 @@ void ScheduleTreeElemBand::drop(size_t pos, size_t n) { coincident_.resize(nBegin - n); std::copy(unroll_.begin() + pos + n, unroll_.end(), unroll_.begin() + pos); unroll_.resize(nBegin - n); - CHECK_EQ(nBegin - n, nMember()); + TC_CHECK_EQ(nBegin - n, nMember()); } bool ScheduleTreeElemBand::operator==(const ScheduleTreeElemBand& other) const { diff --git a/tc/core/polyhedral/schedule_tree_elem.h b/tc/core/polyhedral/schedule_tree_elem.h index f1ab3153b..8f1164b53 100644 --- a/tc/core/polyhedral/schedule_tree_elem.h +++ b/tc/core/polyhedral/schedule_tree_elem.h @@ -22,6 +22,7 @@ #include "tc/external/isl.h" +#include "tc/core/check.h" #include "tc/core/polyhedral/mapping_types.h" namespace tc { @@ -152,11 +153,11 @@ struct ScheduleTreeElemMappingFilter : public ScheduleTreeElemFilter { : ScheduleTreeElemFilter(eb.filter_), mapping(eb.mapping) {} ScheduleTreeElemMappingFilter(const Mapping& mapping) : ScheduleTreeElemFilter(isl::union_set()), mapping(mapping) { - CHECK_GT(mapping.size(), 0u) << "empty mapping filter"; + TC_CHECK_GT(mapping.size(), 0u) << "empty mapping filter"; auto domain = mapping.cbegin()->second.domain(); for (auto& kvp : mapping) { - CHECK(domain.is_equal(kvp.second.domain())); + TC_CHECK(domain.is_equal(kvp.second.domain())); } filter_ = domain.universe(); for (auto& kvp : mapping) { diff --git a/tc/core/polyhedral/schedule_tree_matcher-inl.h b/tc/core/polyhedral/schedule_tree_matcher-inl.h index c4f7bb6ef..4364804cc 100644 --- a/tc/core/polyhedral/schedule_tree_matcher-inl.h +++ b/tc/core/polyhedral/schedule_tree_matcher-inl.h @@ -15,6 +15,7 @@ */ #pragma once +#include "tc/core/check.h" #include "tc/core/polyhedral/schedule_tree.h" #include "tc/core/polyhedral/schedule_tree_elem.h" @@ -258,7 +259,8 @@ inline bool matchOne( // We still need to check well-formedness of the matcher (i.e. no wildcards // except in the last position) for (size_t i = 0; i < matcher.children_.size(); ++i) { - CHECK(!matcher.children_[i].wildcard || i == matcher.children_.size() - 1) + TC_CHECK( + !matcher.children_[i].wildcard || i == matcher.children_.size() - 1) << "Error in matcher structure, wildcard must be the last child!"; if (!matchOne(matcher.children_[i], tree->child({i}))) { return false; diff --git a/tc/core/polyhedral/scop.cc b/tc/core/polyhedral/scop.cc index f6686fd73..13d20cffa 100644 --- a/tc/core/polyhedral/scop.cc +++ b/tc/core/polyhedral/scop.cc @@ -23,6 +23,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/halide2isl.h" #include "tc/core/polyhedral/functional.h" #include "tc/core/polyhedral/memory_promotion.h" @@ -43,7 +44,7 @@ using ScopUPtr = std::unique_ptr; ScopUPtr Scop::makeScop( isl::ctx ctx, const tc2halide::HalideComponents& components) { - CHECK(components.stmt.defined()); + TC_CHECK(components.stmt.defined()); halide2isl::SymbolTable sym = halide2isl::makeSymbolTable(components); @@ -78,14 +79,14 @@ ScopUPtr Scop::makeScop(isl::ctx ctx, const lang::TreeRef& treeRef) { isl::union_set& Scop::domainRef() { auto dom = scheduleRoot()->elemAs(); - CHECK(dom) << "root is not a domain in: " << *scheduleRoot(); + TC_CHECK(dom) << "root is not a domain in: " << *scheduleRoot(); // TODO: activate this when the invariant has a chance of working (i.e. we // don't use a Context node for specifying parameter values that iterate in // spacetime). // TODO: find a proper place for the invariant. // auto noCont = // scheduleRoot()->child({0})->elemAs(); - // CHECK(!noCont) << "root is not a domain in: " << *scheduleRoot(); + // TC_CHECK(!noCont) << "root is not a domain in: " << *scheduleRoot(); return dom->domain_; } @@ -139,7 +140,7 @@ void checkFiltersDisjointStatements(const ScheduleTree* root) { isl::union_set alreadyVisitedStmts; for (auto child : node->children()) { auto filterNode = child->elemAsBase(); - CHECK(filterNode) << "expected children of seqence to be filters"; + TC_CHECK(filterNode) << "expected children of seqence to be filters"; auto filter = filterNode->filter_.universe(); if (!alreadyVisitedStmts.get()) { alreadyVisitedStmts = filter; @@ -149,7 +150,7 @@ void checkFiltersDisjointStatements(const ScheduleTree* root) { // but only to a part of it. Possible solution -- introduce "scope" // mark nodes into the schedule tree that will contain information // about the promotion and process these marks when generating the AST. - CHECK(alreadyVisitedStmts.intersect(filter).is_empty()) + TC_CHECK(alreadyVisitedStmts.intersect(filter).is_empty()) << "filters are expected to be disjoint as stmt level"; alreadyVisitedStmts = alreadyVisitedStmts.unite(filter); } @@ -207,7 +208,7 @@ void Scop::insertSyncsAroundCopies(ScheduleTree* tree) { // Insert syncs before and after copies (FIXME: this is excessive) auto seqNode = tree->child({0, 0}); - CHECK(seqNode->elemAs()) + TC_CHECK(seqNode->elemAs()) << "unexpected tree structure"; int foundMainComputations = 0; @@ -215,7 +216,7 @@ void Scop::insertSyncsAroundCopies(ScheduleTree* tree) { for (size_t i = 0; i < seqNode->numChildren(); ++i) { auto filterNode = seqNode->child({i})->elemAs(); - CHECK(filterNode) << "expected filters below sequence"; + TC_CHECK(filterNode) << "expected filters below sequence"; auto filters = isl::UnionAsVector(filterNode->filter_); bool isCopyFilter = filters.size() == 1 && filters[0].has_tuple_name() && (filters[0].get_tuple_name() == kReadIdName || @@ -227,7 +228,7 @@ void Scop::insertSyncsAroundCopies(ScheduleTree* tree) { if (!isCopyFilter) { ++foundMainComputations; } - CHECK_LT(foundMainComputations, 2) + TC_CHECK_LT(foundMainComputations, 2) << "copies are interleaved with computation" << *seqNode; if (filters[0].get_tuple_name() != lastTupleName) { lastTupleName = filters[0].get_tuple_name(); @@ -269,7 +270,7 @@ std::vector Scop::getParameterValues() const { std::vector paramValues; for (auto const& param : halide.params) { auto name = param.name(); - CHECK(parameterValues.count(name) == 1); + TC_CHECK(parameterValues.count(name) == 1); paramValues.push_back(parameterValues.at(name)); } return paramValues; @@ -409,7 +410,7 @@ namespace { */ detail::ScheduleTree* setPermutable(detail::ScheduleTree* tree) { auto band = tree->elemAs(); - CHECK(band); + TC_CHECK(band); band->permutable_ = true; return tree; } @@ -490,7 +491,7 @@ const Halide::OutputImageParam& Scop::findArgument(isl::id id) const { } } - CHECK(false) << "name \"" << name << "\" not found"; + TC_CHECK(false) << "name \"" << name << "\" not found"; return *halide.inputs.begin(); } diff --git a/tc/core/polyhedral/scop.h b/tc/core/polyhedral/scop.h index e163dbe90..8eaf5c12c 100644 --- a/tc/core/polyhedral/scop.h +++ b/tc/core/polyhedral/scop.h @@ -22,6 +22,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/constants.h" #include "tc/core/halide2isl.h" #include "tc/core/mapping_options.h" @@ -239,7 +240,7 @@ struct Scop { return makeSyncId(); break; default: - CHECK(level != SyncLevel::None); + TC_CHECK(level != SyncLevel::None); return isl::id(); } } @@ -292,8 +293,8 @@ struct Scop { domain().get_ctx(), std::string("red_update") + std::to_string(uid)); auto reductionInitId = isl::id( domain().get_ctx(), std::string("red_init") + std::to_string(uid)); - CHECK_EQ(0u, treeSyncUpdateMap.count(treeSyncId)); - CHECK_EQ(0u, defaultReductionInitMap.count(treeSyncId)); + TC_CHECK_EQ(0u, treeSyncUpdateMap.count(treeSyncId)); + TC_CHECK_EQ(0u, defaultReductionInitMap.count(treeSyncId)); treeSyncUpdateMap.emplace(treeSyncId, updateId); defaultReductionInitMap.emplace(treeSyncId, reductionInitId); @@ -319,7 +320,7 @@ struct Scop { return treeSyncUpdateMap.at(p.first); } } - CHECK(false) << "not found"; + TC_CHECK(false) << "not found"; return id; } @@ -334,7 +335,7 @@ struct Scop { size_t reductionUpdatePos(isl::id id) const { size_t pos = 0; - CHECK(isReductionUpdate(id)); + TC_CHECK(isReductionUpdate(id)); for (const auto& kvp : treeSyncUpdateMap) { if (id == kvp.second) { return pos; diff --git a/tc/core/polyhedral/separation.cc b/tc/core/polyhedral/separation.cc index 7027bfb89..c5d38436e 100644 --- a/tc/core/polyhedral/separation.cc +++ b/tc/core/polyhedral/separation.cc @@ -15,6 +15,7 @@ */ #include "tc/core/polyhedral/separation.h" +#include "tc/core/check.h" #include "tc/external/isl.h" @@ -39,7 +40,7 @@ isl::union_set partialTargetTiles( // Mapping between prefix values and target values // for some common domain element // P -> T - CHECK(domain.is_subset(scheduleMap.domain())); + TC_CHECK(domain.is_subset(scheduleMap.domain())); auto target = domain.apply(scheduleMap).unwrap(); // Mapping between prefix values and target values // for some common domain element, extended to complete target tiles. diff --git a/tc/core/tc2halide.cc b/tc/core/tc2halide.cc index 76adea3d5..01da89edf 100644 --- a/tc/core/tc2halide.cc +++ b/tc/core/tc2halide.cc @@ -15,6 +15,7 @@ */ #include +#include "tc/core/check.h" #include "tc/core/flags.h" #include "tc/core/tc2halide.h" #include "tc/lang/parser.h" @@ -91,7 +92,7 @@ void translateParam( } dims.push_back(Variable::make(Int(32), p.name(), p)); } else { - CHECK(d_->kind() == lang::TK_CONST); + TC_CHECK(d_->kind() == lang::TK_CONST); int32_t value = lang::Const(d_).value(); dims.push_back(Expr(value)); } @@ -279,7 +280,7 @@ void forwardBoundsInference( // Create inequalities that assert this is not an out-of-bounds access. if (op->call_type == Call::Halide) { - CHECK(op->func.defined()) + TC_CHECK(op->func.defined()) << "Expected a Call of type Halide to have an associated Function\n"; const auto& it = bounds.find(Function(op->func)); if (it != bounds.end()) { @@ -289,7 +290,7 @@ void forwardBoundsInference( const auto& it = b.find(dim); if (it != b.end()) { Interval interval = it->second; - CHECK(interval.is_bounded()) + TC_CHECK(interval.is_bounded()) << "Expected explicit constraints on every dimension of every Func\n"; result.push_back(op->args[i] >= interval.min); result.push_back(op->args[i] <= interval.max); @@ -297,10 +298,10 @@ void forwardBoundsInference( } } } else if (op->call_type == Call::Image) { - CHECK(op->param.defined()) + TC_CHECK(op->param.defined()) << "Expected a Call of type Image to have an associated Parameter\n"; for (size_t i = 0; i < op->args.size(); i++) { - CHECK( + TC_CHECK( op->param.min_constraint(i).defined() && op->param.extent_constraint(i).defined()) << "Expected explicit constraints on every dimension of every input\n"; diff --git a/tc/core/tensor-inl.h b/tc/core/tensor-inl.h index 0abef002b..32d970133 100644 --- a/tc/core/tensor-inl.h +++ b/tc/core/tensor-inl.h @@ -18,9 +18,9 @@ #include #include -#include #include +#include "tc/core/check.h" #include "tc/proto/compcache.pb.h" namespace tc { @@ -215,7 +215,7 @@ inline std::vector extractRawPtrs( inline std::string toString(const DLDataType& t) { if (t.lanes != 1) { - CHECK(false) << "NYI: toString for >1 lanes"; + TC_CHECK(false) << "NYI: toString for >1 lanes"; } switch (t.code) { case DLDataTypeCode::kDLFloat: @@ -247,7 +247,8 @@ inline std::string toString(const DLDataType& t) { } break; } - CHECK(false) << "NYI: toString for type: " << t.code << ", bits: " << t.bits; + TC_CHECK(false) << "NYI: toString for type: " << t.code + << ", bits: " << t.bits; return ""; } diff --git a/tc/core/utils/type_traits.h b/tc/core/utils/type_traits.h new file mode 100644 index 000000000..5a93d0c8f --- /dev/null +++ b/tc/core/utils/type_traits.h @@ -0,0 +1,23 @@ +#include + +namespace tc { +// WG21 N3911 2.3 Implementation workaround +// (Make sure template argument is always used.) +template +struct voider { + using type = void; +}; +template +using void_t = typename voider::type; + +template +struct is_std_container : std::false_type {}; + +template +struct is_std_container< + T, + void_t< + decltype(std::declval().begin()), + decltype(std::declval().end()), + typename T::value_type>> : std::true_type {}; +} // namespace tc diff --git a/tc/examples/blockdiagperm.cc b/tc/examples/blockdiagperm.cc index 92e7fe12f..a6a6d0aaf 100644 --- a/tc/examples/blockdiagperm.cc +++ b/tc/examples/blockdiagperm.cc @@ -25,6 +25,7 @@ #include "tc/aten/aten_autotuner.h" #include "tc/aten/aten_compiler.h" #include "tc/autotuner/genetic_search.h" +#include "tc/core/check.h" #include "tc/core/cuda/cuda_mapping_options.h" #include "tc/core/cuda/cuda_tc_executor.h" #include "tc/core/flags.h" @@ -79,7 +80,7 @@ def blockdiagperm2dfissioned_2(float(B, N) I, int32(N) Idx) -> (O) { ATenGeneticCudaTuner geneticAutotuneATen(tc); auto bestOption = geneticAutotuneATen.tune( "blockdiagperm2dfissioned_1", {I, W}, options, FLAGS_proto_path); - CHECK_GT(bestOption.size(), 0u); + TC_CHECK_GT(bestOption.size(), 0u); auto pExecutor = tc::aten::compile( tc, "blockdiagperm2dfissioned_1", {I, W}, bestOption[0]); @@ -92,7 +93,7 @@ def blockdiagperm2dfissioned_2(float(B, N) I, int32(N) Idx) -> (O) { at::Tensor Idx = at::CPU(at::kInt).randperm({500}).toBackend(at::kCUDA); auto bestOption2 = geneticAutotuneATen.tune( "blockdiagperm2dfissioned_2", {O, Idx}, options, FLAGS_proto_path); - CHECK_GT(bestOption2.size(), 0u); + TC_CHECK_GT(bestOption2.size(), 0u); auto pExecutor2 = tc::aten::compile( tc, "blockdiagperm2dfissioned_2", {O, Idx}, bestOption2[0]); diff --git a/tc/examples/group_normalization.cc b/tc/examples/group_normalization.cc index aa7373fae..05823e6a9 100644 --- a/tc/examples/group_normalization.cc +++ b/tc/examples/group_normalization.cc @@ -27,6 +27,7 @@ #include "tc/aten/aten_autotuner.h" #include "tc/aten/aten_compiler.h" #include "tc/autotuner/genetic_search.h" +#include "tc/core/check.h" #include "tc/core/cpu/cpu_mapping_options.h" #include "tc/core/cpu/cpu_tc_executor.h" #include "tc/core/cuda/cuda_mapping_options.h" @@ -119,7 +120,7 @@ def group_normalization( geneticAutotuneATen(tc); auto bestOption = geneticAutotuneATen.tune( "group_normalization", {I, gamma, beta}, baseOptions, FLAGS_proto_path); - CHECK_GT(bestOption.size(), 0u); + TC_CHECK_GT(bestOption.size(), 0u); // 5. Compile and run the TC with the best option. // Outputs get allocated; could also be pre-allocated and passed. diff --git a/tc/examples/tensordot.cc b/tc/examples/tensordot.cc index 4c49806ef..fa15d2a5b 100644 --- a/tc/examples/tensordot.cc +++ b/tc/examples/tensordot.cc @@ -27,6 +27,7 @@ #include "tc/aten/aten_autotuner.h" #include "tc/aten/aten_compiler.h" #include "tc/autotuner/genetic_search.h" +#include "tc/core/check.h" #include "tc/core/cpu/cpu_mapping_options.h" #include "tc/core/cpu/cpu_tc_executor.h" #include "tc/core/cuda/cuda_mapping_options.h" @@ -57,7 +58,7 @@ def tensordot(float(N, C1, C2, H, W) I0, geneticAutotuneATen(tc); auto bestOption = geneticAutotuneATen.tune( "tensordot", {I0, I1}, naiveOptions, FLAGS_proto_path); - CHECK_GT(bestOption.size(), 0u); + TC_CHECK_GT(bestOption.size(), 0u); // 4. Compile and run the TC with the best option. // Outputs get allocated; could also be pre-allocated and passed. diff --git a/tc/external/detail/islpp.h b/tc/external/detail/islpp.h index 5983c2263..c7dae3fe7 100644 --- a/tc/external/detail/islpp.h +++ b/tc/external/detail/islpp.h @@ -26,6 +26,7 @@ #include +#include "tc/core/check.h" #include "tc/core/islpp_wrap.h" namespace isl { @@ -284,7 +285,7 @@ inline bool operator!=(const isl::id& id1, const isl::id& id2) { template isl::multi_val makeMultiVal(isl::space s, const std::vector& vals) { isl::multi_val mv = isl::multi_val::zero(s); - CHECK_EQ(vals.size(), s.dim(isl::dim_type::set)); + TC_CHECK_EQ(vals.size(), s.dim(isl::dim_type::set)); for (size_t i = 0; i < vals.size(); ++i) { mv = mv.set_val(i, isl::val(s.get_ctx(), vals[i])); } diff --git a/tc/lang/gfg/parser.cc b/tc/lang/gfg/parser.cc index 976b14879..f4449943c 100644 --- a/tc/lang/gfg/parser.cc +++ b/tc/lang/gfg/parser.cc @@ -15,18 +15,15 @@ */ #include -// #include - #include #include +#include "tc/core/check.h" using namespace std; namespace tc { namespace parser { -#define CHECK(x) ; - int uid() { static int id = 0; return ++id; @@ -34,15 +31,15 @@ int uid() { Node& GFG::addNode(const string& name) { nodes.emplace_back(new Node(name)); - CHECK(nodes.size() == nodes.back().id); + TC_CHECK(nodes.size() == nodes.back().id); name2NodeId.insert( make_pair(string(name), int(nodes.back()->id))); return *nodes.back(); } const Edge& GFG::addEdge(Node& s, const string& transition, const Node& t) { - CHECK(name2Node.at(s.name) != name2Node.end()); - CHECK(name2Node.at(t.name) != name2Node.end()); + TC_CHECK(name2Node.at(s.name) != name2Node.end()); + TC_CHECK(name2Node.at(t.name) != name2Node.end()); s.outEdges.emplace_back(s, transition, t); return s.outEdges.back(); } @@ -212,6 +209,5 @@ GFG GFG::makeGFG(const string& grammar) { return res; } -#undef CHECK -} +} // namespace parser } // namespace tc diff --git a/tc/library/convolution.h b/tc/library/convolution.h index 85e1a3c30..b542cf87d 100644 --- a/tc/library/convolution.h +++ b/tc/library/convolution.h @@ -15,6 +15,7 @@ */ #pragma once +#include "tc/core/check.h" #include "tc/library/common.h" namespace tc { @@ -31,7 +32,7 @@ constexpr static auto CONVOLUTION2D_TC = R"TC( } // namespace std::string makeConvolution2DTc(int strideH, int strideW) { - CHECK(strideH > 0 && strideW > 0) << "Stride must be greater than 0"; + TC_CHECK(strideH > 0 && strideW > 0) << "Stride must be greater than 0"; std::string tcStr; tcStr = CONVOLUTION2D_TC; tcStr = replaceString(tcStr, "${sh}", std::to_string(strideH)); diff --git a/tc/library/group_convolution.h b/tc/library/group_convolution.h index af81a4dab..5181c0399 100644 --- a/tc/library/group_convolution.h +++ b/tc/library/group_convolution.h @@ -15,6 +15,7 @@ */ #pragma once +#include "tc/core/check.h" #include "tc/library/common.h" namespace tc { @@ -34,7 +35,7 @@ constexpr static auto GROUP_CONVOLUTION2D_TC = R"TC( } // namespace std::string makeGroupConvolution2DTc(int strideH, int strideW) { - CHECK(strideH > 0 && strideW > 0) << "Stride must be greater than 0"; + TC_CHECK(strideH > 0 && strideW > 0) << "Stride must be greater than 0"; std::string tcStr; tcStr = GROUP_CONVOLUTION2D_TC; tcStr = replaceString(tcStr, "", std::to_string(strideH)); diff --git a/tensor_comprehensions/pybinds/tc.cc b/tensor_comprehensions/pybinds/tc.cc index 978c3d51e..a085c9bce 100644 --- a/tensor_comprehensions/pybinds/tc.cc +++ b/tensor_comprehensions/pybinds/tc.cc @@ -244,7 +244,7 @@ PYBIND11_MODULE(tc, m) { const std::string& entryPoint, py::list& inputs, py::list& outputs) { - CHECK_GE(outputs.size(), 1u); + TC_CHECK_GE(outputs.size(), 1u); auto atOutputs = getATenTensors(outputs, dlpack); tc::aten::uncheckedRun( *instance.compiled.at(entryPoint), diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d7539e96e..94c843f9e 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -32,6 +32,9 @@ foreach(i ${CORE_TEST_FILES}) target_link_libraries(${i} ${GTEST_LIBRARIES} tc_core pthread) endforeach() +add_executable(test_check test_check.cc) +target_link_libraries(test_check ${GTEST_LIBRARIES} tc_core) + ################################################################################ # CPP cuda_mapper crosscompile tests ################################################################################ diff --git a/test/caffe2/test_harness.cc b/test/caffe2/test_harness.cc index 04ab01b95..7653362da 100644 --- a/test/caffe2/test_harness.cc +++ b/test/caffe2/test_harness.cc @@ -45,8 +45,8 @@ void CheckEqual( << " for Tensor " << Texpected.DebugString() << " at position " << i; } else { // From glog's glog/src/glog/logging.h.in - // #define CHECK_NEAR(val1, val2, margin) - // CHECK_NEAR is actualy absolute!!! + // #define TC_CHECK_NEAR(val1, val2, margin) + // TC_CHECK_NEAR is actualy absolute!!! ASSERT_NEAR( Texpected.data()[i + offsetInExpected], Ttested.data()[i + offsetInTested], diff --git a/test/cuda/test_compile_and_run.cc b/test/cuda/test_compile_and_run.cc index 3d7ae89e8..9dc2bde61 100644 --- a/test/cuda/test_compile_and_run.cc +++ b/test/cuda/test_compile_and_run.cc @@ -24,6 +24,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 "tc/core/cuda/cuda_tc_executor.h" #include "tc/library/common.h" @@ -271,7 +272,7 @@ def cast(float(M,N) A, int32 four) -> (int32(M,N) output) { tc::CudaMappingOptions::makeNaiveMappingOptions(), {a, b}); auto r = outputs[0].sub(at::CUDA(at::kInt).ones({2, 4}) + 4).max().toCFloat(); - CHECK_EQ(r, 0); + TC_CHECK_EQ(r, 0); } int main(int argc, char** argv) { diff --git a/test/cuda/test_corner_cases.cc b/test/cuda/test_corner_cases.cc index 5c0aa2fe8..6cbf60751 100644 --- a/test/cuda/test_corner_cases.cc +++ b/test/cuda/test_corner_cases.cc @@ -24,6 +24,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 "tc/core/cuda/cuda_tc_executor.h" #include "tc/library/common.h" @@ -185,7 +186,7 @@ TEST(TestCornerCases, E15){ {a, b}, \ {c}); \ auto r = at::Scalar(a).toFloat() op at::Scalar(b).toFloat(); \ - CHECK_EQ(r, at::Scalar(c[0]).toFloat()); \ + TC_CHECK_EQ(r, at::Scalar(c[0]).toFloat()); \ } GEN_COMPARATOR(<=) GEN_COMPARATOR(>=) GEN_COMPARATOR(==) GEN_COMPARATOR(!=) @@ -206,7 +207,7 @@ TEST(TestCornerCases, E16){ {c}); \ auto r = !(at::Scalar(a).toFloat() < .5) op at::Scalar(b).toFloat() > .5; \ ; \ - CHECK_EQ(r, at::Scalar(c[0]).toFloat()); \ + TC_CHECK_EQ(r, at::Scalar(c[0]).toFloat()); \ } GEN_BOOLS(||) GEN_BOOLS(&&)} @@ -215,7 +216,7 @@ TEST(TestCornerCases, E17) { auto r = F(1); Succeed( "def f(float(1) a) -> (b) { b(i) = 4.0 where exists a(i) }", {F(1)}, {r}); - CHECK_EQ(at::Scalar(r[0]).toFloat(), 4); + TC_CHECK_EQ(at::Scalar(r[0]).toFloat(), 4); } TEST(TestCornerCases, E18) { @@ -223,7 +224,7 @@ TEST(TestCornerCases, E18) { auto r = F(1); Succeed( "def f(float(1) a) -> (b) { b(i) = 2*foo where foo = a(i) }", {a}, {r}); - CHECK_EQ(at::Scalar(r[0]).toFloat(), at::Scalar(a[0]).toFloat() * 2); + TC_CHECK_EQ(at::Scalar(r[0]).toFloat(), at::Scalar(a[0]).toFloat() * 2); } TEST(TestCornerCases, E19) { Fail( @@ -248,7 +249,7 @@ TEST(TestCornerCases, E21) { "def f(float(1) a, float(1) b) -> (c) { c(i) = max(a(i), b(i)) }", {a, b}, {c}); - CHECK_EQ( + TC_CHECK_EQ( fmaxf(at::Scalar(a[0]).toFloat(), at::Scalar(b[0]).toFloat()), at::Scalar(c[0]).toFloat()); } @@ -261,7 +262,7 @@ TEST(TestCornerCases, E22) { "def f(float(1) a, float(1) b) -> (c) { c(i) = min(a(i), b(i)) }", {a, b}, {c}); - CHECK_EQ( + TC_CHECK_EQ( fminf(at::Scalar(a[0]).toFloat(), at::Scalar(b[0]).toFloat()), at::Scalar(c[0]).toFloat()); } @@ -275,7 +276,7 @@ TEST(TestCornerCases, E23) { "def f(float(1) a, float(1) b, float(1) c) -> (d) { d(i) = min(a(i), max(b(i), c(i))) }", {a, b, c}, {d}); - CHECK_EQ( + TC_CHECK_EQ( fminf( at::Scalar(a[0]).toFloat(), fmaxf(at::Scalar(b[0]).toFloat(), at::Scalar(c[0]).toFloat())), @@ -312,7 +313,7 @@ i)) )TC", {a, b, c}, {d}); - CHECK_EQ( + TC_CHECK_EQ( fminf( at::Scalar(a[0]).toFloat(), fmaxf(at::Scalar(b[0]).toFloat(), at::Scalar(c[0]).toFloat())) + diff --git a/test/cuda/test_tc_mapper.cc b/test/cuda/test_tc_mapper.cc index e89756aea..a1f303bd7 100644 --- a/test/cuda/test_tc_mapper.cc +++ b/test/cuda/test_tc_mapper.cc @@ -20,6 +20,7 @@ #include "tc/aten/aten.h" #include "tc/aten/aten_compiler.h" +#include "tc/core/check.h" #include "tc/core/cuda/cuda.h" #include "tc/core/cuda/cuda_tc_executor.h" #include "tc/core/exceptions.h" @@ -393,7 +394,7 @@ def fun(float(B, R) LUT, int32(B, N) I) -> (O) { float correct = 0; for (int r = 0; r < R; r++) { int idx = IAccessor[b][n]; - CHECK(idx >= 0 && idx < B); + TC_CHECK(idx >= 0 && idx < B); correct += LUTAccessor[idx][r]; } OAccessor[b][n] -= correct; diff --git a/test/cuda/test_tc_mapper_bugs.cc b/test/cuda/test_tc_mapper_bugs.cc index 86b764ec7..28bfbae5a 100644 --- a/test/cuda/test_tc_mapper_bugs.cc +++ b/test/cuda/test_tc_mapper_bugs.cc @@ -20,6 +20,7 @@ #include "tc/aten/aten.h" #include "tc/aten/aten_compiler.h" +#include "tc/core/check.h" #include "tc/core/cuda/cuda.h" #include "tc/core/cuda/cuda_tc_executor.h" #include "tc/core/flags.h" @@ -771,7 +772,7 @@ TEST(Convolution, NestedExpressions) { auto outputs = tc::aten::prepareOutputs(TC, convolution, inputs); tc::aten::run(*pExecutor, inputs, outputs); auto B = outputs[0]; - CHECK_EQ(at::Scalar(B[10]).toFloat(), 1); + TC_CHECK_EQ(at::Scalar(B[10]).toFloat(), 1); } // Previous versions of TC would map the reduction in the code below diff --git a/test/isl_cli_strategy.h b/test/isl_cli_strategy.h index be5b606fb..dfaeeccda 100644 --- a/test/isl_cli_strategy.h +++ b/test/isl_cli_strategy.h @@ -18,6 +18,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/cuda/cuda_mapping_options.h" #define DEFAULT_FUSION_STRATEGY "Preserve3Coincident" @@ -80,8 +81,8 @@ namespace tc { // (at a minimum: tile, mapToThreads and mapToBlocks) // 3. call makeCliStrategy with the overridden options tc::CudaMappingOptions makeBaseCliStrategy() { - tc::FusionStrategy fs; - CHECK(tc::FusionStrategy_Parse(DEFAULT_FUSION_STRATEGY, &fs)); + tc::FusionStrategy fs(FusionStrategy::Max); + TC_CHECK(tc::FusionStrategy_Parse(DEFAULT_FUSION_STRATEGY, &fs)); CudaMappingOptions options = CudaMappingOptions::makeNaiveMappingOptions() .mapToThreads(DEFAULT_BLOCK) @@ -104,11 +105,11 @@ tc::CudaMappingOptions makeBaseCliStrategy() { tc::CudaMappingOptions makeCliStrategy(tc::CudaMappingOptions options) { if (FLAGS_fusion_strategy != std::string(DEFAULT_FUSION_STRATEGY)) { - tc::FusionStrategy fs; + tc::FusionStrategy fs(FusionStrategy::Max); if (tc::FusionStrategy_Parse(FLAGS_fusion_strategy, &fs)) { options.scheduleFusionStrategy(fs); } else { - CHECK(false) << "Unknown fusion_strategy: " << FLAGS_fusion_strategy; + TC_CHECK(false) << "Unknown fusion_strategy: " << FLAGS_fusion_strategy; } } options.generic.outerScheduleOptions.proto.set_allow_skewing( diff --git a/test/test_check.cc b/test/test_check.cc new file mode 100644 index 000000000..b05247cd9 --- /dev/null +++ b/test/test_check.cc @@ -0,0 +1,209 @@ +/** + * Copyright (c) 2018-present, Facebook, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include + +#include "tc/core/check.h" + +// gtest doesn't define a macro that inspects the contents of the exception +#define ASSERT_THROW_WHAT(x, y) \ + try { \ + (x); \ + ASSERT_TRUE(false); \ + } catch (std::runtime_error & e) { \ + ASSERT_EQ(y, e.what()); \ + } + +TEST(CHECK, Plain) { + ASSERT_NO_THROW(TC_CHECK(true)); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 1 << ']'; + ASSERT_THROW_WHAT(TC_CHECK(false), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 << ']' + << ": 1+1=3"; + ASSERT_THROW_WHAT(TC_CHECK(false) << "1+1=3", expected.str()); + } +} + +TEST(CHECK, Vector) { + std::stringstream expected; + auto v = std::vector{1, 2, 3, 4}; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 << ']' + << ": 1,2,3,4"; + ASSERT_THROW_WHAT((TC_CHECK(false) << v), expected.str()); +} + +TEST(CHECK, EQ) { + ASSERT_NO_THROW(TC_CHECK_EQ(1, 1)); + ASSERT_NO_THROW(TC_CHECK_EQ(std::string("aaa"), std::string("aaa"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 << "] " + << "1 not equal to 2"; + ASSERT_THROW_WHAT(TC_CHECK_EQ(1, 2), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 1 not equal to 2: 2+2=5"; + ASSERT_THROW_WHAT(TC_CHECK_EQ(1, 2) << "2+2=5", expected.str()); + } +} + +TEST(CHECK, NE) { + ASSERT_NO_THROW(TC_CHECK_NE(1, 2)); + ASSERT_NO_THROW(TC_CHECK_NE(std::string("aaa"), std::string("baa"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 1 equal to 1"; + ASSERT_THROW_WHAT(TC_CHECK_NE(1, 1), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 1 equal to 1: 2+2=5"; + ASSERT_THROW_WHAT(TC_CHECK_NE(1, 1) << "2+2=5", expected.str()); + } +} + +TEST(CHECK, LT) { + ASSERT_NO_THROW(TC_CHECK_LT(1, 2)); + ASSERT_NO_THROW(TC_CHECK_LT(std::string("aaa"), std::string("baa"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 1 not less than 1"; + ASSERT_THROW_WHAT(TC_CHECK_LT(1, 1), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 4 not less than " << 1 << ": 4+3=8"; + ASSERT_THROW_WHAT(TC_CHECK_LT(4, 1) << "4+3=8", expected.str()); + } +} + +TEST(CHECK, GT) { + ASSERT_NO_THROW(TC_CHECK_GT(2, 1)); + ASSERT_NO_THROW(TC_CHECK_GT(std::string("ca"), std::string("baa"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 1 not greater than " << 1; + ASSERT_THROW_WHAT(TC_CHECK_GT(1, 1), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 2 not greater than 4: 3+3=7"; + ASSERT_THROW_WHAT(TC_CHECK_GT(2, 4) << "3+3=7", expected.str()); + } +} + +TEST(CHECK, LE) { + ASSERT_NO_THROW(TC_CHECK_LE(1, 2)); + ASSERT_NO_THROW(TC_CHECK_LE(1, 1)); + ASSERT_NO_THROW(TC_CHECK_LE(std::string("aaa"), std::string("baa"))); + ASSERT_NO_THROW(TC_CHECK_LE(std::string("aa"), std::string("aa"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 2 not less than or equal to 1"; + ASSERT_THROW_WHAT(TC_CHECK_LE(2, 1), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 4 not less than or equal to 1: 4+5=10"; + ASSERT_THROW_WHAT(TC_CHECK_LE(4, 1) << "4+5=10", expected.str()); + } +} + +TEST(CHECK, GE) { + ASSERT_NO_THROW(TC_CHECK_GE(2, 1)); + ASSERT_NO_THROW(TC_CHECK_GE(2, 2)); + ASSERT_NO_THROW(TC_CHECK_GE(std::string("ca"), std::string("baa"))); + ASSERT_NO_THROW(TC_CHECK_GE(std::string("ba"), std::string("ba"))); + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 7 not greater than or equal to 9"; + ASSERT_THROW_WHAT(TC_CHECK_GE(7, 9), expected.str()); + } + + { + std::stringstream expected; + expected << "Check failed [" << __FILE__ << ':' << __LINE__ + 2 + << "] 2 not greater than or equal to 6: 9+3=13"; + ASSERT_THROW_WHAT(TC_CHECK_GE(2, 6) << "9+3=13", expected.str()); + } +} + +TEST(CHECK, CustomException) { + ASSERT_THROW(TC_CHECK(false, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK(false, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK(true, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK(true, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_EQ(1, 2, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_EQ(1, 2, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_EQ(1, 1, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_EQ(1, 1, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_NE(1, 1, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_NE(1, 1, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_NE(1, 2, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_NE(1, 2, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_LT(1, 1, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_LT(1, 1, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_LT(1, 2, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_LT(1, 2, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_GT(1, 1, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_GT(1, 1, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_GT(2, 1, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_GT(2, 1, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_LE(2, 1, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_LE(2, 1, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_LE(1, 2, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_LE(1, 2, std::out_of_range) << "aa"); + + ASSERT_THROW(TC_CHECK_GE(1, 2, std::out_of_range), std::out_of_range); + ASSERT_THROW(TC_CHECK_GE(1, 2, std::out_of_range) << "aa", std::out_of_range); + ASSERT_NO_THROW(TC_CHECK_GE(2, 1, std::out_of_range)); + ASSERT_NO_THROW(TC_CHECK_GE(2, 1, std::out_of_range) << "aa"); +} + +int main(int argc, char** argv) { + ::testing::InitGoogleTest(&argc, argv); + return RUN_ALL_TESTS(); +} diff --git a/test/test_cuda_mapper_memory_promotion.cc b/test/test_cuda_mapper_memory_promotion.cc index a688295c8..6f2ed85cc 100644 --- a/test/test_cuda_mapper_memory_promotion.cc +++ b/test/test_cuda_mapper_memory_promotion.cc @@ -17,6 +17,7 @@ #include #include +#include "tc/core/check.h" #include "tc/core/polyhedral/cuda/codegen.h" #include "tc/core/polyhedral/cuda/mapped_scop.h" #include "tc/core/polyhedral/cuda/memory_promotion_heuristic.h" @@ -299,7 +300,7 @@ def fun(float(N, M) A, float(N, M) B) -> (C) { auto schedule = partialSchedule( scop.scheduleRoot(), scop.scheduleRoot()->child(childPos)); auto scopedAccess = oneGroup->originalAccesses().apply_domain(schedule); - CHECK(scopedAccess.is_equal(oneGroup->scopedAccesses())) + TC_CHECK(scopedAccess.is_equal(oneGroup->scopedAccesses())) << "expected original accesses " << oneGroup->originalAccesses() << " to be equal to scoped accesses " << oneGroup->scopedAccesses() << " after applying the partial schedule " << schedule; @@ -366,7 +367,7 @@ def fun(float(N, M) A) -> (B, C) { auto schedule = partialSchedule(scop.scheduleRoot(), t); auto scopedAccess = groupsB[0]->originalAccesses().apply_domain(schedule); - CHECK(scopedAccess.is_equal(groupsB[0]->scopedAccesses())) + TC_CHECK(scopedAccess.is_equal(groupsB[0]->scopedAccesses())) << "expected original accesses " << groupsB[0]->originalAccesses() << " to be equal to scoped accesses " << groupsB[0]->scopedAccesses() << " after applying the partial schedule " << schedule; diff --git a/test/test_mapper_llvm.cc b/test/test_mapper_llvm.cc index 1eac81fb8..8a051b426 100644 --- a/test/test_mapper_llvm.cc +++ b/test/test_mapper_llvm.cc @@ -19,6 +19,7 @@ #include #include "tc/aten/aten.h" +#include "tc/core/check.h" #include "tc/core/cpu/cpu_mapping_options.h" #include "tc/core/cpu/cpu_tc_executor.h" #include "tc/core/polyhedral/codegen_llvm.h" @@ -208,7 +209,7 @@ def convolution(float(N,C,H,W) I, float(O,C,KH,KW) W1, float(O) B) -> (tmp, O1) B.data(), tmp.data(), output.data()); - CHECK_EQ(output.ndimension(), 4); + TC_CHECK_EQ(output.ndimension(), 4); checkRtol(output - expected, {I, W1, B}, C * KH * KW, 1e-6); } diff --git a/test/test_tc_mapper_harness-inl.h b/test/test_tc_mapper_harness-inl.h index 231e95c68..a4a1c6e3c 100644 --- a/test/test_tc_mapper_harness-inl.h +++ b/test/test_tc_mapper_harness-inl.h @@ -20,6 +20,7 @@ #include "tc/aten/aten.h" #include "tc/aten/aten_compiler.h" +#include "tc/core/check.h" #include "tc/core/scope_guard.h" #include "tc/core/tensor.h" #include "tc/lang/canonicalize.h" @@ -91,7 +92,8 @@ def sum1D(float(M) A) -> (C) { } )TC"}; - CHECK_LE(version, 3u) << "Versions [0-3] supported, asked for: " << version; + TC_CHECK_LE(version, 3u) + << "Versions [0-3] supported, asked for: " << version; auto refOutput = A.sum(); auto checkFun = [&, refOutput]( const std::vector& inputs,