#include "Dialect/NVGPU/IR/Dialect.h"
#include "Dialect/NVWS/IR/Dialect.h"
#include "NVGPUToLLVM/Passes.h"
#include "TritonNVIDIAGPUToLLVM/Passes.h"
#include "cublas_instance.h"
#include "mlir/Pass/PassManager.h"
#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "nvidia/hopper/include/Transforms/Passes.h"
#include "nvidia/include/Dialect/NVWS/Transforms/Passes.h"
#include "passes.h"
#include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"
#include "triton/Dialect/TritonNvidiaGPU/Transforms/Passes.h"
#include "llvm/IR/Constants.h"
#include <dlfcn.h>
#include <nanobind/nanobind.h>
#include <nanobind/stl/optional.h>
#include <nanobind/stl/pair.h>
#include <nanobind/stl/string.h>
#include <nanobind/stl/vector.h>

namespace py = nanobind;
namespace ttng = mlir::triton::nvidia_gpu;

namespace {

using CUresult = int;
using CUdeviceptr = unsigned long long;

constexpr CUresult CUDA_SUCCESS = 0;

class CudaExampleUtils {
private:
  using cuMemAlloc_t = CUresult (*)(CUdeviceptr *, size_t);
  using cuMemFree_t = CUresult (*)(CUdeviceptr);
  using cuMemcpyHtoD_t = CUresult (*)(CUdeviceptr, const void *, size_t);
  using cuMemcpyDtoH_t = CUresult (*)(void *, CUdeviceptr, size_t);
  using cuCtxSynchronize_t = CUresult (*)();
  using cuGetErrorString_t = CUresult (*)(CUresult, const char **);

  void *dylibHandle = nullptr;
  cuMemAlloc_t cuMemAlloc = nullptr;
  cuMemFree_t cuMemFree = nullptr;
  cuMemcpyHtoD_t cuMemcpyHtoD = nullptr;
  cuMemcpyDtoH_t cuMemcpyDtoH = nullptr;
  cuCtxSynchronize_t cuCtxSynchronize = nullptr;
  cuGetErrorString_t cuGetErrorString = nullptr;

  template <typename T> T loadSymbol(const char *name) {
    dlerror();
    auto *symbol = dlsym(dylibHandle, name);
    if (const char *error = dlerror()) {
      throw std::runtime_error(
          "Could not load CUDA driver symbol `" + std::string(name) +
          "`. Initialize the NVIDIA runtime before using this helper: " +
          std::string(error));
    }
    return reinterpret_cast<T>(symbol);
  }

  void check(CUresult status, const char *call) {
    if (status == CUDA_SUCCESS) {
      return;
    }
    const char *error = nullptr;
    if (cuGetErrorString) {
      cuGetErrorString(status, &error);
    }
    throw std::runtime_error(
        "Triton Error [CUDA]: " + std::string(call) + " failed: " +
        (error ? std::string(error) : std::string("Unknown error")));
  }

public:
  static CudaExampleUtils &get() {
    static CudaExampleUtils instance;
    return instance;
  }

  CudaExampleUtils() {
    dylibHandle = dlopen("libcuda.so.1", RTLD_NOLOAD | RTLD_LAZY);
    if (dylibHandle == nullptr) {
      throw std::runtime_error(
          "Could not find an already-loaded `libcuda.so.1`. Initialize "
          "Triton's NVIDIA runtime before using this helper.");
    }
    cuMemAlloc = loadSymbol<cuMemAlloc_t>("cuMemAlloc_v2");
    cuMemFree = loadSymbol<cuMemFree_t>("cuMemFree_v2");
    cuMemcpyHtoD = loadSymbol<cuMemcpyHtoD_t>("cuMemcpyHtoD_v2");
    cuMemcpyDtoH = loadSymbol<cuMemcpyDtoH_t>("cuMemcpyDtoH_v2");
    cuCtxSynchronize = loadSymbol<cuCtxSynchronize_t>("cuCtxSynchronize");
    cuGetErrorString = loadSymbol<cuGetErrorString_t>("cuGetErrorString");
  }

  ~CudaExampleUtils() {
    if (dylibHandle) {
      dlclose(dylibHandle);
    }
  }

  uint64_t deviceMalloc(size_t size) {
    CUdeviceptr ptr = 0;
    check(cuMemAlloc(&ptr, size), "cuMemAlloc");
    return ptr;
  }

