[C++ API] Step 1: drop manual and compilation caches

This is the first PR towards implementing the API proposed in #307.
As part of refactoring the OptionsCache we realized that compilation caches may
have been quite useful in the beginning but are a premature optimization in the
new API.

Removing them has multiple benefits such as:
1. simplify the codebase CRTP + inheritance + custom traversals and
serialization are not fun to maintain
2. simplify porting to the new API where ExecutionEngine is dead
3. caches seem to add overhead in the context of cooperative groups
4. in the new API we can cache executors which removes the need to rely on
compilation caches
5. tuning and options caches is the expensive bit that needs to be memoized

This seems corroborated by the [discussion](#339 (comment)).

Note that some tests in test_tc_mapper need to be temporarily degraded because
we do not have a simple way to extract generated code before the new API.

Author: nicolasvasilache

docs/source/framework/pytorch_integration/writing_layers.rst

@@ -219,48 +219,6 @@ adopt whatever feels more convenient.
     inp = torch.ones(1, 1, 4, 4).cuda()
     out = avgpool(inp)
 
-
-Manually injecting external CUDA code
--------------------------------------
-
-If you have an external efficient CUDA code that you want to use rather than
-the CUDA code that TC generates, you can inject your code easily. For this,
-you need to create a string which has the CUDA code you want to inject and you
-need to pass the name of the kernel and the CUDA code string to the :code:`tc.define`
-call. For example:
-
-.. code-block:: python
-
-    import tensor_comprehensions as tc
-    import torch
-    lang = """
-    def add(float(N) A, float(N) B) -> (output) {
-        output(n) = A(n) + B(n)
-    }
-    """
-
-    cuda_code = """
-    extern "C"{
-    __global__ void my_add(float* __restrict__ output, const float* __restrict__ A, const float* __restrict B)
-    {
-        int t = threadIdx.x;
-        output[t] = A[t] + B[t];
-    }
-    }
-    """
-
-    add = tc.define(lang, name="add", inject_kernel="my_add", cuda_code=cuda_code)
-    a, b = torch.randn(100).cuda(), torch.randn(100).cuda()
-    out = add(a, b, grid=[1, 1, 1], block=[100, 1, 1])
-
-.. note::
-
-    In such cases, please note that TC doesn't modify the injected CUDA kernel. It will
-    simply run the kernel injected as is and TC will also not guarantee the performance
-    of the kernel. User needs to specify the :code:`grid` and :code:`block` values
-    when running the layer and TC will simply use those settings.
-
-
 Built-in Functions
 ------------------
 

tc/autotuner/genetic_autotuner.cc

@@ -43,10 +43,8 @@ namespace {
 
 void enableOrLoadCache(const std::string& filename) {
   tc::OptionsCache::enableCache();
-  tc::CudaCache::enableCache();
   if (!filename.empty()) {
     tc::OptionsCache::loadCacheFromProtobuf(tc::makeOptionsFilename(filename));
-    tc::CudaCache::loadCacheFromProtobuf(tc::makeCudaFilename(filename));
   }
 }
 } // namespace
@@ -62,9 +60,6 @@ void GeneticAutotuner::storeCaches(const std::string& filename) {
     tc::OptionsCache::dumpCacheToProtobuf(tc::makeOptionsFilename(filename));
 
     tc::OptionsCache::getCache()->keepOnlyBestCandidates(1);
-    tc::removeFromCudaCacheEntriesNotInOptionsCache(
-        *tc::CudaCache::getCache(), *tc::OptionsCache::getCache());
-    tc::CudaCache::dumpCacheToProtobuf(tc::makeCudaFilename(filename));
   }
 }
 

tc/autotuner/genetic_autotuner_aten.cc

@@ -21,7 +21,6 @@
 #include <thread>
 
 #include "tc/core/cuda/cuda.h"
-#include "tc/core/cuda/cuda_compilation_cache.h"
 #include "tc/core/cuda/cuda_tc_executor.h"
 #include "tc/core/flags.h"
 #include "tc/core/scope_guard.h"

tc/autotuner/genetic_tuning_harness.cc

@@ -27,7 +27,6 @@
 #include "tc/autotuner/utils/printer.h"
 #include "tc/autotuner/utils/utils.h"
 #include "tc/core/cuda/cuda.h"
-#include "tc/core/cuda/cuda_compilation_cache.h"
 #include "tc/core/cuda/cuda_mapping_options_cpp_printer.h"
 #include "tc/core/cuda/cuda_tc_executor.h"
 #include "tc/core/execution_engine.h"

tc/benchmarks/benchmark_fixture.h

