Convert more NVF_ERROR(false, to use NVF_THROW (#2955)

PR #2941 did a lot of conversions but only hit the cases where the
arguments to `NVF_ERROR` were not wrapped to the next line. This PR
picks up those line-wrapped cases.

For those interested, these were found using
```
rg --multiline 'NVF_ERROR\((\n|\s)*false' -l
```
and those files were edited in vim and processed using the vim command
`:bufdo %s/NVF_ERROR(\_s*false,\_s*/NVF_THROW(/g`

benchmarks/cpp/utils.cpp

@@ -124,8 +124,7 @@ std::string toString(const std::shared_ptr<HeuristicParams>& params) {
   if (tparams) {
     return toString(*tparams);
   }
-  NVF_ERROR(
-      false,
+  NVF_THROW(
       "Unknown heuristic parameter type. Did you just added a new heuristic parameter type but forget to update here?");
 }
 

csrc/codegen.cpp

@@ -1974,8 +1974,7 @@ class CudaKernelGenerator : private kir::ConstIrVisitor {
       }
       return;
     } else {
-      NVF_ERROR(
-          false, "Non-allreduce grouped grid welford is not yet supported");
+      NVF_THROW("Non-allreduce grouped grid welford is not yet supported");
     }
   }
 
@@ -2858,8 +2857,7 @@ class CudaKernelGenerator : private kir::ConstIrVisitor {
   }
 
   void handle(const GroupedWelfordOp* grouped_wop) final {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Should not reach here as grouped welford is only enabled for grid welford,",
         " which is handled by its own handler");
   }

csrc/device_lower/analysis/fused_reduction.cpp

@@ -161,8 +161,7 @@ class FusionInspector : private IterVisitor {
       } else if (preceding_expr->isA<WelfordOp>()) {
         fusion_list_.emplace_back(preceding_expr->as<WelfordOp>(), true);
       } else {
-        NVF_ERROR(
-            false, "Invalid preceding expr: ", preceding_expr->toString());
+        NVF_THROW("Invalid preceding expr: ", preceding_expr->toString());
       }
 
       fused_exprs_.insert(preceding_expr);

csrc/device_lower/analysis/predicate_elimination.cpp

@@ -988,8 +988,7 @@ bool PredicateElimination::setReductionInitValue(
   } else if (existing_val->sameAs(reduction_init)) {
     return true;
   } else {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Inconsistent setting of initialization value for t",
         tv->name(),
         ". Prev: ",

csrc/device_lower/analysis/sync_information.cpp

@@ -780,8 +780,7 @@ SyncMap::SyncMap(Fusion* fusion) {
                 continue;
               }
               // Can this happen?
-              NVF_ERROR(
-                  false,
+              NVF_THROW(
                   "Unexpected case. Producer: ",
                   producer->toString(),
                   ", consumer: ",

csrc/device_lower/pass/alias_memory.cpp

@@ -483,8 +483,7 @@ class ScopeMap : private kir::IrVisitor {
   }
 
   void handle(kir::IfThenElse* ite) final {
-    NVF_ERROR(
-        false, "lower_alias_memory: no support for IfThenElse at this phase.");
+    NVF_THROW("lower_alias_memory: no support for IfThenElse at this phase.");
   }
 
   //! Factory function for internal loop information data
@@ -786,8 +785,7 @@ class AllocationInfoMap : private kir::IrVisitor {
   }
 
   void handle(kir::IfThenElse* ite) final {
-    NVF_ERROR(
-        false, "lower_alias_memory: no support for IfThenElse at this phase.");
+    NVF_THROW("lower_alias_memory: no support for IfThenElse at this phase.");
   }
 
   // Generate allocation info for allocation after some pre-filtering

csrc/device_lower/pass/allocation.cpp

@@ -599,8 +599,7 @@ class AllocationInserter : public kir::ExprMutator {
   }
 
   void handle(kir::IfThenElse*) final {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Pass does not support conditional statements, ",
         "this pass should be run before any conditionals are placed in code.");
   }

csrc/device_lower/pass/index.cpp

@@ -1166,8 +1166,7 @@ void IndexLowering::handle(const GroupedWelfordOp* grouped_wop) {
     handleGroupedGridWelford(
         grouped_wop, indexed_outputs, indexed_inputs, grouped_wop->initVals());
   } else {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Only grid welford is supported. Validation should have caught non-grid welford grouping.");
   }
 }
@@ -1431,8 +1430,7 @@ void IndexLowering::handle(const kir::MBarrierInvalidate* minval) {
     smem_address_ptr = lower_utils::u32IndexScalarSmemTv(
         minval->mbarrier()->as<kir::TensorIndex>());
   } else {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Unexpected MBarrierInvalidate barrier value: ",
         minval->mbarrier()->toString());
   }