  void deviceFree(uint64_t ptr) {
    if (ptr == 0) {
      return;
    }
    check(cuMemFree(static_cast<CUdeviceptr>(ptr)), "cuMemFree");
  }

  void copyHostToDevice(uint64_t dst, const py::bytes &src) {
    check(cuMemcpyHtoD(static_cast<CUdeviceptr>(dst), src.c_str(), src.size()),
          "cuMemcpyHtoD");
  }

  py::bytes copyDeviceToHost(uint64_t src, size_t size) {
    std::string payload(size, '\0');
    check(cuMemcpyDtoH(payload.data(), static_cast<CUdeviceptr>(src), size),
          "cuMemcpyDtoH");
    return py::bytes(payload.c_str(), payload.size());
  }

  void synchronize() { check(cuCtxSynchronize(), "cuCtxSynchronize"); }
};

void init_triton_nvidia_passes_ttgpuir(py::module_ &m) {
  using namespace mlir::triton;
  // TODO: it is weird to pass mlir::triton::NVVM here since the conversion is
  // nvidia-specificontext
  m.def("add_allocate_shared_memory_nv",
        [](mlir::PassManager &pm, int32_t capability, int32_t ptxVersion) {
          pm.addPass(mlir::triton::createAllocateSharedMemoryNvPass(
              capability, ptxVersion));
        });
  m.def("add_to_llvmir",
        [](mlir::PassManager &pm, int32_t capability, int32_t ptxVersion,
           bool enableConcurrencySanitizer) {
          pm.addPass(mlir::triton::createConvertTritonGPUToLLVMPass(
              capability, ptxVersion, enableConcurrencySanitizer));
        });
}

std::unique_ptr<mlir::Pass>
createTritonGPUFenceInsertionWrapper(int32_t capability) {
  ttng::TritonGPUFenceInsertionOptions options;
  options.computeCapability = capability;
  return ttng::createTritonGPUFenceInsertion(options);
}

std::unique_ptr<mlir::Pass>
createTritonGPUProxyFenceInsertionWrapper(int32_t capability) {
  ttng::TritonGPUProxyFenceInsertionOptions options;
  options.computeCapability = capability;
  return ttng::createTritonGPUProxyFenceInsertion(options);
}

std::unique_ptr<mlir::Pass>
createInitializeWSClusterBarriersWrapper(int32_t capability,
                                         int32_t ptxVersion) {
  mlir::triton::InitializeWSClusterBarriersOptions options;
  options.computeCapability = capability;
  options.ptxVersion = ptxVersion;
  return mlir::triton::createInitializeWSClusterBarriers(options);
}

void init_triton_nvidia_passes_ttnvgpuir(py::module_ &m) {
  ADD_PASS_WRAPPER_0("add_plan_cta", ttng::createTritonNvidiaGPUPlanCTAPass);
  ADD_PASS_WRAPPER_0("add_to_clc", ttng::createTritonNvidiaGPUToCLCPass);
  ADD_PASS_WRAPPER_0("add_lower_clc", ttng::createTritonNvidiaGPULowerCLCPass);
  ADD_PASS_WRAPPER_1("add_fence_insertion",
                     createTritonGPUFenceInsertionWrapper, int32_t);
  ADD_PASS_WRAPPER_1("add_proxy_fence_insertion",
                     createTritonGPUProxyFenceInsertionWrapper, int32_t);
  ADD_PASS_WRAPPER_0("add_tmem_barrier_insertion",
                     ttng::createTritonNvidiaGPUTMemBarrierInsertionPass);
  ADD_PASS_WRAPPER_0("add_tmem_load_reduce",
                     ttng::createTritonNvidiaGPUFuseTMEMLoadReducePass);
  ADD_PASS_WRAPPER_0("add_tma_lowering",
                     ttng::createTritonNvidiaGPUTMALoweringPass);
  ADD_PASS_WRAPPER_0("add_promote_lhs_to_tmem",
                     ttng::createTritonNvidiaGPUPromoteLHSToTMemPass);
  ADD_PASS_WRAPPER_0("add_remove_tmem_tokens",
                     ttng::createTritonNvidiaGPURemoveTMEMTokensPass);
  ADD_PASS_WRAPPER_0("add_check_matmul_two_cta",
                     ttng::createTritonNvidiaGPUCheckMatmulTwoCTAPass);
  ADD_PASS_WRAPPER_0("add_nvgpu_to_llvm",
                     mlir::triton::createConvertNVGPUToLLVM);
  ADD_PASS_WRAPPER_2("add_initialize_ws_cluster_barriers",
                     createInitializeWSClusterBarriersWrapper, int32_t,
                     int32_t);
  ADD_PASS_WRAPPER_0("add_warp_specialize_to_llvm",
                     mlir::triton::createConvertWarpSpecializeToLLVM);
  ADD_PASS_WRAPPER_0("add_allocate_tensor_memory",
                     ttng::createTritonTensorMemoryAllocationPass);
  ADD_PASS_WRAPPER_0("add_lower_mma",
                     ttng::createTritonNvidiaGPUMMALoweringPass);
  ADD_PASS_WRAPPER_0("add_optimize_descriptor_encoding",
                     ttng::createTritonNvidiaGPUOptimizeDescriptorEncodingPass);
  ADD_PASS_WRAPPER_0("add_optimize_tmem_layouts",
                     ttng::createTritonNvidiaGPUOptimizeTMemLayoutsPass);
  ADD_PASS_WRAPPER_0("add_interleave_tmem",
                     ttng::createTritonNvidiaGPUInterleaveTMemPass);
  ttng::registerConSanNVIDIAHooks();
}

void init_triton_nvidia_passes_nvws(py::module_ &m) {
  ADD_PASS_WRAPPER_0("add_lower_warp_group",
                     mlir::triton::createNVWSLowerWarpGroup);
  ADD_PASS_WRAPPER_0("add_lower_aref", mlir::triton::createNVWSLowerAref);
  ADD_PASS_WRAPPER_0("add_assign_stage_phase",
                     mlir::triton::createNVWSAssignStagePhase);
  ADD_PASS_WRAPPER_0("add_insert_tmem_aref",
                     mlir::triton::createNVWSInsertTmemAref);
}

void init_triton_hopper_passes(py::module_ &m) {
  // Meta's autoWS
  ADD_PASS_OPTION_WRAPPER_2("add_hopper_warpspec",
                            mlir::createNVGPUWarpSpecialization, int, bool);
}

void checkMatmulConstraints(const std::string &A_dtype,
                            const std::string &B_dtype,
                            const std::string &C_dtype,
                            const std::vector<int> &A_shape,
                            const std::vector<int> &B_shape,
                            const std::vector<int> &C_shape) {
  if (A_dtype != B_dtype || A_dtype != C_dtype) {
    throw std::runtime_error("Data types do not match.");
  }
  if (A_dtype != "torch.float8_e4m3fn" && A_dtype != "torch.float16" &&
      A_dtype != "torch.float32" && A_dtype != "torch.bfloat16") {
    throw std::runtime_error("Unsupported data type.");
  }

  if (A_shape.size() != 2 || B_shape.size() != 2 || C_shape.size() != 2) {
    throw std::runtime_error("Only 2D matrices are supported.");
  }

  int k = A_shape[1];
  if (k != B_shape[1]) {
    throw std::runtime_error(
        "Matrix dimensions do not match. A is [" + std::to_string(A_shape[0]) +
        ", " + std::to_string(A_shape[1]) + "], B is [" +
        std::to_string(B_shape[0]) + ", " + std::to_string(B_shape[1]) +
        "]. Expected A.shape[1] == B.shape[1]. Note "
        "that B needs to be transposed.");
  }

  int m = A_shape[0];
  if (m != C_shape[0]) {
    throw std::runtime_error(
        "Matrix dimensions do not match. A is [" + std::to_string(A_shape[0]) +
        ", " + std::to_string(A_shape[1]) + "], C is [" +
        std::to_string(C_shape[0]) + ", " + std::to_string(C_shape[1]) +
        "]. Expected A.shape[0] == C.shape[0].");
  }

  int n = B_shape[0];
  if (n != C_shape[1]) {
    throw std::runtime_error(
        "Matrix dimensions do not match. B is [" + std::to_string(B_shape[0]) +
        ", " + std::to_string(B_shape[1]) + "], C is [" +
        std::to_string(C_shape[0]) + ", " + std::to_string(C_shape[1]) +
        "]. Expected B.shape[0] == C.shape[1]. Note "
        "that B needs to be transposed.");
  }
}

} // namespace