@@ -260,8 +260,6 @@ struct Benchmark : public ::testing::Test {
 
     tc::OptionsCache::enableCache();
     tc::OptionsCache::loadCacheFromProtobuf(cacheFilename + ".options");
-    tc::CudaCache::enableCache();
-    tc::CudaCache::loadCacheFromProtobuf(tc::makeCudaFilename(cacheFilename));
     tc::FLAGS_tuner_gen_restore_number = 1;
 
     tc::ATenCompilationUnit<tc::CudaTcExecutor> atCompl;

tc/core/cuda/cuda_compilation_cache.cc

@@ -141,152 +141,6 @@ CachedEntryType* searchKernel(
 }
 } // namespace
 
-////////////////////////////////////////////////////////////////////////////////
-// CudaCache
-////////////////////////////////////////////////////////////////////////////////
-std::shared_ptr<CudaCache>& CudaCache::getGlobalSharedCache() {
-  static std::shared_ptr<CudaCache> cudaCache_;
-  return cudaCache_;
-}
-
-CudaCachedEntry::CudaCachedEntry(
-    const std::string& id,
-    const std::string& kernelSpecializedName,
-    const std::vector<int>& kernelParameters,
-    const Grid& grid,
-    const Block& block,
-    const CudaMappingOptions& mappingOptions,
-    const std::vector<const DLTensor*>& inputs,
-    const std::vector<const DLTensor*>& outputs,
-    const std::string& cudaSource,
-    const std::string& deviceStr)
-    : key{id,
-          mappingOptions,
-          DLTensorToTensorInfoVector(inputs),
-          DLTensorToTensorInfoVector(outputs),
-          deviceStr,
-          git_version},
-      values{cudaSource, kernelSpecializedName, kernelParameters, grid, block} {
-}
-
-CudaCachedEntry::CudaCachedEntry(const CudaCacheEntryProto& buf)
-    : key{buf.id(),
-          CudaMappingOptions{buf.kernel_options()},
-          ProtoToTensorInfoVector(buf.inputs()),
-          ProtoToTensorInfoVector(buf.outputs()),
-          buf.device_str(),
-          buf.git_version()},
-      values{buf.cuda_source(),
-             buf.specialized_name(),
-             std::vector<int>{buf.parameters().begin(), buf.parameters().end()},
-             Grid(buf.grid_dims()),
-             Block(buf.block_dims())} {}
-
-CudaCache::CudaCache(const CudaCacheProto& buf) {
-  entries_.reserve(buf.entries_size());
-  for (const auto& entry_buf : buf.entries())
-    entries_.emplace_back(entry_buf);
-}
-
-void CudaCache::cacheKernel(CudaCachedEntry&& entry) {
-  std::lock_guard<std::mutex> lock(mtx_);
-  ++numberCacheAttemps;
-  auto retrievedEntry = searchKernel(
-      entries_,
-      entry.key.id,
-      entry.key.mappingOptions,
-      entry.key.inputs,
-      entry.key.outputs);
-  if (retrievedEntry) {
-    if (retrievedEntry->values.cudaSource != entry.values.cudaSource or
-        retrievedEntry->values.grid != entry.values.grid or
-        retrievedEntry->values.block != entry.values.block) {
-      throw CacheEntrySameKeyDifferentValue(
-          "CudaCache::CacheKernel: a kernel matching the id, options and "
-          "inputs was previously cached with different cuda source or block "
-          "or grid dimensions.");
-    }
-    return;
-  }
-  entries_.emplace_back(std::move(entry));
-}
-
-std::unique_ptr<CudaCacheRetrievalResult> CudaCache::retrieveKernel(
-    const std::string& id,
-    const CudaMappingOptions& options,
-    const std::vector<const DLTensor*>& inputs,
-    const std::vector<const DLTensor*>& outputs) const {
-  std::lock_guard<std::mutex> lock(mtx_);
-  ++numberAttemptedRetrievals;
-  auto entry = searchKernel(entries_, id, options, inputs, outputs);
-  if (not entry) {
-    return nullptr;
-  }
-  ++numberSuccessfulRetrievals;
-  return std::unique_ptr<CudaCacheRetrievalResult>(
-      new CudaCacheRetrievalResult{entry->values.cudaSource,
-                                   entry->values.kernelSpecializedName,
-                                   entry->values.kernelParameters,
-                                   entry->values.grid,
-                                   entry->values.block});
-}
-
-void CudaCache::removeEntriesNotInOptionsCache(const OptionsCache& oc) {
-  std::vector<CudaCachedEntry> newEntries;
-  for (const auto& entry : oc) {
-    for (const auto& options : entry.values) {
-      auto cudaEntry = searchKernel(
-          entries_,
-          entry.key.id,
-          options.mappingOptions,
-          entry.key.inputs,
-          entry.key.outputs);
-      if (cudaEntry) {
-        newEntries.push_back(std::move(*cudaEntry));
-      }
-    }
-  }
-  entries_ = std::move(newEntries);
-}
-
-CudaCacheProto CudaCache::toProtobuf() const {
-  CudaCacheProto buf;
-  auto* entriesBuf = buf.mutable_entries();
-  entriesBuf->Reserve(entries_.size());
-  std::transform(
-      entries_.begin(),
-      entries_.end(),
-      google::protobuf::RepeatedPtrFieldBackInserter(entriesBuf),
-      [](const CudaCachedEntry& entry) { return entry.toProtobuf(); });
-  return buf;
-}
-
-CudaCacheEntryProto CudaCachedEntry::toProtobuf() const {
-  CudaCacheEntryProto buf;
-  buf.set_id(key.id);
-  *buf.mutable_kernel_options() = key.mappingOptions.proto();
-  std::transform(
-      key.inputs.begin(),
-      key.inputs.end(),
-      google::protobuf::RepeatedPtrFieldBackInserter(buf.mutable_inputs()),
-      [](const detail::TensorInfo& input) { return input.toProtobuf(); });
-  std::transform(
-      key.outputs.begin(),
-      key.outputs.end(),
-      google::protobuf::RepeatedPtrFieldBackInserter(buf.mutable_outputs()),
-      [](const detail::TensorInfo& output) { return output.toProtobuf(); });
-  buf.set_device_str(key.deviceStr);
-  buf.set_git_version(key.gitVersion);
-
-  buf.set_cuda_source(values.cudaSource);
-  *buf.mutable_grid_dims() = values.grid.view.proto;
-  *buf.mutable_block_dims() = values.block.view.proto;
-  buf.set_specialized_name(values.kernelSpecializedName);
-  WriteProtobufArray(values.kernelParameters, buf.mutable_parameters());
-
-  return buf;
-}
-
 ////////////////////////////////////////////////////////////////////////////////
 // OptionsCache
 ////////////////////////////////////////////////////////////////////////////////