csrc/device_lower/pass/insert_syncs.cpp

@@ -508,8 +508,7 @@ class ReadAfterWriteSyncs : public kir::ExprMutator {
   }
 
   void handle(kir::IfThenElse*) final {
-    NVF_ERROR(
-        false,
+    NVF_THROW(
         "Pass does not support conditional statements, ",
         "this pass should be run before any conditionals are placed in code.");
   }

csrc/dispatch.cpp

@@ -70,8 +70,7 @@ void Val::dispatch(T handler, Val* val) {
       ptr(handler)->handle(val);
       return;
   }
-  NVF_ERROR(
-      false,
+  NVF_THROW(
       "Unknown valtype in dispatch! val: ",
       val->toString(),
       " = ",
@@ -141,8 +140,7 @@ void Val::constDispatch(T handler, const Val* val) {
       ptr(handler)->handle(val);
       return;
   }
-  NVF_ERROR(
-      false,
+  NVF_THROW(
       "Unknown valtype in dispatch! val: ",
       val->toString(),
       " = ",
@@ -303,29 +301,21 @@ void OptOutConstDispatch::dispatch(const Val* v) {
 
 void OptInConstDispatch::unhandled(const Statement* stmt) {
   if (stmt->isExpr()) {
-    NVF_ERROR(
-        false,
-        "Handle not overriden for ",
-        stmt->as<Expr>()->getOpString(),
-        ".");
+    NVF_THROW(
+        "Handle not overriden for ", stmt->as<Expr>()->getOpString(), ".");
   } else if (stmt->isVal()) {
-    NVF_ERROR(
-        false, "Handle not overriden for ", stmt->getValType().value(), ".");
+    NVF_THROW("Handle not overriden for ", stmt->getValType().value(), ".");
   } else {
     NVF_THROW("Unrecognized statement type.");
   }
 }
 
 void OptInDispatch::unhandled(Statement* stmt) {
   if (stmt->isExpr()) {
-    NVF_ERROR(
-        false,
-        "Handle not overriden for ",
-        stmt->as<Expr>()->getOpString(),
-        ".");
+    NVF_THROW(
+        "Handle not overriden for ", stmt->as<Expr>()->getOpString(), ".");
   } else if (stmt->isVal()) {
-    NVF_ERROR(
-        false, "Handle not overriden for ", stmt->getValType().value(), ".");
+    NVF_THROW("Handle not overriden for ", stmt->getValType().value(), ".");
   } else {
     NVF_THROW("Unrecognized statement type.");
   }

csrc/evaluator_common.cpp

@@ -565,8 +565,7 @@ void NaiveValueMachine::runUnaryOp(int index) {
       } else if (data_type_[index] == DataType::Bool) {
         dest = PolymorphicValue((bool)src);
       } else {
-        NVF_ERROR(
-            false, "dtype not supported in evaluator: ", data_type_[index]);
+        NVF_THROW("dtype not supported in evaluator: ", data_type_[index]);
       }
       break;
     case UnaryOpType::Abs:

csrc/fusion_executor/executor_kernel_arg.cpp

@@ -227,8 +227,7 @@ std::vector<std::byte> polymorphicValueToBytes(
       int32_t v32 = (int32_t)v;
       return std::vector<std::byte>((std::byte*)&v32, (std::byte*)&v32 + 4);
     } else {
-      NVF_ERROR(
-          false,
+      NVF_THROW(
           "Cannot convert int64_t to ",
           dtype,
           " type: only int32 and int64 are supported.");
@@ -265,8 +264,7 @@ std::vector<std::byte> polymorphicValueToBytes(
       return std::vector<std::byte>(
           (std::byte*)&v8, (std::byte*)&v8 + sizeof(at::Float8_e5m2));
     } else {
-      NVF_ERROR(
-          false,
+      NVF_THROW(
           "Cannot convert double to ",
           dtype,
           " type: only half, bfloat16, float and double are supported.");
@@ -282,8 +280,7 @@ std::vector<std::byte> polymorphicValueToBytes(
       return std::vector<std::byte>(
           (std::byte*)&v32, (std::byte*)&v32 + sizeof(std::complex<float>));
     } else {
-      NVF_ERROR(
-          false,
+      NVF_THROW(
           "Cannot convert complex double to ",
           dtype,
           " type: only complex float and complex double are supported.");
@@ -350,11 +347,8 @@ std::vector<std::byte> polymorphicValueToBytes(
     // FUSER_PERF_SCOPE("polymorphicValueToBytes(Opaque)");
     return argument.as<Opaque>().bytes();
   } else {
-    NVF_ERROR(
-        false,
-        "Cannot convert ",
-        argument.type().name(),
-        " to kernel argument data.");
+    NVF_THROW(
+        "Cannot convert ", argument.type().name(), " to kernel argument data.");
   }
 }
 

csrc/fusion_executor/executor_params.cpp

@@ -69,10 +69,8 @@ void LaunchParams::bind(int64_t val, ParallelType p_type) {
       checkAndSet(val, gdimz_, "gridDim.z");
       break;
     default:
-      NVF_ERROR(
-          false,
-          "Tried to bind invalid parallel type in launch config: ",
-          p_type);
+      NVF_THROW(
+          "Tried to bind invalid parallel type in launch config: ", p_type);
   }
   assertValid();
 }
@@ -92,10 +90,8 @@ int64_t LaunchParams::getDim(ParallelType p_type) const {
     case ParallelType::BIDz:
       return gdimz();
     default:
-      NVF_ERROR(
-          false,
-          "Tried to get with invalid parallel type in launch config: ",
-          p_type);
+      NVF_THROW(
+          "Tried to get with invalid parallel type in launch config: ", p_type);
   }
 }
 
@@ -118,10 +114,8 @@ const int64_t& LaunchParams::getRawVal(ParallelType p_type) const {
     case ParallelType::BIDz:
       return gdimz_;
     default:
-      NVF_ERROR(
-          false,
-          "Tried to get with invalid parallel type in launch config: ",
-          p_type);
+      NVF_THROW(
+          "Tried to get with invalid parallel type in launch config: ", p_type);
   }
 }
 

csrc/fusion_executor/executor_utils.cpp

@@ -317,8 +317,7 @@ std::unique_ptr<caching::VectorizedTensorInfo> getVectorizedTensorValidationInfo
         vectorized_tensor_info_ptr->global_inp_misaligned_tv.insert(
             producer_tv);
       } else {
-        NVF_ERROR(
-            false,
+        NVF_THROW(
             "Unsupported memory configuration for misaligned vectorization.");
       }
     }
@@ -850,8 +849,7 @@ class NvrtcCompileDriver {
     if (result != NVRTC_SUCCESS) {
       // Print CUDA starting at first global function
       size_t kernel_start = src.find("__global__");
-      NVF_ERROR(
-          false,
+      NVF_THROW(
           "\n",
           src.substr(kernel_start),
           "\nCUDA NVRTC compile error: ",

csrc/id_model/circular_buffer_indexing.cpp

@@ -174,10 +174,8 @@ CircularBufferLoopStage getCircularBufferLoopStage(
     }
   }
 
-  NVF_ERROR(
-      false,
-      "Circular buffer loop not found for ",
-      circular_buffer_tv->toString());
+  NVF_THROW(
+      "Circular buffer loop not found for ", circular_buffer_tv->toString());
 }
 
 } // namespace nvfuser

csrc/id_model/indexing.cpp

@@ -741,8 +741,7 @@ ParallelType getParallelType(const ValGroup& loop_group) {
       common_pt = pt;
     } else {
       // Inconsistent parallelization
-      NVF_ERROR(
-          false,
+      NVF_THROW(
           "Inconsistent parallelization detected. ",
           "Known type: ",
           common_pt,

csrc/index_compute.cpp

@@ -2174,8 +2174,7 @@ kir::TensorIndex* Index::getProducerIndex(
       } else if (items_per_thread == 4) {
         op = UnaryOpType::AdjustPartialLdMatrixAddrInTuring16;
       } else {
-        NVF_ERROR(
-            false,
+        NVF_THROW(
             "Unexpected output vectorizaiton for ldmatrix, expect 2, 4, or 8, get ",
             items_per_thread);
       }

csrc/ir/base_nodes.cpp

@@ -60,16 +60,12 @@ bool Statement::lessThan(const Statement* stmt1, const Statement* stmt2) {
 }
 
 std::string Statement::toString(int indent_size) const {
-  NVF_ERROR(
-      false, "toString for IR node ", typeid(*this).name(), " is not defined");
+  NVF_THROW("toString for IR node ", typeid(*this).name(), " is not defined");
 }
 
 std::string Statement::toInlineString(int indent_size) const {
-  NVF_ERROR(
-      false,
-      "toInlineString for IR node ",
-      typeid(*this).name(),
-      " is not defined");
+  NVF_THROW(
+      "toInlineString for IR node ", typeid(*this).name(), " is not defined");
 }
 
 Fusion* Statement::fusion() const {
@@ -386,8 +382,7 @@ Expr* Expr::withWritePredicate(kir::Predicate* predicate) {
 std::vector<PolymorphicValue> Expr::evaluate(
     const ExpressionEvaluator& ee,
     const std::vector<PolymorphicValue>& inputs) const {
-  NVF_ERROR(
-      false,
+  NVF_THROW(
       "`evaluate` method for expression ",
       getOpString(),
       " is not defined. ",

csrc/ir/nodes.cpp

@@ -413,8 +413,7 @@ std::vector<PolymorphicValue> UnaryOp::evaluate(
       } else if (isComplexType(*out()->getDataType())) {
         return {PolymorphicValue((std::complex<double>)in)};
       } else {
-        NVF_ERROR(
-            false, "dtype not supported in evaluator: ", *out()->getDataType());
+        NVF_THROW("dtype not supported in evaluator: ", *out()->getDataType());
       }
     case UnaryOpType::Reciprocal:
       return {1.0 / in};
@@ -442,8 +441,7 @@ std::vector<PolymorphicValue> UnaryOp::evaluate(
       if (*out()->getDataType() == DataType::Float) {
         return {PolymorphicValue((double)*(float*)in)};
       } else {
-        NVF_ERROR(
-            false, "dtype not supported in evaluator: ", *out()->getDataType());
+        NVF_THROW("dtype not supported in evaluator: ", *out()->getDataType());
       }
       break;
     case UnaryOpType::Sigmoid:
@@ -1568,8 +1566,7 @@ int GroupedReductionOp::getExprIndexOfOutput(Val* output_val) const {
     return (int)std::distance(outputs().begin(), it);
   }
 
-  NVF_ERROR(
-      false, "Not an output, ", output_val->toString(), ", of ", toString());
+  NVF_THROW("Not an output, ", output_val->toString(), ", of ", toString());
 }
 
 std::vector<PolymorphicValue> GroupedReductionOp::evaluate(
@@ -1968,8 +1965,7 @@ int GroupedWelfordOp::getExprIndexOfOutput(Val* output_val) const {
     }
   }
 
-  NVF_ERROR(
-      false, "Not an output, ", output_val->toString(), ", of ", toString());
+  NVF_THROW("Not an output, ", output_val->toString(), ", of ", toString());
 }
 
 Val* GroupedWelfordOp::getInitValOfOutput(Val* output_val) const {