void init_triton_nvidia(py::module_ &m) {
  auto passes = m.def_submodule("passes");
  auto nvws_m = passes.def_submodule("nvws");
  init_triton_nvidia_passes_nvws(nvws_m);
  auto ttgpuir_m = passes.def_submodule("ttgpuir");
  init_triton_nvidia_passes_ttgpuir(ttgpuir_m);
  auto ttnvgpuir_m = passes.def_submodule("ttnvgpuir");
  init_triton_nvidia_passes_ttnvgpuir(ttnvgpuir_m);
  auto hopper_m = passes.def_submodule("hopper");
  init_triton_hopper_passes(hopper_m);

  // load dialects
  m.def("load_dialects", [](mlir::MLIRContext &context) {
    mlir::DialectRegistry registry;
    registry.insert<mlir::triton::nvidia_gpu::TritonNvidiaGPUDialect,
                    mlir::triton::nvgpu::NVGPUDialect,
                    mlir::triton::nvws::NVWSDialect>();
    mlir::registerNVVMDialectTranslation(registry);
    context.appendDialectRegistry(registry);
    context.loadAllAvailableDialects();
  });

  // Set short point option, this needs to be set before setting the data
  // layout.
  m.def("set_short_ptr", []() {
    auto options = llvm::cl::getRegisteredOptions();
    const char *flag = "nvptx-short-ptr";
    auto *shortPtr = static_cast<llvm::cl::opt<bool> *>(options[flag]);
    assert(shortPtr);
    shortPtr->setValue(true);
  });

  // TODO: could be done in python if we had a generic interface to set metadata
  m.def("set_nvvm_reflect_ftz", [](llvm::Module *mod) {
    // please check https://llvm.org/docs/NVPTXUsage.html#reflection-parameters
    // this will enable fast math path in libdevice
    // for example, when enable nvvm-reflect-ftz, sqrt.approx.f32 will change to
    // sqrt.approx.ftz.f32
    using namespace llvm;
    auto &ctx = mod->getContext();
    Type *i32 = Type::getInt32Ty(ctx);
    auto *mdFour = ConstantAsMetadata::get(ConstantInt::getSigned(i32, 4));
    auto *mdName = MDString::get(ctx, "nvvm-reflect-ftz");
    auto *mdOne = ConstantAsMetadata::get(ConstantInt::getSigned(i32, 1));
    auto *reflect = MDNode::get(ctx, {mdFour, mdName, mdOne});
    mod->addModuleFlag(reflect);
  });

  m.def("device_malloc",
        [](size_t size) { return CudaExampleUtils::get().deviceMalloc(size); });
  m.def("device_free",
        [](uint64_t ptr) { CudaExampleUtils::get().deviceFree(ptr); });
  m.def("copy_host_to_device", [](uint64_t dst, const py::bytes &src) {
    CudaExampleUtils::get().copyHostToDevice(dst, src);
  });
  m.def("copy_device_to_host", [](uint64_t src, size_t size) {
    return CudaExampleUtils::get().copyDeviceToHost(src, size);
  });
  m.def("synchronize", []() { CudaExampleUtils::get().synchronize(); });

  // cublas
  auto cublas = m.def_submodule("cublas");

  py::class_<CublasLtInstance>(cublas, "CublasLt")
      .def(py::new_([](py::object workspace) {
        auto wrk_ptr = py::cast<uint64_t>(workspace.attr("data_ptr")());
        auto wrk_size = py::cast<size_t>(workspace.attr("numel")()) *
                        py::cast<size_t>(workspace.attr("element_size")());
        return new CublasLtInstance(wrk_ptr, wrk_size);
      }))
      .def("matmul",
           [](CublasLtInstance &self, py::object &A, py::object &B,
              py::object &C) {
             auto A_ptr = py::cast<uint64_t>(A.attr("data_ptr")());
             auto B_ptr = py::cast<uint64_t>(B.attr("data_ptr")());
             auto C_ptr = py::cast<uint64_t>(C.attr("data_ptr")());

             auto A_shape = py::cast<std::vector<int>>(A.attr("shape"));
             auto B_shape = py::cast<std::vector<int>>(B.attr("shape"));
             auto C_shape = py::cast<std::vector<int>>(C.attr("shape"));

             auto A_dtype =
                 py::cast<std::string>(A.attr("dtype").attr("__str__")());
             auto B_dtype =
                 py::cast<std::string>(B.attr("dtype").attr("__str__")());
             auto C_dtype =
                 py::cast<std::string>(C.attr("dtype").attr("__str__")());

             checkMatmulConstraints(A_dtype, B_dtype, C_dtype, A_shape, B_shape,
                                    C_shape);

             std::string dtype_str =
                 A_dtype.substr(A_dtype.find_last_of('.') + 1);
             cudaDataType_t dtype;
             if (dtype_str == "float8_e4m3fn") {
               dtype = CUDA_R_8F_E4M3;
             } else if (dtype_str == "float16") {
               dtype = CUDA_R_16F;
             } else if (dtype_str == "float32") {
               // Use FP32 inputs with TF32 compute in cublasLt (set in compute
               // type)
               dtype = CUDA_R_32F;
             } else if (dtype_str == "bfloat16") {
               dtype = CUDA_R_16BF;
             } else {
               throw std::runtime_error(
                   "Unsupported dtype for cublasLt.matmul: " + dtype_str);
             }

             self.matmul(A_shape[0], B_shape[0], A_shape[1], A_ptr, B_ptr,
                         C_ptr, dtype);
           })
      .def("gemm",
           [](CublasLtInstance &self, py::object &A, py::object &B,
              py::object &C, py::object &D, float alpha, float beta) {
             auto A_ptr = py::cast<uint64_t>(A.attr("data_ptr")());
             auto B_ptr = py::cast<uint64_t>(B.attr("data_ptr")());
             auto C_ptr = py::cast<uint64_t>(C.attr("data_ptr")());
             auto D_ptr = py::cast<uint64_t>(D.attr("data_ptr")());

             auto A_shape = py::cast<std::vector<int>>(A.attr("shape"));
             auto B_shape = py::cast<std::vector<int>>(B.attr("shape"));
             auto C_shape = py::cast<std::vector<int>>(C.attr("shape"));
             auto D_shape = py::cast<std::vector<int>>(D.attr("shape"));

             auto A_dtype =
                 py::cast<std::string>(A.attr("dtype").attr("__str__")());
             auto B_dtype =
                 py::cast<std::string>(B.attr("dtype").attr("__str__")());
             auto C_dtype =
                 py::cast<std::string>(C.attr("dtype").attr("__str__")());
             auto D_dtype =
                 py::cast<std::string>(D.attr("dtype").attr("__str__")());

             checkMatmulConstraints(A_dtype, B_dtype, D_dtype, A_shape, B_shape,
                                    D_shape);
             if (C_dtype != "torch.float16") {
               throw std::runtime_error("C dtype must be float16, got " +
                                        C_dtype);
             }
             if (C_shape != D_shape) {
               throw std::runtime_error("C and D shapes must match");
             }

             std::string dtype_str =
                 A_dtype.substr(A_dtype.find_last_of('.') + 1);
             cudaDataType_t dtype;
             if (dtype_str == "float8_e4m3fn") {
               dtype = CUDA_R_8F_E4M3;
             } else if (dtype_str == "float16") {
               dtype = CUDA_R_16F;
             } else if (dtype_str == "float32") {
               dtype = CUDA_R_32F;
             } else if (dtype_str == "bfloat16") {
               dtype = CUDA_R_16BF;
             } else {
               throw std::runtime_error(
                   "Unsupported dtype for cublasLt.gemm: " + dtype_str);
             }

             self.gemm(A_shape[0], B_shape[0], A_shape[1], A_ptr, B_ptr, C_ptr,
                       D_ptr, dtype, alpha, beta);
           })
      .def("block_scaled_matmul_mxfp8",
           [](CublasLtInstance &self, py::object &A, py::object &B,
              py::object &output, py::object &scale_A, py::object &scale_B) {
             auto A_ptr = py::cast<uint64_t>(A.attr("data_ptr")());
             auto B_ptr = py::cast<uint64_t>(B.attr("data_ptr")());
             auto output_ptr = py::cast<uint64_t>(output.attr("data_ptr")());
             auto scale_A_ptr = py::cast<uint64_t>(scale_A.attr("data_ptr")());
             auto scale_B_ptr = py::cast<uint64_t>(scale_B.attr("data_ptr")());

             auto A_shape = py::cast<std::vector<int>>(A.attr("shape"));
             auto B_shape = py::cast<std::vector<int>>(B.attr("shape"));

             auto A_dtype =
                 py::cast<std::string>(A.attr("dtype").attr("__str__")());
             auto B_dtype =
                 py::cast<std::string>(B.attr("dtype").attr("__str__")());
             auto output_dtype =
                 py::cast<std::string>(output.attr("dtype").attr("__str__")());

             // Only support MXFP8: FP8 E4M3 inputs, FP16 output
             if (A_dtype != "torch.float8_e4m3fn" ||
                 B_dtype != "torch.float8_e4m3fn") {
               throw std::runtime_error(
                   "block_scaled_matmul_mxfp8 only supports float8_e4m3fn "
                   "inputs (MXFP8)");
             }

             if (output_dtype != "torch.float16") {
               throw std::runtime_error(
                   "block_scaled_matmul_mxfp8 output must be float16, got " +
                   output_dtype);
             }

             int K = A_shape[1];

             self.block_scaled_matmul_mxfp8(A_shape[0], B_shape[0], K, A_ptr,
                                            B_ptr, output_ptr, scale_A_ptr,
                                            scale_B_ptr);
           })
      .def("block_scaled_matmul_nvfp4", [](CublasLtInstance &self,
                                           py::object &A, py::object &B,
                                           py::object &output,
                                           py::object &scale_A,
                                           py::object &scale_B) {
        auto A_ptr = py::cast<uint64_t>(A.attr("data_ptr")());
        auto B_ptr = py::cast<uint64_t>(B.attr("data_ptr")());
        auto output_ptr = py::cast<uint64_t>(output.attr("data_ptr")());
        auto scale_A_ptr = py::cast<uint64_t>(scale_A.attr("data_ptr")());
        auto scale_B_ptr = py::cast<uint64_t>(scale_B.attr("data_ptr")());

        auto A_shape = py::cast<std::vector<int>>(A.attr("shape"));
        auto B_shape = py::cast<std::vector<int>>(B.attr("shape"));

        auto A_dtype = py::cast<std::string>(A.attr("dtype").attr("__str__")());
        auto B_dtype = py::cast<std::string>(B.attr("dtype").attr("__str__")());
        auto output_dtype =
            py::cast<std::string>(output.attr("dtype").attr("__str__")());

        // NVFP4: uint8 packed FP4 inputs (2 elements per byte), FP8 E4M3
        // scales, FP16 output
        if (A_dtype != "torch.uint8" || B_dtype != "torch.uint8") {
          throw std::runtime_error("block_scaled_matmul_nvfp4 only supports "
                                   "uint8 packed FP4 inputs (NVFP4), got A=" +
                                   A_dtype + ", B=" + B_dtype);
        }

        if (output_dtype != "torch.float16") {
          throw std::runtime_error(
              "block_scaled_matmul_nvfp4 output must be float16, got " +
              output_dtype);
        }

        // For packed FP4, shape[1] is in bytes, but K dimension should be in
        // elements So K = A_shape[1] * 2 (2 elements per byte)
        int K = A_shape[1] * 2;
        if (B_shape[1] * 2 != K) {
          throw std::runtime_error("K dimensions must match. A has " +
                                   std::to_string(K) + " elements, B has " +
                                   std::to_string(B_shape[1] * 2) +
                                   " elements");
        }

        self.block_scaled_matmul_nvfp4(A_shape[0], B_shape[0], K, A_ptr, B_ptr,
                                       output_ptr, scale_A_ptr, scale_B_ptr);
      });

  m.def("has_extern_deps", [](llvm::Module *dstMod) -> bool {
    // `global_smem` is special cased in Triton, so we ignore it here.
    for (const auto &g : dstMod->globals()) {
      if (g.hasExternalLinkage() && g.getName() != "global_smem") {
        return true;
      }
    }
    for (const auto &f : *dstMod) {
      if (f.hasExternalLinkage() && !f.hasExactDefinition() &&
          !f.isIntrinsic()) {
        return true;
      }
    }
    return false;
  });
}
