Skip to content

Remove validateMinimumArch #173

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
zasdfgbnm merged 4 commits into from
Apr 17, 2023
Merged

Remove validateMinimumArch #173

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
5a60745
Remove `validateMinimumArch`
zasdfgbnm Apr 17, 2023
de54e73
remove ArchCheck
zasdfgbnm Apr 17, 2023
2011265
save
zasdfgbnm Apr 17, 2023
1ee7e69
Merge branch 'main' into disable-arch-validate
zasdfgbnm Apr 17, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 0 additions & 8 deletions csrc/executor_utils.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1378,10 +1378,6 @@ std::tuple<NvrtcFunction, std::string, std::vector<char>> getCompiledKernel(
const int max_register_heuristic,
bool return_compiled_binary) {
FUSER_PERF_SCOPE("executor_utils::NVRTC");
if (isOptionDisabled(DisableOption::ArchCheck)) {
TORCH_WARN(
"NVFuser Compile: arch check disabled, should not compile any kernel");
}

at::cuda::jit::initializeCudaContext();

Expand Down Expand Up @@ -1477,10 +1473,6 @@ std::tuple<NvrtcFunction, std::string, std::vector<char>> getCompiledKernel(
object_code.clear();
}

TORCH_CHECK(
!isOptionDisabled(DisableOption::ArchCheck),
"NVFuser Compile: arch check disabled, should not return any compiled kernel");

return {compiled_kernel, log.str(), object_code};
}

Expand Down
22 changes: 0 additions & 22 deletions csrc/lower_validation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -844,21 +844,6 @@ void validatePartialSplit(Fusion* fusion) {

namespace {

//! Utility to make sure targeted gpu capability is
//! higher than provided major.minor.
void validateMinimumArch(int major, int minor) {
// Skip checking arch if disabled.
if (isOptionDisabled(DisableOption::ArchCheck)) {
Copy link
Collaborator

@naoyam naoyam Apr 17, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this option still used?

Copy link
Collaborator Author

@zasdfgbnm zasdfgbnm Apr 17, 2023 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, in executor_utils.cpp getCompiledKernel, there is a

  TORCH_CHECK(
      !isOptionDisabled(DisableOption::ArchCheck),
      "NVFuser Compile: arch check disabled, should not return any compiled kernel");

don't think this makes much sense either.... Let me remove it.

return;
}

auto prop = at::cuda::getCurrentDeviceProperties();
TORCH_INTERNAL_ASSERT(prop->major >= major);
if (prop->major == major) {
TORCH_INTERNAL_ASSERT(prop->minor >= minor);
}
}

//! Validates that the operand and result tensors
//! of mma ops are swizzled and also validates
//! specialization of tidx as lane id.
Expand Down Expand Up @@ -1024,12 +1009,10 @@ void validateArchMemoryOp(LoadStoreOp* ldst) {
switch (ldst->opType()) {
case LoadStoreOpType::LdMatrix:
case LoadStoreOpType::LdMatrixTranspose:
validateMinimumArch(7, 5);
validateLdMatrixOutput(ldst->out()->as<TensorView>());
return;
case LoadStoreOpType::CpAsyncCg:
case LoadStoreOpType::CpAsyncCa:
validateMinimumArch(8, 0);
return;
default:
return;
Expand All @@ -1049,21 +1032,16 @@ void validateMma(Fusion* fusion) {

switch (mma->options().macro) {
case MmaOptions::MacroType::Volta_16_16_4:
validateMinimumArch(7, 0);
break;
case MmaOptions::MacroType::Turing_16_8_16:
case MmaOptions::MacroType::Turing_16_16_16:
validateMinimumArch(7, 5);

// Check that operands come from ldmatrix, can be
// relaxed once swizzles can be labeled on iterdomains.
validateTuringMmaInput(mma->inA()->as<TensorView>());
validateTuringMmaInput(mma->inB()->as<TensorView>());
break;
case MmaOptions::MacroType::Ampere_16_8_16:
case MmaOptions::MacroType::Ampere_16_16_16:
validateMinimumArch(8, 0);

// Check that operands come from ldmatrix, can be
// relaxed once swizzles can be labeled on iterdomains.
validateTuringMmaInput(mma->inA()->as<TensorView>());
Expand Down
1 change: 0 additions & 1 deletion csrc/utils.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -150,7 +150,6 @@ const auto& getDebugDumpOptions() {

auto parseDisableOptions() {
const std::unordered_map<std::string, DisableOption> available_options = {
{"arch_check", DisableOption::ArchCheck},
{"compile_to_sass", DisableOption::CompileToSass},
{"fallback", DisableOption::Fallback},
{"fma", DisableOption::Fma},
Expand Down
1 change: 0 additions & 1 deletion csrc/utils.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,7 +94,6 @@ TORCH_CUDA_CU_API const std::vector<std::string>& getDebugDumpArguments(
//! These can be set through the `PYTORCH_NVFUSER_DISABLE` environment variable
//!
enum class DisableOption {
ArchCheck, //! Disable hardware-specific checks to enable cross arch debug
CompileToSass, //! Disable direct compilation to sass so the ptx can be
//! examined
Fallback, //! Disable fallback
Expand Down