@@ -566,66 +420,4 @@ void OptionsCache::keepOnlyBestCandidates(size_t numberToKeep) {
     }
   }
 }
-
-////////////////////////////////////////////////////////////////////////////////
-// ManualCudaCache
-////////////////////////////////////////////////////////////////////////////////
-std::shared_ptr<ManualCudaCache>& ManualCudaCache::getGlobalSharedCache() {
-  static std::shared_ptr<ManualCudaCache> manualCudaCache_;
-  return manualCudaCache_;
-}
-
-ManualCudaCachedEntry::ManualCudaCachedEntry(
-    const std::string& id,
-    const std::string& kernelSpecializedName,
-    const std::vector<int>& kernelParameters,
-    const Grid& grid,
-    const Block& block,
-    const std::vector<const DLTensor*>& inputs,
-    const std::vector<const DLTensor*>& outputs,
-    const std::string& cudaSource,
-    const std::string& deviceStr)
-    : key{id,
-          DLTensorToTensorInfoVector(inputs),
-          DLTensorToTensorInfoVector(outputs),
-          deviceStr,
-          git_version},
-      values{cudaSource, kernelSpecializedName, kernelParameters, grid, block} {
-}
-
-void ManualCudaCache::cacheKernel(ManualCudaCachedEntry&& entry) {
-  std::lock_guard<std::mutex> lock(mtx_);
-  ++numberCacheAttemps;
-  auto retrievedEntry =
-      searchKernel(entries_, entry.key.id, entry.key.inputs, entry.key.outputs);
-  if (retrievedEntry) {
-    retrievedEntry->values.grid = entry.values.grid;
-    retrievedEntry->values.block = entry.values.block;
-    retrievedEntry->values.cudaSource = entry.values.cudaSource;
-    retrievedEntry->values.kernelSpecializedName =
-        entry.values.kernelSpecializedName;
-    retrievedEntry->values.kernelParameters = entry.values.kernelParameters;
-    return;
-  }
-  entries_.emplace_back(std::move(entry));
-}
-
-std::unique_ptr<ManualCudaCacheRetrievalResult> ManualCudaCache::retrieveKernel(
-    const std::string& id,
-    const std::vector<const DLTensor*>& inputs,
-    const std::vector<const DLTensor*>& outputs) const {
-  std::lock_guard<std::mutex> lock(mtx_);
-  ++numberAttemptedRetrievals;
-  auto entry = searchKernel(entries_, id, inputs, outputs);
-  if (not entry) {
-    return nullptr;
-  }
-  ++numberSuccessfulRetrievals;
-  return std::unique_ptr<ManualCudaCacheRetrievalResult>(
-      new ManualCudaCacheRetrievalResult{entry->values.cudaSource,
-                                         entry->values.kernelSpecializedName,
-                                         entry->values.kernelParameters,
-                                         entry->values.grid,
-                                         entry->values.block});
-}
 } // namespace tc