diff --git a/tc/aten/aten_compiler-inl.h b/tc/aten/aten_compiler-inl.h index b33854dc6..b67d75538 100644 --- a/tc/aten/aten_compiler-inl.h +++ b/tc/aten/aten_compiler-inl.h @@ -111,6 +111,26 @@ Duration ATenCompilationUnit::run( handle, inputDLTensorsPair.first, outputDLTensorsPair.first, profile); } +template +typename ExecutorType::ProfilingInfoType +ATenCompilationUnit::profile( + const std::string& name, + const std::vector& inputs, + std::vector& outputs, + size_t handle) { + at::Backend backend = inputs[0].type().backend(); + auto inputDLTensorsPair = toConstDlpackTensors(inputs); + ScopeGuard g1([&]() { deleteDlmTensors(inputDLTensorsPair.second); }); + auto outTensorInfo = + executionEngine_->inferOutputTensorInfo(name, inputDLTensorsPair.first); + prepareOutputs( + executionEngine_->treeForFunction(name), outTensorInfo, backend, outputs); + auto outputDLTensorsPair = toDlpackTensors(outputs); + ScopeGuard g2([&]() { deleteDlmTensors(outputDLTensorsPair.second); }); + return executionEngine_->profile( + handle, inputDLTensorsPair.first, outputDLTensorsPair.first); +} + template void ATenCompilationUnit::uncheckedRun( const std::vector& inputs, diff --git a/tc/aten/aten_compiler.h b/tc/aten/aten_compiler.h index 9a24ea2af..013d76141 100644 --- a/tc/aten/aten_compiler.h +++ b/tc/aten/aten_compiler.h @@ -62,6 +62,12 @@ class ATenCompilationUnit { size_t handle, bool profile = false); + typename ExecutorType::ProfilingInfoType profile( + const std::string& name, + const std::vector& inputs, + std::vector& outputs, + size_t handle); + /// This is the "low-latency" mode in which we just propagate ATen tensors /// Sizes are not checked and it is the user's responsibility to ensure that /// they match. If the user doesn't then segfault will likely occur. diff --git a/tc/autotuner/genetic_autotuner.cc b/tc/autotuner/genetic_autotuner.cc index 332cff21f..92b9588d5 100644 --- a/tc/autotuner/genetic_autotuner.cc +++ b/tc/autotuner/genetic_autotuner.cc @@ -57,8 +57,10 @@ void GeneticAutotuner::storeCaches(const std::string& filename) { } else { std::cout << "Dumping cache to " << filename << ".cuda/options" << std::endl; - tc::OptionsCache::getCache()->keepOnlyBestCandidates( - tc::FLAGS_tuner_save_best_candidates_count); + if (not FLAGS_tuner_gen_profiled_run) { + tc::OptionsCache::getCache()->keepOnlyBestCandidates( + tc::FLAGS_tuner_save_best_candidates_count); + } tc::OptionsCache::dumpCacheToProtobuf(tc::makeOptionsFilename(filename)); tc::OptionsCache::getCache()->keepOnlyBestCandidates(1); diff --git a/tc/autotuner/genetic_tuning_harness.cc b/tc/autotuner/genetic_tuning_harness.cc index f5196635a..2ab79cd47 100644 --- a/tc/autotuner/genetic_tuning_harness.cc +++ b/tc/autotuner/genetic_tuning_harness.cc @@ -408,7 +408,11 @@ void GeneticTunerHarness::doGpuWork( } else { runtimes.reserve(kReducedBenchmarkIterations); for (size_t i = 0; i < kReducedBenchmarkIterations; ++i) { - runtimes.push_back(engine.run(handle, inputs, outputs, true)); + if (FLAGS_tuner_gen_profiled_run) { + runtimes.push_back(engine.profile(handle, inputs, outputs).runtime); + } else { + runtimes.push_back(engine.run(handle, inputs, outputs, true)); + } } engine.clear(handle); } diff --git a/tc/core/CMakeLists.txt b/tc/core/CMakeLists.txt index 19eed9577..248e871ce 100644 --- a/tc/core/CMakeLists.txt +++ b/tc/core/CMakeLists.txt @@ -157,8 +157,9 @@ if (WITH_CUDA) cuda/cuda_compilation_cache.cc cuda/cuda_rtc.cc cuda/cuda_tc_executor.cc + cuda/cuda_profile.cc ) - target_include_directories(tc_cuda PUBLIC ${LLVM_INCLUDE_DIRS}) + target_include_directories(tc_cuda PUBLIC ${LLVM_INCLUDE_DIRS} ${CUDA_TOOLKIT_ROOT_DIR}/extras/CUPTI/include) target_link_libraries( tc_cuda @@ -166,6 +167,7 @@ if (WITH_CUDA) ${CUDA_curand_LIBRARY} ${CUDA_LIBRARIES} ${CUDA_NVRTC_LIBRARIES} + ${CUDA_cupti_LIBRARY} ${ISL_LIBRARIES} tc_lang diff --git a/tc/core/cpu/cpu_tc_executor.h b/tc/core/cpu/cpu_tc_executor.h index b74f58e51..e17aed80f 100644 --- a/tc/core/cpu/cpu_tc_executor.h +++ b/tc/core/cpu/cpu_tc_executor.h @@ -33,10 +33,12 @@ struct CpuRTCFunction { void clear() {} }; +struct CpuProfilingInfo {}; + class CpuTcExecutor : public ::tc::TcExecutor { public: using MappingOptionsType = CpuMappingOptions; - + using ProfilingInfoType = CpuProfilingInfo; CpuTcExecutor( std::string id, const std::vector& inputsInfo, diff --git a/tc/core/cuda/cuda.h b/tc/core/cuda/cuda.h index 2168df1c0..06a3fdab4 100644 --- a/tc/core/cuda/cuda.h +++ b/tc/core/cuda/cuda.h @@ -57,6 +57,19 @@ } \ } while (0) +#define TC_CUPTI_CHECK(condition) \ + do { \ + CUptiResult result = condition; \ + if (result != CUPTI_SUCCESS) { \ + const char* msg; \ + cuptiGetResultString(result, &msg); \ + std::stringstream ss; \ + ss << "Error at: " << __FILE__ << ":" << __LINE__ << ": " << msg; \ + LOG(WARNING) << ss.str(); \ + throw std::runtime_error(ss.str().c_str()); \ + } \ + } while (0) + #define TC_CUDA_RUNTIMEAPI_ENFORCE(condition) \ do { \ cudaError_t result = condition; \ diff --git a/tc/core/cuda/cuda_compilation_cache.cc b/tc/core/cuda/cuda_compilation_cache.cc index 7123aba2d..e0ac6cb1c 100644 --- a/tc/core/cuda/cuda_compilation_cache.cc +++ b/tc/core/cuda/cuda_compilation_cache.cc @@ -302,6 +302,29 @@ OptionsCachedEntry::OptionsCachedEntry( values.emplace_back(options, runtime); } +OptionsCachedEntry::OptionsCachedEntry( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs, + const std::string& deviceStr, + const CudaMappingOptions& options, + const CudaProfilingInfo& pInfo) + : key(id, inputs, outputs, deviceStr, git_version) { + values.emplace_back(options, pInfo); +} + +OptionsCachedEntry::OptionsCachedEntry( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs, + const std::string& deviceStr, + const CudaMappingOptions& options, + Duration runtime, + const CudaProfilingInfo& pInfo) + : key(id, inputs, outputs, deviceStr, git_version) { + values.emplace_back(options, runtime, pInfo); +} + OptionsCachedEntry::Key::Key( const std::string& id, const std::vector& inputs_, @@ -333,8 +356,37 @@ OptionsCachedEntry::Values::Values( OptionsCachedEntry::Values::Values( const CudaMappingOptions& options, - std::vector&& runtimes) - : mappingOptions(options), recordedRuntimes(std::move(runtimes)) {} + const CudaProfilingInfo& pInfo) + : mappingOptions(options), profiles{pInfo} {} + +OptionsCachedEntry::Values::Values( + const CudaMappingOptions& options, + Duration runtime, + const CudaProfilingInfo& pInfo) + : mappingOptions(options), recordedRuntimes{runtime}, profiles{pInfo} {} + +OptionsCachedEntry::Values::Values( + const CudaMappingOptions& options, + std::vector&& runtimes, + std::vector&& pInfos) + : mappingOptions(options), + recordedRuntimes(std::move(runtimes)), + profiles(std::move(pInfos)) {} + +namespace { +tc::CudaProfilingInfo fromProto(const tc::CudaProfilingProto& buf) { + tc::CudaProfilingInfo pInfo; + pInfo.runtime = std::chrono::microseconds(buf.runtime()); + pInfo.ipc = buf.ipc(); + pInfo.globalLoadEfficiency = buf.globalloadefficiency(); + pInfo.globalStoreEfficiency = buf.globalstoreefficiency(); + pInfo.sharedMemoryEfficiency = buf.sharedmemoryefficiency(); + pInfo.localMemoryOverhead = buf.localmemoryoverhead(); + pInfo.achievedOccupancy = buf.achievedoccupancy(); + pInfo.warpExecutionEfficiency = buf.warpexecutionefficiency(); + return pInfo; +} +} // namespace OptionsCachedEntry::OptionsCachedEntry(const OptionsCacheEntryProto& buf) : key(buf.id(), @@ -348,7 +400,7 @@ OptionsCachedEntry::OptionsCachedEntry(const OptionsCacheEntryProto& buf) } for (const auto& value : buf.values()) { - if (value.recorded_runtimes_size() == 0) { + if (value.recorded_runtimes_size() == 0 and value.profiles_size() == 0) { throw std::invalid_argument( "OptionsCachedEntry invalid protobuf: each entry value should have at least one recorded runtime."); } @@ -359,11 +411,39 @@ OptionsCachedEntry::OptionsCachedEntry(const OptionsCacheEntryProto& buf) value.recorded_runtimes().end(), std::back_inserter(runtimes), [](int64_t us) { return std::chrono::microseconds(us); }); + std::vector profiles; + profiles.reserve(value.profiles_size()); + std::transform( + value.profiles().begin(), + value.profiles().end(), + std::back_inserter(profiles), + [](const CudaProfilingProto& buf) { return fromProto(buf); }); + values.emplace_back( - CudaMappingOptions(value.kernel_options()), std::move(runtimes)); + CudaMappingOptions(value.kernel_options()), + std::move(runtimes), + std::move(profiles)); } } +namespace { +tc::CudaProfilingProto toProto(const tc::CudaProfilingInfo& pInfo) { + tc::CudaProfilingProto buf; + buf.set_runtime( + std::chrono::duration_cast(pInfo.runtime) + .count()); + buf.set_ipc(pInfo.ipc); + buf.set_globalloadefficiency(pInfo.globalLoadEfficiency); + buf.set_globalstoreefficiency(pInfo.globalStoreEfficiency); + buf.set_sharedmemoryefficiency(pInfo.sharedMemoryEfficiency); + buf.set_localmemoryoverhead(pInfo.localMemoryOverhead); + buf.set_achievedoccupancy(pInfo.achievedOccupancy); + buf.set_warpexecutionefficiency(pInfo.warpExecutionEfficiency); + + return buf; +} +} // namespace + OptionsCacheEntryProto OptionsCachedEntry::toProtobuf() const { OptionsCacheEntryProto buf; buf.set_id(key.id); @@ -392,6 +472,9 @@ OptionsCacheEntryProto OptionsCachedEntry::toProtobuf() const { buf.add_recorded_runtimes( std::chrono::duration_cast(r).count()); } + for (const auto& p : v.profiles) { + *buf.add_profiles() = toProto(p); + } return buf; }); return buf; @@ -460,6 +543,69 @@ void OptionsCache::recordRuntime( v->recordedRuntimes.push_back(runtime); } +void OptionsCache::recordProfilingInfo( + const std::string& id, + const CudaMappingOptions& options, + const std::vector& inputs, + const std::vector& outputs, + const CudaProfilingInfo pInfo) { + std::lock_guard lock(mtx_); + ++numberCacheAttemps; + auto gpuStr = CudaGPUInfo::GPUInfo().GetCudaDeviceStr(); + + auto kernel = searchKernel(entries_, id, inputs, outputs); + if (not kernel) { + entries_.emplace_back( + id, inputs, outputs, gpuStr, options, pInfo.runtime, pInfo); + return; + } + auto v = std::find_if( + kernel->values.begin(), + kernel->values.end(), + [&options](const CachedEntry::Values& v) { + return v.mappingOptions == options; + }); + if (v == kernel->values.end()) { + kernel->values.emplace_back(options, pInfo); + return; + } + + v->recordedRuntimes.push_back(pInfo.runtime); + v->profiles.push_back(pInfo); +} + +std::vector +OptionsCache::retrieveOptionsAndProfilingInfo( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs) const { + std::lock_guard lock(mtx_); + ++numberAttemptedRetrievals; + auto ret = searchKernel(entries_, id, inputs, outputs); + if (not ret) { + return {}; + } + ++numberSuccessfulRetrievals; + std::vector res; + res.reserve(ret->values.size()); + std::transform( + ret->values.begin(), + ret->values.end(), + std::back_inserter(res), + [](const CachedEntry::Values& v) -> OptionsCacheRetrievalResult { + return {v.mappingOptions, v.recordedRuntimes, v.profiles}; + }); + res.erase( + std::remove_if( + res.begin(), + res.end(), + [](const OptionsCacheRetrievalResult& rr) { + return rr.profilingInfo.empty(); + }), + res.end()); + return res; +} + std::vector OptionsCache::retrieveOptionsAndRuntimes( const std::string& id, diff --git a/tc/core/cuda/cuda_compilation_cache.h b/tc/core/cuda/cuda_compilation_cache.h index f2e949e51..2f904b467 100644 --- a/tc/core/cuda/cuda_compilation_cache.h +++ b/tc/core/cuda/cuda_compilation_cache.h @@ -53,6 +53,22 @@ struct OptionsCachedEntry { const std::string& deviceStr, const CudaMappingOptions& options, Duration runtime); + OptionsCachedEntry( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs, + const std::string& deviceStr, + const CudaMappingOptions& options, + const CudaProfilingInfo& pInfo); + OptionsCachedEntry( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs, + const std::string& deviceStr, + const CudaMappingOptions& options, + Duration runtime, + const CudaProfilingInfo& pInfo); + OptionsCachedEntry(const OptionsCacheEntryProto& buf); OptionsCacheEntryProto toProtobuf() const; @@ -78,9 +94,19 @@ struct OptionsCachedEntry { struct Values { Values(const CudaMappingOptions& options, Duration runtime); + Values(const CudaMappingOptions& options, const CudaProfilingInfo& pInfo); + Values( + const CudaMappingOptions& options, + Duration runtime, + const CudaProfilingInfo& pInfo); Values(const CudaMappingOptions& options, std::vector&& runtimes); + Values( + const CudaMappingOptions& options, + std::vector&& runtimes, + std::vector&& pInfos); CudaMappingOptions mappingOptions; std::vector recordedRuntimes; + std::vector profiles; }; Key key; std::vector values; @@ -89,6 +115,7 @@ struct OptionsCachedEntry { struct OptionsCacheRetrievalResult { CudaMappingOptions options; std::vector recordedRuntimes; + std::vector profilingInfo; }; class OptionsCache : public Cache { @@ -114,6 +141,18 @@ class OptionsCache : public Cache { const std::vector& outputs, Duration runtime); + void recordProfilingInfo( + const std::string& id, + const CudaMappingOptions& options, + const std::vector& inputs, + const std::vector& outputs, + CudaProfilingInfo pIfno); + + std::vector retrieveOptionsAndProfilingInfo( + const std::string& id, + const std::vector& inputs, + const std::vector& outputs) const; + std::vector retrieveOptionsAndRuntimes( const std::string& id, const std::vector& inputs, diff --git a/tc/core/cuda/cuda_profile.cc b/tc/core/cuda/cuda_profile.cc new file mode 100644 index 000000000..9e8ce58b9 --- /dev/null +++ b/tc/core/cuda/cuda_profile.cc @@ -0,0 +1,509 @@ +/** + * 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 "tc/core/cuda/cuda_profile.h" + +#include +#include +#include +#include + +#include "tc/core/cuda/cuda.h" +#include "tc/core/flags.h" +#include "tc/core/scope_guard.h" + +namespace tc { + +class CudaDeviceLockableRuntimes { + public: + CudaDeviceLockableRuntimes() + : runtimes_(CudaGPUInfo::GPUInfo().NumberGPUs()), + mutexes_(CudaGPUInfo::GPUInfo().NumberGPUs()) {} + + std::mutex& getDeviceMutex(uint32_t deviceID) { + return mutexes_.at(deviceID); + } + + Duration getRuntime(uint32_t deviceID) { + return runtimes_.at(deviceID); + } + + void setRuntime(const Duration& d, uint32_t deviceID) { + runtimes_.at(deviceID) = d; + } + + private: + std::vector runtimes_; + std::vector mutexes_; +}; + +namespace { +CudaDeviceLockableRuntimes& runtimes() { + static CudaDeviceLockableRuntimes runtimes; + return runtimes; +} +} // namespace + +bool operator==(const CudaProfilingInfo& a, const CudaProfilingInfo& b) { + return std::tie( + a.runtime, + a.ipc, + a.globalLoadEfficiency, + a.globalStoreEfficiency, + a.sharedMemoryEfficiency, + a.localMemoryOverhead, + a.achievedOccupancy, + a.warpExecutionEfficiency) == + std::tie( + b.runtime, + b.ipc, + b.globalLoadEfficiency, + b.globalStoreEfficiency, + b.sharedMemoryEfficiency, + b.localMemoryOverhead, + b.achievedOccupancy, + b.warpExecutionEfficiency); +} + +CudaCuptiProfiler::CudaCuptiProfiler(KernelType kernel, CUdevice device) + : kernel_{std::move(kernel)}, + device_{device}, + metrics{{"ipc", device_}, + {"gld_efficiency", device_}, + {"gst_efficiency", device_}, + {"achieved_occupancy", device_}, + {"shared_efficiency", device_}, + {"warp_execution_efficiency", device_}, + {"local_memory_overhead", device_}} {} + +namespace { +void CUPTIAPI +bufferRequested(uint8_t** buffer, size_t* size, size_t* maxNumRecords) { + *size = 16 * 1024; + *buffer = static_cast(aligned_alloc(8, *size)); + *maxNumRecords = 0; + CHECK(*buffer != nullptr) << "Could not allocate memory."; +} + +void CUPTIAPI bufferCompleted( + CUcontext ctx, + uint32_t streamId, + uint8_t* buffer, + size_t size, + size_t validSize) { + CUpti_Activity* record = nullptr; + // since we launched only 1 kernel, we should have only 1 kernel record + TC_CUPTI_CHECK(cuptiActivityGetNextRecord(buffer, validSize, &record)); + +#if (CUPTI_API_VERSION >= 10) + using ActivityKernelType = CUpti_ActivityKernel4; +#else + using ActivityKernelType = CUpti_ActivityKernel3; +#endif + + auto kernel = reinterpret_cast(record); + CHECK_EQ(kernel->kind, CUPTI_ACTIVITY_KIND_KERNEL) + << "Expected kernel activity record, got " << kernel->kind; + + auto kernelDuration = kernel->end - kernel->start; + auto device = kernel->deviceId; + runtimes().setRuntime(std::chrono::nanoseconds(kernelDuration), device); + free(buffer); +} + +struct MetricData { + CUdevice device; + CUpti_EventGroupSet* eventGroups; + uint32_t numEvents; + std::vector eventIds; + std::vector eventValues; +}; + +std::pair logEvent( + CUpti_EventID eventId, + uint64_t numInstances, + uint64_t numTotalInstances, + uint64_t sum, + uint64_t normalized, + uint64_t* values) { + constexpr size_t keventNameArraySize = 128; + char eventName[keventNameArraySize]; + size_t eventNameSize = keventNameArraySize - 1; + CHECK_GT(eventNameSize, 0); + TC_CUPTI_CHECK(cuptiEventGetAttribute( + eventId, CUPTI_EVENT_ATTR_NAME, &eventNameSize, eventName)); + eventName[eventNameSize] = '\0'; + + std::stringstream ss; + ss << eventName << " = " << sum << " ("; + if (numInstances > 1) { + for (uint64_t k = 0; k < numInstances; k++) { + if (k != 0) { + ss << ", "; + } + ss << values[k]; + } + } + ss << ')'; + auto line1 = ss.str(); + LOG(INFO) << ss.str(); + + ss.str(""); + ss.clear(); + + ss << eventName << " (normalized) (" << sum << " * " << numTotalInstances + << ") / " << numInstances << " = " << normalized; + auto line2 = ss.str(); + return std::make_pair(std::move(line1), std::move(line2)); +} + +void eventCollectionStart( + MetricData& metricData, + const CUpti_CallbackData* cbInfo) { + cudaDeviceSynchronize(); + TC_CUPTI_CHECK(cuptiSetEventCollectionMode( + cbInfo->context, CUPTI_EVENT_COLLECTION_MODE_KERNEL)); + for (uint64_t i = 0; i < metricData.eventGroups->numEventGroups; i++) { + uint32_t all = + 1; // 1 means that all instances of the event will be collected + TC_CUPTI_CHECK(cuptiEventGroupSetAttribute( + metricData.eventGroups->eventGroups[i], + CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES, + sizeof(all), + &all)); + TC_CUPTI_CHECK( + cuptiEventGroupEnable(metricData.eventGroups->eventGroups[i])); + } +} + +void eventCollectionEnd(MetricData& metricData) { + cudaDeviceSynchronize(); + + // for each group, read the event values from the group and record + // in metricData + for (uint64_t i = 0; i < metricData.eventGroups->numEventGroups; i++) { + CUpti_EventGroup group = metricData.eventGroups->eventGroups[i]; + CUpti_EventDomainID groupDomain; + uint32_t numEvents, numInstances, numTotalInstances; + CUpti_EventID* eventIds; + size_t groupDomainSize = sizeof(groupDomain); + size_t numEventsSize = sizeof(numEvents); + size_t numInstancesSize = sizeof(numInstances); + size_t numTotalInstancesSize = sizeof(numTotalInstances); + size_t valuesSize, eventIdsSize; + + TC_CUPTI_CHECK(cuptiEventGroupGetAttribute( + group, + CUPTI_EVENT_GROUP_ATTR_EVENT_DOMAIN_ID, + &groupDomainSize, + &groupDomain)); + + TC_CUPTI_CHECK(cuptiDeviceGetEventDomainAttribute( + metricData.device, + groupDomain, + CUPTI_EVENT_DOMAIN_ATTR_TOTAL_INSTANCE_COUNT, + &numTotalInstancesSize, + &numTotalInstances)); + TC_CUPTI_CHECK(cuptiEventGroupGetAttribute( + group, + CUPTI_EVENT_GROUP_ATTR_INSTANCE_COUNT, + &numInstancesSize, + &numInstances)); + TC_CUPTI_CHECK(cuptiEventGroupGetAttribute( + group, CUPTI_EVENT_GROUP_ATTR_NUM_EVENTS, &numEventsSize, &numEvents)); + eventIdsSize = numEvents * sizeof(CUpti_EventID); + eventIds = (CUpti_EventID*)malloc(eventIdsSize); + TC_CUPTI_CHECK(cuptiEventGroupGetAttribute( + group, CUPTI_EVENT_GROUP_ATTR_EVENTS, &eventIdsSize, eventIds)); + + std::vector values(numInstances); + + for (uint32_t j = 0; j < numEvents; j++) { + values.clear(); + valuesSize = sizeof(uint64_t) * numInstances; + TC_CUPTI_CHECK(cuptiEventGroupReadEvent( + group, + CUPTI_EVENT_READ_FLAG_NONE, + eventIds[j], + &valuesSize, + values.data())); + CHECK_EQ(numInstances, valuesSize / sizeof(uint64_t)); + CHECK_LE(metricData.eventValues.size(), metricData.numEvents) + << "Too many events collected, metric expects only " + << metricData.numEvents; + + // sum collect event values from all instances + auto sum = std::accumulate( + values.begin(), values.begin() + numInstances, uint64_t(0)); + + // normalize the event value to represent the total number of + // domain instances on the device + auto normalized = (sum * numTotalInstances) / numInstances; + + metricData.eventIds.push_back(eventIds[j]); + metricData.eventValues.push_back(normalized); + + if (FLAGS_cuda_profile_verbose_events) { + std::string line1, line2; + std::tie(line1, line2) = logEvent( + eventIds[j], + numInstances, + numTotalInstances, + sum, + normalized, + values.data()); + LOG(INFO) << line1; + LOG(INFO) << line2; + } + } + } + for (uint64_t i = 0; i < metricData.eventGroups->numEventGroups; i++) + TC_CUPTI_CHECK( + cuptiEventGroupDisable(metricData.eventGroups->eventGroups[i])); +} + +void CUPTIAPI getMetricValueCallback( + void* userdata, + CUpti_CallbackDomain domain, + CUpti_CallbackId cbid, + const CUpti_CallbackData* cbInfo) { + MetricData& metricData = *static_cast(userdata); + + // This callback is enabled only for launch so we shouldn't see + // anything else. + CHECK_EQ(cbid, CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel) + << "Unexpected cbid: " << cbid; + + // on entry, enable all the event groups being collected this pass, + // for metrics we collect for all instances of the event + if (cbInfo->callbackSite == CUPTI_API_ENTER) { + eventCollectionStart(metricData, cbInfo); + } + + // on exit, read and record event values + if (cbInfo->callbackSite == CUPTI_API_EXIT) { + eventCollectionEnd(metricData); + } +} +CUpti_MetricValueKind getValueKind(const CudaMetric& metric) { + CUpti_MetricValueKind valueKind; + size_t valueKindSize = sizeof(valueKind); + TC_CUPTI_CHECK(cuptiMetricGetAttribute( + metric.id, CUPTI_METRIC_ATTR_VALUE_KIND, &valueKindSize, &valueKind)); + return valueKind; +} + +void logMetric(const CudaMetric& metric) { + auto valueKind = getValueKind(metric); + std::stringstream ss; + ss << "Metric " << metric.name << " = "; + switch (valueKind) { + case CUPTI_METRIC_VALUE_KIND_DOUBLE: + ss << metric.value.metricValueDouble; + break; + case CUPTI_METRIC_VALUE_KIND_UINT64: + ss << metric.value.metricValueUint64; + break; + case CUPTI_METRIC_VALUE_KIND_INT64: + ss << metric.value.metricValueInt64; + break; + case CUPTI_METRIC_VALUE_KIND_PERCENT: + ss << metric.value.metricValuePercent << '%'; + break; + case CUPTI_METRIC_VALUE_KIND_THROUGHPUT: + ss << metric.value.metricValueThroughput << " bytes/sec"; + break; + case CUPTI_METRIC_VALUE_KIND_UTILIZATION_LEVEL: + ss << "utilization level " << metric.value.metricValueUtilizationLevel; + break; + default: + CHECK(false) << "Unknown metric value kind"; + } + LOG(INFO) << ss.str(); +} + +template +uint64_t vectorSizeBytes(const std::vector& v) { + return v.size() * sizeof(T); +} + +} // namespace + +CudaMetric::operator double() const { + auto valueKind = getValueKind(*this); + if (valueKind == CUPTI_METRIC_VALUE_KIND_DOUBLE) { + return value.metricValueDouble; + } else if (valueKind == CUPTI_METRIC_VALUE_KIND_PERCENT) { + return value.metricValuePercent; + } else { + CHECK(false) << "Invalid metric value conversion."; + return 0.0; + } +} + +CudaMetric::operator uint64_t() const { + auto valueKind = getValueKind(*this); + if (valueKind == CUPTI_METRIC_VALUE_KIND_UINT64) { + return value.metricValueUint64; + } else if (valueKind == CUPTI_METRIC_VALUE_KIND_THROUGHPUT) { + return value.metricValueThroughput; + } else { + CHECK(false) << "Invalid metric value conversion."; + return 0; + } +} + +CudaMetric::operator int64_t() const { + auto valueKind = getValueKind(*this); + if (valueKind == CUPTI_METRIC_VALUE_KIND_INT64) { + return value.metricValueInt64; + } else { + CHECK(false) << "Invalid metric value conversion."; + return 0; + } +} + +CudaMetric::CudaMetric(const char* name_, CUdevice device) : name{name_} { + TC_CUPTI_CHECK(cuptiMetricGetIdFromName(device, name.c_str(), &id)); + TC_CUPTI_CHECK(cuptiMetricGetNumEvents(id, &numberEvents)); +} + +CudaProfilingInfo CudaCuptiProfiler::Profile() { + std::lock_guard lock(runtimes().getDeviceMutex(device_)); + + TC_CUPTI_CHECK(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_KERNEL)); + ScopeGuard activityGuard{[]() { + TC_CUPTI_CHECK(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_KERNEL)); + }}; + TC_CUPTI_CHECK( + cuptiActivityRegisterCallbacks(bufferRequested, bufferCompleted)); + kernel_(); + cudaDeviceSynchronize(); + TC_CUPTI_CHECK(cuptiActivityFlushAll(0)); + + CudaProfilingInfo pi; + pi.runtime = runtimes().getRuntime(device_); + CHECK_GT(pi.runtime.count(), 0); + + CUpti_SubscriberHandle subscriber; + MetricData metricData; + TC_CUPTI_CHECK(cuptiSubscribe( + &subscriber, (CUpti_CallbackFunc)getMetricValueCallback, &metricData)); + ScopeGuard subscriberGuard{ + [&]() { TC_CUPTI_CHECK(cuptiUnsubscribe(subscriber)); }}; + TC_CUPTI_CHECK(cuptiEnableCallback( + 1, + subscriber, + CUPTI_CB_DOMAIN_DRIVER_API, + CUPTI_DRIVER_TRACE_CBID_cuLaunchKernel)); + + metricData.numEvents = std::accumulate( + metrics.begin(), + metrics.end(), + uint32_t(0), + [](uint32_t sum, const CudaMetric& metric) { + return sum + metric.numberEvents; + }); + metricData.device = device_; + + CUcontext context = 0; + TC_CUDA_DRIVERAPI_ENFORCE(cuCtxGetCurrent(&context)); + + CUpti_EventGroupSets* passData; + { + std::vector ids(metrics.size()); + std::transform( + metrics.begin(), + metrics.end(), + ids.begin(), + [](const CudaMetric& metric) { return metric.id; }); + TC_CUPTI_CHECK(cuptiMetricCreateEventGroupSets( + context, sizeof(CUpti_MetricID) * ids.size(), ids.data(), &passData)); + } + + ScopeGuard passDataGuard{ + [&]() { TC_CUPTI_CHECK(cuptiEventGroupSetsDestroy(passData)); }}; + + for (uint32_t pass = 0; pass < passData->numSets; pass++) { + LOG_IF(INFO, FLAGS_cuda_profile_verbose) << "Profiling Pass: " << pass; + metricData.eventGroups = passData->sets + pass; + kernel_(); + } + + auto getMetricValue = [&](const CudaMetric& metric) { + CUpti_MetricValue value; + TC_CUPTI_CHECK(cuptiMetricGetValue( + device_, + metric.id, + vectorSizeBytes(metricData.eventIds), + metricData.eventIds.data(), + vectorSizeBytes(metricData.eventValues), + metricData.eventValues.data(), + std::chrono::duration_cast(pi.runtime) + .count(), + &value)); + return value; + }; + + for (auto& metric : metrics) { + metric.value = getMetricValue(metric); + } + writeMetricValues(pi); + + if (FLAGS_cuda_profile_verbose) { + for (const auto& metric : metrics) { + logMetric(metric); + } + } + + return pi; +} + +void CudaCuptiProfiler::writeMetricValues(CudaProfilingInfo& pinfo) const { + for (const auto& metric : metrics) { + if (metric.name == "ipc") { + pinfo.ipc = metric; + continue; + } + if (metric.name == "gld_efficiency") { + pinfo.globalLoadEfficiency = metric; + continue; + } + if (metric.name == "gst_efficiency") { + pinfo.globalStoreEfficiency = metric; + continue; + } + if (metric.name == "shared_efficiency") { + pinfo.sharedMemoryEfficiency = metric; + continue; + } + if (metric.name == "achieved_occupancy") { + pinfo.achievedOccupancy = metric; + continue; + } + if (metric.name == "warp_execution_efficiency") { + pinfo.warpExecutionEfficiency = metric; + continue; + } + if (metric.name == "local_memory_overhead") { + pinfo.localMemoryOverhead = metric; + continue; + } + + CHECK(false) << "NYI: " << metric.name; + } +} + +} // namespace tc diff --git a/tc/core/cuda/cuda_profile.h b/tc/core/cuda/cuda_profile.h new file mode 100644 index 000000000..44b9ad8e5 --- /dev/null +++ b/tc/core/cuda/cuda_profile.h @@ -0,0 +1,71 @@ +/** + * 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. + */ +#pragma once + +#include +#include + +#include +#include + +#include "tc/core/utils/time.h" + +namespace tc { +struct CudaProfilingInfo { + Duration runtime; + double ipc; + double globalLoadEfficiency; + double globalStoreEfficiency; + double sharedMemoryEfficiency; + double localMemoryOverhead; + double achievedOccupancy; + double warpExecutionEfficiency; + + friend bool operator==( + const CudaProfilingInfo& a, + const CudaProfilingInfo& b); +}; + +struct CudaMetric { + CudaMetric(const char* name_, CUdevice device); + + std::string name; + CUpti_MetricID id; + uint32_t numberEvents; + + CUpti_MetricValue value; + + operator double() const; + operator uint64_t() const; + operator int64_t() const; +}; + +using KernelType = std::function; + +class CudaCuptiProfiler { + public: + CudaCuptiProfiler(KernelType kernel, CUdevice device); + CudaProfilingInfo Profile(); + + private: + void writeMetricValues(CudaProfilingInfo& pinfo) const; + + std::function kernel_; + CUdevice device_; + std::vector metrics; +}; + +} // namespace tc diff --git a/tc/core/cuda/cuda_rtc.cc b/tc/core/cuda/cuda_rtc.cc index a6a883905..47f18dbfd 100644 --- a/tc/core/cuda/cuda_rtc.cc +++ b/tc/core/cuda/cuda_rtc.cc @@ -18,6 +18,7 @@ #include #include +#include #include #include "tc/core/cuda/cuda.h" @@ -126,17 +127,15 @@ std::ostream& operator<<(std::ostream& os, const std::array& a) { } // namespace -Duration CudaRTCFunction::Launch( +KernelType CudaRTCFunction::MakeLaunch( + int dev, const std::array& grid, const std::array& block, unsigned int shared_mem, cudaStream_t stream, - std::vector params, - std::vector outputs, - std::vector inputs, - bool profile) const { - int dev; - TC_CUDA_RUNTIMEAPI_ENFORCE(cudaGetDevice(&dev)); + std::vector& params, + std::vector& outputs, + std::vector& inputs) const { if (perGpuModule_.count(dev) == 0) { CUmodule module; CUfunction function; @@ -169,7 +168,7 @@ Duration CudaRTCFunction::Launch( unsigned int bx = block[0]; unsigned int by = block[1]; unsigned int bz = block[2]; - auto launch = [&]() { + auto launch = [=]() mutable { TC_CUDA_DRIVERAPI_ENFORCE(cuLaunchKernel( perGpuKernel_.at(dev), gx, @@ -184,6 +183,22 @@ Duration CudaRTCFunction::Launch( 0)); }; + return launch; +} + +Duration CudaRTCFunction::Launch( + const std::array& grid, + const std::array& block, + unsigned int shared_mem, + cudaStream_t stream, + std::vector params, + std::vector outputs, + std::vector inputs, + bool profile) const { + int dev; + TC_CUDA_RUNTIMEAPI_ENFORCE(cudaGetDevice(&dev)); + auto launch = + MakeLaunch(dev, grid, block, shared_mem, stream, params, outputs, inputs); if (not profile) { launch(); return Duration::max(); @@ -204,4 +219,20 @@ Duration CudaRTCFunction::Launch( TC_CUDA_RUNTIMEAPI_ENFORCE(cudaEventDestroy(stop)); return std::chrono::microseconds(static_cast(milliseconds * 1000)); } + +CudaProfilingInfo CudaRTCFunction::Profile( + const std::array& grid, + const std::array& block, + unsigned int shared_mem, + cudaStream_t stream, + std::vector params, + std::vector outputs, + std::vector inputs) const { + int dev; + TC_CUDA_RUNTIMEAPI_ENFORCE(cudaGetDevice(&dev)); + auto launch = + MakeLaunch(dev, grid, block, shared_mem, stream, params, outputs, inputs); + CudaCuptiProfiler profiler(launch, dev); + return profiler.Profile(); +} } // namespace tc diff --git a/tc/core/cuda/cuda_rtc.h b/tc/core/cuda/cuda_rtc.h index 4df57430f..f97abb5d7 100644 --- a/tc/core/cuda/cuda_rtc.h +++ b/tc/core/cuda/cuda_rtc.h @@ -25,6 +25,7 @@ #include #include // cuda driver types +#include "tc/core/cuda/cuda_profile.h" #include "tc/core/utils/time.h" namespace tc { @@ -59,9 +60,28 @@ class CudaRTCFunction { std::vector inputs, bool profile = false) const; + CudaProfilingInfo Profile( + const std::array& grid, + const std::array& block, + unsigned int shared_mem, + cudaStream_t stream, + std::vector params, + std::vector outputs, + std::vector inputs) const; + void clear(); private: + KernelType MakeLaunch( + int dev, + const std::array& grid, + const std::array& block, + unsigned int shared_mem, + cudaStream_t stream, + std::vector& params, + std::vector& outputs, + std::vector& inputs) const; + mutable std::unordered_map perGpuModule_; mutable std::unordered_map perGpuKernel_; std::string specializedName; diff --git a/tc/core/cuda/cuda_tc_executor.cc b/tc/core/cuda/cuda_tc_executor.cc index 8590123a7..88b06d0dd 100644 --- a/tc/core/cuda/cuda_tc_executor.cc +++ b/tc/core/cuda/cuda_tc_executor.cc @@ -178,10 +178,9 @@ void CudaTcExecutor::compileWithTcMapper() { LOG_IF(INFO, FLAGS_dump_cuda) << "generatedCuda: " << cudaSource; } -Duration CudaTcExecutor::run( +void CudaTcExecutor::preRunChecks( const std::vector& inputs, - const std::vector& outputs, - bool profile) const { + const std::vector& outputs) const { CHECK(rtcFun) << "Can't launch uncompiled: " << executionInfo_.kernelName; CHECK_NE(executionInfo_.options, ""); checkSizesAndStridesAreCompliant( @@ -191,6 +190,14 @@ Duration CudaTcExecutor::run( executionInfo_.outputsInfo, halideComponents_.getDef().returns()); + CHECK_NE(grid.view[0], 0) << "Grid dims are not set up"; + CHECK_NE(block.view[0], 0) << "Block dims are not set up"; +} + +std::pair, std::vector> +CudaTcExecutor::prepareCudaArgs( + const std::vector& inputs, + const std::vector& outputs) const { std::vector I; std::vector O; for (size_t i = 0; i < inputs.size(); ++i) { @@ -199,6 +206,15 @@ Duration CudaTcExecutor::run( for (size_t i = 0; i < outputs.size(); ++i) { O.push_back(outputs[i]->data); } + return std::make_pair(std::move(I), std::move(O)); +} + +Duration CudaTcExecutor::run( + const std::vector& inputs, + const std::vector& outputs, + bool profile) const { + preRunChecks(inputs, outputs); + auto IO = prepareCudaArgs(inputs, outputs); 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"; @@ -208,9 +224,10 @@ Duration CudaTcExecutor::run( 0, stream, executionInfo_.kernelParams, - O, - I, + IO.second, + IO.first, profile); + if (profile and OptionsCache::cacheEnabled()) { OptionsCache::getCache()->recordRuntime( cacheKeyId_, @@ -222,6 +239,31 @@ Duration CudaTcExecutor::run( return res; } +CudaProfilingInfo CudaTcExecutor::profile( + const std::vector& inputs, + const std::vector& outputs) const { + preRunChecks(inputs, outputs); + auto IO = prepareCudaArgs(inputs, outputs); + cudaStream_t stream = 0; + auto res = rtcFun->Profile( + grid.view.extractDefaultedArray(), + block.view.extractDefaultedArray(), + 0, + stream, + executionInfo_.kernelParams, + IO.second, + IO.first); + if (OptionsCache::cacheEnabled()) { + OptionsCache::getCache()->recordProfilingInfo( + cacheKeyId_, + CudaMappingOptions(executionInfo_.options), + inputs, + constPtrs(outputs), + res); + } + return res; +} + void CudaTcExecutor::uncheckedRun( const std::vector& inputs, const std::vector& outputs) const { diff --git a/tc/core/cuda/cuda_tc_executor.h b/tc/core/cuda/cuda_tc_executor.h index 2f9b68cca..dfb3066e0 100644 --- a/tc/core/cuda/cuda_tc_executor.h +++ b/tc/core/cuda/cuda_tc_executor.h @@ -33,7 +33,7 @@ namespace tc { class CudaTcExecutor : public ::tc::TcExecutor { public: using MappingOptionsType = CudaMappingOptions; - + using ProfilingInfoType = CudaProfilingInfo; CudaTcExecutor( std::string id, const std::vector& inputsInfo, @@ -73,7 +73,11 @@ class CudaTcExecutor : public ::tc::TcExecutor { Duration run( const std::vector& inputs, const std::vector& outputs, - bool profile = false) const; + bool profile = false) const override; + + CudaProfilingInfo profile( + const std::vector& inputs, + const std::vector& outputs) const; // This is the "low-latency" mode in which we just propagate raw pointers to // data in GPU address space. @@ -82,7 +86,7 @@ class CudaTcExecutor : public ::tc::TcExecutor { // doesn't then segfault will likely occur. void uncheckedRun( const std::vector& inputs, - const std::vector& outputs) const; + const std::vector& outputs) const override; bool hasRuntimeCompiledFunction() override { return rtcFun.get() != nullptr; @@ -102,6 +106,12 @@ class CudaTcExecutor : public ::tc::TcExecutor { } private: + void preRunChecks( + const std::vector& inputs, + const std::vector& outputs) const; + std::pair, std::vector> prepareCudaArgs( + const std::vector& inputs, + const std::vector& outputs) const; void compileWithTcMapper(); public: diff --git a/tc/core/execution_engine-inl.h b/tc/core/execution_engine-inl.h index c7e3be92a..fb3c1fc9a 100644 --- a/tc/core/execution_engine-inl.h +++ b/tc/core/execution_engine-inl.h @@ -114,6 +114,68 @@ size_t ExecutionEngine::compile( return handle; } +template +std::unique_ptr ExecutionEngine::borrowExecutor( + size_t handle) { + std::unique_ptr executorUPtr(nullptr); + { + std::lock_guard lg(tcExecutorMutex_); + std::swap(executorUPtr, executors_[handle]); + } + return std::move(executorUPtr); +} + +namespace { +template +struct MaybeVoid { + template + void getValue(const F& f) { + value = f(); + } + + T get() { + return value; + } + + T value; +}; + +template <> +struct MaybeVoid { + template + void getValue(const F& f) { + f(); + } + + void get() {} +}; + +} // namespace + +template +template +R ExecutionEngine::execute( + size_t handle, + std::unique_ptr& executor, + std::function func) { + CHECK(executor); + CHECK(executor->hasRuntimeCompiledFunction()); + MaybeVoid res; + try { + // Must catch and swap to avoid exception in destructor! + res.getValue([&]() { return func(*executor); }); + } catch (std::exception& e) { + std::lock_guard lg(tcExecutorMutex_); + std::swap(executor, executors_[handle]); + throw; + } + { + std::lock_guard lg(tcExecutorMutex_); + std::swap(executor, executors_[handle]); + } + return res.get(); +} + // Steal the executor instance and give it back under lock. // Run outside of lock on owning ExecutorType. template @@ -123,36 +185,41 @@ Duration ExecutionEngine::run( const std::vector& outputs, bool profile, std::function pruningFunction) { - std::unique_ptr executorUPtr(nullptr); - { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - } + auto executorUPtr = borrowExecutor(handle); // It turns out someone else may already be running this configuration in // some unexpected cases: there is no guarantee of no-redundancy in // compilation options. In that case, we swapped 2 nullptrs and we just // exit. - Duration res(Duration::max()); if (executorUPtr) { if (pruningFunction(static_cast(executorUPtr.get()))) { return Duration::max(); } - CHECK(executorUPtr->hasRuntimeCompiledFunction()); - try { - // Must catch and swap to avoid exception in destructor! - res = executorUPtr->run(inputs, outputs, profile); - } catch (std::exception& e) { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - throw; - } - { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - } + return execute(handle, executorUPtr, [&](ExecutorType& exec) { + return exec.run(inputs, outputs, profile); + }); } - return res; + return Duration::max(); +} + +// Steal the executor instance and give it back under lock. +// Run outside of lock on owning ExecutorType. +template +typename ExecutorType::ProfilingInfoType ExecutionEngine::profile( + size_t handle, + const std::vector& inputs, + const std::vector& outputs) { + auto executorUPtr = borrowExecutor(handle); + + using Ptype = typename ExecutorType::ProfilingInfoType; + if (executorUPtr) { + return execute(handle, executorUPtr, [&](ExecutorType& exec) { + return exec.profile(inputs, outputs); + }); + } + Ptype pInfo; + pInfo.runtime = Duration::max(); + return pInfo; } // Steal ExecutorType and give it back under lock @@ -162,30 +229,16 @@ void ExecutionEngine::uncheckedRun( size_t handle, const std::vector& inputs, const std::vector& outputs) { - std::unique_ptr executorUPtr(nullptr); - { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - } + auto executorUPtr = borrowExecutor(handle); // It turns out someone else may already be running this configuration in // some unexpected cases: there is no guarantee of no-redundancy in // compilation options. In that case, we swapped 2 nullptrs and we just // exit. if (executorUPtr) { - CHECK(executorUPtr->hasRuntimeCompiledFunction()); - try { - // Must catch and swap to avoid exception in destructor! - executorUPtr->uncheckedRun(inputs, outputs); - } catch (std::exception& e) { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - throw; - } - { - std::lock_guard lg(tcExecutorMutex_); - std::swap(executorUPtr, executors_[handle]); - } + execute(handle, executorUPtr, [&](ExecutorType& exec) { + return exec.uncheckedRun(inputs, outputs); + }); } } diff --git a/tc/core/execution_engine.h b/tc/core/execution_engine.h index 6ec252103..cdfd8f02e 100644 --- a/tc/core/execution_engine.h +++ b/tc/core/execution_engine.h @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include @@ -72,10 +73,20 @@ class ExecutionEngine { std::function pruningFunction = [](const ExecutorType*) { return false; }); - /// "Low-latency" execution mode in which we just propagate raw pointers to - /// data in GPU address space. - /// No tensor-related information can be checked so it is the user's - /// responsibility to ensure that shapes and strides match. + /// Run a compiled TC kernel given its handle, on the given input tensors, and + /// profile it. All tensors must be allocated and have appropriate + /// shapes (inputs same as for copmilation, outputs same as returned by + /// inferOutputTensorInfo). + /// \returns The kernel's profiling results. + typename ExecutorType::ProfilingInfoType profile( + size_t handle, + const std::vector& inputs, + const std::vector& outputs); + + /// "Low-latency" execution mode in which we just propagate raw pointers + /// to data in GPU address space. No tensor-related information can be + /// checked so it is the user's responsibility to ensure that shapes and + /// strides match. void uncheckedRun( size_t handle, const std::vector& inputs, @@ -85,6 +96,14 @@ class ExecutionEngine { void clear(size_t handle); protected: + template + R execute( + size_t handle, + std::unique_ptr& executor, + std::function func); + + std::unique_ptr borrowExecutor(size_t handle); + size_t emplaceExecutor(std::unique_ptr p); size_t getHandle( diff --git a/tc/core/flags.cc b/tc/core/flags.cc index e2f4f5a25..c54f75977 100644 --- a/tc/core/flags.cc +++ b/tc/core/flags.cc @@ -37,6 +37,9 @@ DEFINE_bool( "Print debug spew for the tc_mapper like cuda code, mapping options etc"); DEFINE_bool(dump_cuda, false, "Print the generated cudaSource"); +DEFINE_bool(cuda_profile_verbose, false, "Verbose profiling"); +DEFINE_bool(cuda_profile_verbose_events, false, "Verbose event profiling"); + // CPU codegen options DEFINE_bool(llvm_dump_before_opt, false, "Print IR before optimization"); DEFINE_bool(llvm_dump_after_opt, false, "Print IR after optimization"); @@ -97,6 +100,10 @@ DEFINE_bool( tuner_gen_log_generations, false, "Log each generation's runtimes."); +DEFINE_bool( + tuner_gen_profiled_run, + false, + "Collect performance metrics during tuning."); DEFINE_uint64( tuner_min_launch_total_threads, 64, diff --git a/tc/core/flags.h b/tc/core/flags.h index 0324b00ad..83d03b5b4 100644 --- a/tc/core/flags.h +++ b/tc/core/flags.h @@ -30,6 +30,9 @@ DECLARE_bool(debug_cuda); DECLARE_bool(debug_tuner); DECLARE_bool(dump_cuda); +DECLARE_bool(cuda_profile_verbose); +DECLARE_bool(cuda_profile_verbose_events); + // llvm codegen DECLARE_bool(llvm_dump_before_opt); DECLARE_bool(llvm_dump_after_opt); @@ -51,6 +54,7 @@ DECLARE_string(tuner_rng_restore); DECLARE_bool(tuner_gen_restore_from_proto); DECLARE_uint32(tuner_gen_restore_number); DECLARE_bool(tuner_gen_log_generations); +DECLARE_bool(tuner_gen_profiled_run); DECLARE_uint64(tuner_min_launch_total_threads); DECLARE_uint32(tuner_save_best_candidates_count); diff --git a/tc/proto/compcache.proto b/tc/proto/compcache.proto index 64b57b062..79321245f 100644 --- a/tc/proto/compcache.proto +++ b/tc/proto/compcache.proto @@ -43,9 +43,21 @@ message ManualCudaCacheEntryProto { repeated uint64 block_dims = 8; } +message CudaProfilingProto{ + required uint64 runtime = 1; + required double ipc = 2; + required double globalLoadEfficiency = 4; + required double globalStoreEfficiency = 5; + required double sharedMemoryEfficiency = 7; + required double localMemoryOverhead = 9; + required double achievedOccupancy = 10; + required double warpExecutionEfficiency = 11; +} + message OptionsCacheValuesProto{ required CudaMappingOptionsProto kernel_options = 1; repeated uint64 recorded_runtimes = 2; + repeated CudaProfilingProto profiles = 3; } message OptionsCacheEntryProto { diff --git a/test/cuda/test_autotuner.cc b/test/cuda/test_autotuner.cc index 7de32a0c3..581a34ba8 100644 --- a/test/cuda/test_autotuner.cc +++ b/test/cuda/test_autotuner.cc @@ -42,7 +42,7 @@ DEFINE_bool( "load options from previously stored cache (or store them)"); DEFINE_bool(no_memory_promotion, false, "disable memory promotion"); -struct ATenCompilationUnitTest : public ::testing::Test { +struct ATenCompilationUnitTest : public ::testing::TestWithParam { static constexpr uint32_t N = 32, C1 = 512, C2 = 8, C3 = 2, H = 28, W = 28; ATenCompilationUnitTest() { @@ -53,6 +53,7 @@ struct ATenCompilationUnitTest : public ::testing::Test { tc::FLAGS_tuner_threads = std::min(8u, tc::FLAGS_tuner_gen_pop_size); tc::FLAGS_tuner_gen_number_elites = tc::FLAGS_tuner_gen_pop_size / 4; } + tc::FLAGS_tuner_gen_profiled_run = GetParam(); } void Check( @@ -90,7 +91,12 @@ struct ATenCompilationUnitTest : public ::testing::Test { } }; -TEST_F(ATenCompilationUnitTest, LayerNorm) { +INSTANTIATE_TEST_CASE_P( + WithAndWithoutProfiling, + ATenCompilationUnitTest, + ::testing::Bool()); + +TEST_P(ATenCompilationUnitTest, LayerNorm) { at::Tensor mat1 = at::CUDA(at::kFloat).rand({7, 32, 64}); std::vector inputs = {mat1}; std::vector outputs; @@ -113,7 +119,7 @@ def layernorm(float(T, B, C) I) -> (O, mean, centered, var) { autotune(cacheFilename, TC, name, inputs, options, {options}); } -TEST_F(ATenCompilationUnitTest, MatmulA) { +TEST_P(ATenCompilationUnitTest, MatmulA) { at::Tensor mat1 = at::CUDA(at::kFloat).rand({3, 4}); at::Tensor mat2 = at::CUDA(at::kFloat).rand({4, 5}); std::vector inputs = {mat1, mat2}; @@ -132,7 +138,7 @@ def matmul(float(M,N) A, float(N,K) B) -> (output) { autotune(cacheFilename, TC, name, inputs, options, {options}); } -TEST_F(ATenCompilationUnitTest, MatmulB) { +TEST_P(ATenCompilationUnitTest, MatmulB) { at::Tensor mat1 = at::CUDA(at::kFloat).rand({72, 26}); at::Tensor mat2 = at::CUDA(at::kFloat).rand({26, 72}); std::vector inputs = {mat1, mat2}; @@ -151,7 +157,7 @@ def matmul(float(M,N) A, float(N,K) B) -> (output) { autotune(cacheFilename, TC, name, inputs, options, {options}); } -TEST_F(ATenCompilationUnitTest, MatmulC) { +TEST_P(ATenCompilationUnitTest, MatmulC) { at::Tensor mat1 = at::CUDA(at::kFloat).rand({100, 400}); at::Tensor mat2 = at::CUDA(at::kFloat).rand({400, 500}); std::vector inputs = {mat1, mat2}; @@ -170,7 +176,7 @@ def matmul(float(M,N) A, float(N,K) B) -> (output) { autotune(cacheFilename, TC, name, inputs, options, {options}); } -TEST_F(ATenCompilationUnitTest, TensorDot) { +TEST_P(ATenCompilationUnitTest, TensorDot) { at::Tensor I0 = at::CUDA(at::kFloat).rand({N, C1, C2, H, W}); at::Tensor I1 = at::CUDA(at::kFloat).rand({N, C2, C3, H, W}); std::vector inputs = {I0, I1}; diff --git a/test/cuda/test_compilation_cache.cc b/test/cuda/test_compilation_cache.cc index fb5e53edc..c0b5fedbf 100644 --- a/test/cuda/test_compilation_cache.cc +++ b/test/cuda/test_compilation_cache.cc @@ -775,6 +775,18 @@ TEST_F(OptionsCacheTest, Serialization) { inputPtrs, outputPtrs, std::chrono::microseconds(11)); + tc::CudaProfilingInfo pInfoOrig; + pInfoOrig.runtime = std::chrono::microseconds(444); + pInfoOrig.ipc = 1.23; + pInfoOrig.globalLoadEfficiency = 6.546; + pInfoOrig.globalStoreEfficiency = 7.123; + pInfoOrig.sharedMemoryEfficiency = 7.111111; + pInfoOrig.localMemoryOverhead = 888.222; + pInfoOrig.achievedOccupancy = 11231231.14; + pInfoOrig.warpExecutionEfficiency = 910293123918239.1029381; + + tc::OptionsCache::getCache()->recordProfilingInfo( + "kernel0", options0, inputPtrs, outputPtrs, pInfoOrig); tc::OptionsCache::getCache()->recordRuntime( "kernel1", options0, inputPtrs, outputPtrs, std::chrono::microseconds(1)); @@ -791,8 +803,9 @@ TEST_F(OptionsCacheTest, Serialization) { "kernel0", inputPtrs, outputPtrs); ASSERT_EQ(ret.size(), 2u); ASSERT_EQ(ret[0].options, options0); - ASSERT_EQ(ret[0].recordedRuntimes.size(), 1u); + ASSERT_EQ(ret[0].recordedRuntimes.size(), 2u); ASSERT_EQ(ret[0].recordedRuntimes[0], std::chrono::microseconds(10)); + ASSERT_EQ(ret[0].recordedRuntimes[1], std::chrono::microseconds(444)); ASSERT_EQ(ret[1].options, options1); ASSERT_EQ(ret[1].recordedRuntimes.size(), 1u); @@ -809,10 +822,18 @@ TEST_F(OptionsCacheTest, Serialization) { "kernel2", inputPtrs, outputPtrs); ASSERT_EQ(ret.size(), 0u); + + auto ret2 = tc::OptionsCache::getCache()->retrieveOptionsAndProfilingInfo( + "kernel0", inputPtrs, outputPtrs); + ASSERT_EQ(ret2.size(), 1u); + ASSERT_EQ(ret2.front().options, options0); + ASSERT_EQ(ret2.front().profilingInfo.size(), 1u); + ASSERT_EQ(ret2.front().profilingInfo.front(), pInfoOrig); + ASSERT_EQ(tc::OptionsCache::getCache()->size(), 2u); ASSERT_EQ(tc::OptionsCache::getCache()->totalSize(), 3u); - ASSERT_EQ(tc::OptionsCache::getCache()->numberAttemptedRetrievals, 3); - ASSERT_EQ(tc::OptionsCache::getCache()->numberSuccessfulRetrievals, 2); + ASSERT_EQ(tc::OptionsCache::getCache()->numberAttemptedRetrievals, 4); + ASSERT_EQ(tc::OptionsCache::getCache()->numberSuccessfulRetrievals, 3); ASSERT_EQ(tc::OptionsCache::getCache()->numberCacheAttemps, 0); }