diff --git a/cpp/doxygen/main_page.md b/cpp/doxygen/main_page.md
index 497fb3e13e..22bab7c861 100644
--- a/cpp/doxygen/main_page.md
+++ b/cpp/doxygen/main_page.md
@@ -76,14 +76,19 @@ Then run the example:
 ## Runtime Settings
 
 #### Compatibility Mode (KVIKIO_COMPAT_MODE)
-When KvikIO is running in compatibility mode, it doesn't load `libcufile.so`. Instead, reads and writes are done using POSIX. Notice, this is not the same as the compatibility mode in cuFile. That is cuFile can run in compatibility mode while KvikIO is not.
+When KvikIO is running in compatibility mode, it doesn't load `libcufile.so`. Instead, reads and writes are done using POSIX. Notice, this is not the same as the compatibility mode in cuFile. It is possible that KvikIO performs I/O in the non-compatibility mode by using the cuFile library, but the cuFile library itself is configured to operate in its own compatibility mode. For more details, refer to [cuFile compatibility mode](https://docs.nvidia.com/gpudirect-storage/api-reference-guide/index.html#cufile-compatibility-mode) and [cuFile environment variables](https://docs.nvidia.com/gpudirect-storage/troubleshooting-guide/index.html#environment-variables)
 
-Set the environment variable `KVIKIO_COMPAT_MODE` to enable/disable compatibility mode. By default, compatibility mode is enabled:
+The environment variable `KVIKIO_COMPAT_MODE` has three options (case-insensitive):
+  - `ON` (aliases: `TRUE`, `YES`, `1`): Enable the compatibility mode.
+  - `OFF` (aliases: `FALSE`, `NO`, `0`): Disable the compatibility mode, and enforce cuFile I/O. GDS will be activated if the system requirements for cuFile are met and cuFile is properly configured. However, if the system is not suited for cuFile, I/O operations under the `OFF` option may error out, crash or hang.
+  - `AUTO`: Try cuFile I/O first, and fall back to POSIX I/O if the system requirements for cuFile are not met.
+
+Under `AUTO`, KvikIO falls back to the compatibility mode:
   - when `libcufile.so` cannot be found.
   - when running in Windows Subsystem for Linux (WSL).
   - when `/run/udev` isn't readable, which typically happens when running inside a docker image not launched with `--volume /run/udev:/run/udev:ro`.
 
-This setting can also be controlled by `defaults::compat_mode()` and `defaults::compat_mode_reset()`.
+This setting can also be programmatically controlled by `defaults::set_compat_mode()` and `defaults::compat_mode_reset()`.
 
 
 #### Thread Pool (KVIKIO_NTHREADS)
diff --git a/cpp/examples/basic_io.cpp b/cpp/examples/basic_io.cpp
index 1eabd8fdee..4d04391404 100644
--- a/cpp/examples/basic_io.cpp
+++ b/cpp/examples/basic_io.cpp
@@ -65,7 +65,7 @@ int main()
   check(cudaSetDevice(0) == cudaSuccess);
 
   cout << "KvikIO defaults: " << endl;
-  if (kvikio::defaults::compat_mode()) {
+  if (kvikio::defaults::is_compat_mode_preferred()) {
     cout << "  Compatibility mode: enabled" << endl;
   } else {
     kvikio::DriverInitializer manual_init_driver;
@@ -181,7 +181,7 @@ int main()
     cout << "Parallel POSIX read (" << kvikio::defaults::thread_pool_nthreads()
          << " threads): " << read << endl;
   }
-  if (kvikio::is_batch_and_stream_available() && !kvikio::defaults::compat_mode()) {
+  if (kvikio::is_batch_and_stream_available() && !kvikio::defaults::is_compat_mode_preferred()) {
     std::cout << std::endl;
     Timer timer;
     // Here we use the batch API to read "/tmp/test-file" into `b_dev` by
diff --git a/cpp/examples/basic_no_cuda.cpp b/cpp/examples/basic_no_cuda.cpp
index 0d79a52883..42ecb7142d 100644
--- a/cpp/examples/basic_no_cuda.cpp
+++ b/cpp/examples/basic_no_cuda.cpp
@@ -41,7 +41,7 @@ constexpr int LARGE_SIZE = 8 * SIZE;             // LARGE SIZE to test partial s
 int main()
 {
   cout << "KvikIO defaults: " << endl;
-  if (kvikio::defaults::compat_mode()) {
+  if (kvikio::defaults::is_compat_mode_preferred()) {
     cout << "  Compatibility mode: enabled" << endl;
   } else {
     kvikio::DriverInitializer manual_init_driver;
diff --git a/cpp/include/kvikio/batch.hpp b/cpp/include/kvikio/batch.hpp
index 9c58a50b1d..7eebbd4df0 100644
--- a/cpp/include/kvikio/batch.hpp
+++ b/cpp/include/kvikio/batch.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -118,7 +118,7 @@ class BatchHandle {
     std::vector<CUfileIOParams_t> io_batch_params;
     io_batch_params.reserve(operations.size());
     for (const auto& op : operations) {
-      if (op.file_handle.is_compat_mode_on()) {
+      if (op.file_handle.is_compat_mode_preferred()) {
         throw CUfileException("Cannot submit a FileHandle opened in compatibility mode");
       }
 
diff --git a/cpp/include/kvikio/buffer.hpp b/cpp/include/kvikio/buffer.hpp
index c0aa7f9fbc..85c60b3f90 100644
--- a/cpp/include/kvikio/buffer.hpp
+++ b/cpp/include/kvikio/buffer.hpp
@@ -49,7 +49,7 @@ inline void buffer_register(const void* devPtr_base,
                             int flags                                = 0,
                             const std::vector<int>& errors_to_ignore = std::vector<int>())
 {
-  if (defaults::compat_mode()) { return; }
+  if (defaults::is_compat_mode_preferred()) { return; }
   CUfileError_t status = cuFileAPI::instance().BufRegister(devPtr_base, size, flags);
   if (status.err != CU_FILE_SUCCESS) {
     // Check if `status.err` is in `errors_to_ignore`
@@ -67,7 +67,7 @@ inline void buffer_register(const void* devPtr_base,
  */
 inline void buffer_deregister(const void* devPtr_base)
 {
-  if (defaults::compat_mode()) { return; }
+  if (defaults::is_compat_mode_preferred()) { return; }
   CUFILE_TRY(cuFileAPI::instance().BufDeregister(devPtr_base));
 }
 
diff --git a/cpp/include/kvikio/defaults.hpp b/cpp/include/kvikio/defaults.hpp
index 32367686d2..91071cbb28 100644
--- a/cpp/include/kvikio/defaults.hpp
+++ b/cpp/include/kvikio/defaults.hpp
@@ -13,6 +13,11 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
+
+/**
+ * @file
+ */
+
 #pragma once
 
 #include <algorithm>
@@ -27,7 +32,48 @@
 #include <kvikio/shim/cufile.hpp>
 
 namespace kvikio {
+/**
+ * @brief I/O compatibility mode.
+ */
+enum class CompatMode : uint8_t {
+  OFF,  ///< Enforce cuFile I/O. GDS will be activated if the system requirements for cuFile are met
+        ///< and cuFile is properly configured. However, if the system is not suited for cuFile, I/O
+        ///< operations under the OFF option may error out, crash or hang.
+  ON,   ///< Enforce POSIX I/O.
+  AUTO,  ///< Try cuFile I/O first, and fall back to POSIX I/O if the system requirements for cuFile
+         ///< are not met.
+};
+
 namespace detail {
+/**
+ * @brief Parse a string into a CompatMode enum.
+ *
+ * @param compat_mode_str Compatibility mode in string format(case-insensitive). Valid values
+ * include:
+ *   - `ON` (alias: `TRUE`, `YES`, `1`)
+ *   - `OFF` (alias: `FALSE`, `NO`, `0`)
+ *   - `AUTO`
+ * @return A CompatMode enum.
+ */
+inline CompatMode parse_compat_mode_str(std::string_view compat_mode_str)
+{
+  // Convert to lowercase
+  std::string tmp{compat_mode_str};
+  std::transform(
+    tmp.begin(), tmp.end(), tmp.begin(), [](unsigned char c) { return std::tolower(c); });
+
+  CompatMode res{};
+  if (tmp == "on" || tmp == "true" || tmp == "yes" || tmp == "1") {
+    res = CompatMode::ON;
+  } else if (tmp == "off" || tmp == "false" || tmp == "no" || tmp == "0") {
+    res = CompatMode::OFF;
+  } else if (tmp == "auto") {
+    res = CompatMode::AUTO;
+  } else {
+    throw std::invalid_argument("Unknown compatibility mode: " + std::string{tmp});
+  }
+  return res;
+}
 
 template <typename T>
 T getenv_or(std::string_view env_var_name, T default_val)
@@ -77,16 +123,24 @@ inline bool getenv_or(std::string_view env_var_name, bool default_val)
                               std::string{env_val});
 }
 
+template <>
+inline CompatMode getenv_or(std::string_view env_var_name, CompatMode default_val)
+{
+  auto* env_val = std::getenv(env_var_name.data());
+  if (env_val == nullptr) { return default_val; }
+  return parse_compat_mode_str(env_val);
+}
+
 }  // namespace detail
 
 /**
- * @brief Singleton class of default values used thoughtout KvikIO.
+ * @brief Singleton class of default values used throughout KvikIO.
  *
  */
 class defaults {
  private:
   BS::thread_pool _thread_pool{get_num_threads_from_env()};
-  bool _compat_mode;
+  CompatMode _compat_mode;
   std::size_t _task_size;
   std::size_t _gds_threshold;
   std::size_t _bounce_buffer_size;
@@ -104,13 +158,7 @@ class defaults {
   {
     // Determine the default value of `compat_mode`
     {
-      if (std::getenv("KVIKIO_COMPAT_MODE") != nullptr) {
-        // Setting `KVIKIO_COMPAT_MODE` take precedence
-        _compat_mode = detail::getenv_or("KVIKIO_COMPAT_MODE", false);
-      } else {
-        // If `KVIKIO_COMPAT_MODE` isn't set, we infer based on runtime environment
-        _compat_mode = !is_cufile_available();
-      }
+      _compat_mode = detail::getenv_or("KVIKIO_COMPAT_MODE", CompatMode::AUTO);
     }
     // Determine the default value of `task_size`
     {
@@ -163,19 +211,77 @@ class defaults {
    *  - when `/run/udev` isn't readable, which typically happens when running inside a docker
    *    image not launched with `--volume /run/udev:/run/udev:ro`
    *
-   * @return The boolean answer
+   * @return Compatibility mode.
+   */
+  [[nodiscard]] static CompatMode compat_mode() { return instance()->_compat_mode; }
+
+  /**
+   * @brief Reset the value of `kvikio::defaults::compat_mode()`.
+   *
+   * Changing the compatibility mode affects all the new FileHandles whose `compat_mode` argument is
+   * not explicitly set, but it never affects existing FileHandles.
+   *
+   * @param compat_mode Compatibility mode.
+   */
+  static void compat_mode_reset(CompatMode compat_mode) { instance()->_compat_mode = compat_mode; }
+
+  /**
+   * @brief Infer the `AUTO` compatibility mode from the system runtime.
+   *
+   * If the requested compatibility mode is `AUTO`, set the expected compatibility mode to
+   * `ON` or `OFF` by performing a system config check; otherwise, do nothing. Effectively, this
+   * function reduces the requested compatibility mode from three possible states
+   * (`ON`/`OFF`/`AUTO`) to two (`ON`/`OFF`) so as to determine the actual I/O path. This function
+   * is lightweight as the inferred result is cached.
+   */
+  static CompatMode infer_compat_mode_if_auto(CompatMode compat_mode)
+  {
+    if (compat_mode == CompatMode::AUTO) {
+      static auto inferred_compat_mode_for_auto = []() -> CompatMode {
+        return is_cufile_available() ? CompatMode::OFF : CompatMode::ON;
+      }();
+      return inferred_compat_mode_for_auto;
+    }
+    return compat_mode;
+  }
+
+  /**
+   * @brief Given a requested compatibility mode, whether it is expected to reduce to `ON`.
+   *
+   * This function returns true if any of the two condition is satisfied:
+   *   - The compatibility mode is `ON`.
+   *   - It is `AUTO` but inferred to be `ON`.
+   *
+   * Conceptually, the opposite of this function is whether requested compatibility mode is expected
+   * to be `OFF`, which would occur if any of the two condition is satisfied:
+   *   - The compatibility mode is `OFF`.
+   *   - It is `AUTO` but inferred to be `OFF`.
+   *
+   * @param compat_mode Compatibility mode.
+   * @return Boolean answer.
    */
-  [[nodiscard]] static bool compat_mode() { return instance()->_compat_mode; }
+  static bool is_compat_mode_preferred(CompatMode compat_mode)
+  {
+    return compat_mode == CompatMode::ON ||
+           (compat_mode == CompatMode::AUTO &&
+            defaults::infer_compat_mode_if_auto(compat_mode) == CompatMode::ON);
+  }
 
   /**
-   * @brief Reset the value of `kvikio::defaults::compat_mode()`
+   * @brief Whether the global compatibility mode from class defaults is expected to be `ON`.
+   *
+   * This function returns true if any of the two condition is satisfied:
+   *   - The compatibility mode is `ON`.
+   *   - It is `AUTO` but inferred to be `ON`.
    *
-   * Changing compatibility mode, effects all new FileHandles that doesn't sets the
-   * `compat_mode` argument explicitly but it never effect existing FileHandles.
+   * Conceptually, the opposite of this function is whether the global compatibility mode is
+   * expected to be `OFF`, which would occur if any of the two condition is satisfied:
+   *   - The compatibility mode is `OFF`.
+   *   - It is `AUTO` but inferred to be `OFF`.
    *
-   * @param enable Whether to enable compatibility mode or not.
+   * @return Boolean answer.
    */
-  static void compat_mode_reset(bool enable) { instance()->_compat_mode = enable; }
+  static bool is_compat_mode_preferred() { return is_compat_mode_preferred(compat_mode()); }
 
   /**
    * @brief Get the default thread pool.
diff --git a/cpp/include/kvikio/error.hpp b/cpp/include/kvikio/error.hpp
index e84ebd770c..2ecd37b0b3 100644
--- a/cpp/include/kvikio/error.hpp
+++ b/cpp/include/kvikio/error.hpp
@@ -45,8 +45,8 @@ struct CUfileException : public std::runtime_error {
     if (error != CUDA_SUCCESS) {                                                               \
       const char* err_name     = nullptr;                                                      \
       const char* err_str      = nullptr;                                                      \
-      CUresult err_name_status = cudaAPI::instance().GetErrorName(error, &err_name);           \
-      CUresult err_str_status  = cudaAPI::instance().GetErrorString(error, &err_str);          \
+      CUresult err_name_status = kvikio::cudaAPI::instance().GetErrorName(error, &err_name);   \
+      CUresult err_str_status  = kvikio::cudaAPI::instance().GetErrorString(error, &err_str);  \
       if (err_name_status == CUDA_ERROR_INVALID_VALUE) { err_name = "unknown"; }               \
       if (err_str_status == CUDA_ERROR_INVALID_VALUE) { err_str = "unknown"; }                 \
       throw(_exception_type){std::string{"CUDA error at: "} + __FILE__ + ":" +                 \
diff --git a/cpp/include/kvikio/file_handle.hpp b/cpp/include/kvikio/file_handle.hpp
index 141c17371a..abc6660de6 100644
--- a/cpp/include/kvikio/file_handle.hpp
+++ b/cpp/include/kvikio/file_handle.hpp
@@ -20,6 +20,7 @@
 
 #include <cstddef>
 #include <cstdlib>
+#include <stdexcept>
 #include <system_error>
 #include <utility>
 
@@ -46,10 +47,40 @@ class FileHandle {
   int _fd_direct_on{-1};
   int _fd_direct_off{-1};
   bool _initialized{false};
-  bool _compat_mode{false};
+  CompatMode _compat_mode{CompatMode::AUTO};
   mutable std::size_t _nbytes{0};  // The size of the underlying file, zero means unknown.
   CUfileHandle_t _handle{};
 
+  /**
+   * @brief Given a requested compatibility mode, whether it is expected to reduce to `ON` for
+   * asynchronous I/O.
+   *
+   * @param requested_compat_mode Requested compatibility mode.
+   * @return True if POSIX I/O fallback will be used; false for cuFile I/O.
+   * @exception std::runtime_error When the requested compatibility mode is `OFF`, but cuFile
+   * batch/stream library symbol is missing, or cuFile configuration file is missing.
+   */
+  bool is_compat_mode_preferred_for_async(CompatMode requested_compat_mode)
+  {
+    if (!defaults::is_compat_mode_preferred(requested_compat_mode)) {
+      if (!is_batch_and_stream_available()) {
+        if (requested_compat_mode == CompatMode::AUTO) { return true; }
+        throw std::runtime_error("Missing cuFile batch or stream library symbol.");
+      }
+
+      // When checking for availability, we also check if cuFile's config file exist. This is
+      // because even when the stream API is available, it doesn't work if no config file exist.
+      if (config_path().empty()) {
+        if (requested_compat_mode == CompatMode::AUTO) { return true; }
+        throw std::runtime_error("Missing cuFile configuration file.");
+      }
+
+      return false;
+    }
+
+    return true;
+  }
+
  public:
   static constexpr mode_t m644 = S_IRUSR | S_IWUSR | S_IRGRP | S_IWGRP | S_IROTH;
   FileHandle() noexcept        = default;
@@ -68,12 +99,12 @@ class FileHandle {
    *   "a" -> "open for writing, appending to the end of file if it exists"
    *   "+" -> "open for updating (reading and writing)"
    * @param mode Access modes (see `open(2)`).
-   * @param compat_mode Enable KvikIO's compatibility mode for this file.
+   * @param compat_mode Set KvikIO's compatibility mode for this file.
    */
   FileHandle(const std::string& file_path,
              const std::string& flags = "r",
              mode_t mode              = m644,
-             bool compat_mode         = defaults::compat_mode());
+             CompatMode compat_mode   = defaults::compat_mode());
 
   /**
    * @brief FileHandle support move semantic but isn't copyable
@@ -84,7 +115,7 @@ class FileHandle {
     : _fd_direct_on{std::exchange(o._fd_direct_on, -1)},
       _fd_direct_off{std::exchange(o._fd_direct_off, -1)},
       _initialized{std::exchange(o._initialized, false)},
-      _compat_mode{std::exchange(o._compat_mode, false)},
+      _compat_mode{std::exchange(o._compat_mode, CompatMode::AUTO)},
       _nbytes{std::exchange(o._nbytes, 0)},
       _handle{std::exchange(o._handle, CUfileHandle_t{})}
   {
@@ -94,13 +125,18 @@ class FileHandle {
     _fd_direct_on  = std::exchange(o._fd_direct_on, -1);
     _fd_direct_off = std::exchange(o._fd_direct_off, -1);
     _initialized   = std::exchange(o._initialized, false);
-    _compat_mode   = std::exchange(o._compat_mode, false);
+    _compat_mode   = std::exchange(o._compat_mode, CompatMode::AUTO);
     _nbytes        = std::exchange(o._nbytes, 0);
     _handle        = std::exchange(o._handle, CUfileHandle_t{});
     return *this;
   }
   ~FileHandle() noexcept { close(); }
 
+  /**
+   * @brief Whether the file is closed according to its initialization status.
+   *
+   * @return Boolean answer.
+   */
   [[nodiscard]] bool closed() const noexcept { return !_initialized; }
 
   /**
@@ -110,7 +146,8 @@ class FileHandle {
   {
     if (closed()) { return; }
 
-    if (!_compat_mode) { cuFileAPI::instance().HandleDeregister(_handle); }
+    if (!is_compat_mode_preferred()) { cuFileAPI::instance().HandleDeregister(_handle); }
+    _compat_mode = CompatMode::AUTO;
     ::close(_fd_direct_off);
     if (_fd_direct_on != -1) { ::close(_fd_direct_on); }
     _fd_direct_on  = -1;
@@ -122,14 +159,14 @@ class FileHandle {
    * @brief Get the underlying cuFile file handle
    *
    * The file handle must be open and not in compatibility mode i.e.
-   * both `.closed()` and `.is_compat_mode_on()` must be return false.
+   * both `closed()` and `is_compat_mode_preferred()` must be false.
    *
    * @return cuFile's file handle
    */
   [[nodiscard]] CUfileHandle_t handle()
   {
     if (closed()) { throw CUfileException("File handle is closed"); }
-    if (_compat_mode) {
+    if (is_compat_mode_preferred()) {
       throw CUfileException("The underlying cuFile handle isn't available in compatibility mode");
     }
     return _handle;
@@ -202,7 +239,7 @@ class FileHandle {
                    std::size_t devPtr_offset,
                    bool sync_default_stream = true)
   {
-    if (_compat_mode) {
+    if (is_compat_mode_preferred()) {
       return detail::posix_device_read(
         _fd_direct_off, devPtr_base, size, file_offset, devPtr_offset);
     }
@@ -254,7 +291,7 @@ class FileHandle {
   {
     _nbytes = 0;  // Invalidate the computed file size
 
-    if (_compat_mode) {
+    if (is_compat_mode_preferred()) {
       return detail::posix_device_write(
         _fd_direct_off, devPtr_base, size, file_offset, devPtr_offset);
     }
@@ -333,7 +370,7 @@ class FileHandle {
     }
 
     // Let's synchronize once instead of in each task.
-    if (sync_default_stream && !_compat_mode) {
+    if (sync_default_stream && !is_compat_mode_preferred()) {
       PushAndPopContext c(ctx);
       CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(nullptr));
     }
@@ -410,7 +447,7 @@ class FileHandle {
     }
 
     // Let's synchronize once instead of in each task.
-    if (sync_default_stream && !_compat_mode) {
+    if (sync_default_stream && !is_compat_mode_preferred()) {
       PushAndPopContext c(ctx);
       CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(nullptr));
     }
@@ -469,16 +506,14 @@ class FileHandle {
                   ssize_t* bytes_read_p,
                   CUstream stream)
   {
-    // When checking for availability, we also check if cuFile's config file exist. This is because
-    // even when the stream API is available, it doesn't work if no config file exist.
-    if (kvikio::is_batch_and_stream_available() && !_compat_mode && !config_path().empty()) {
+    if (is_compat_mode_preferred_for_async(_compat_mode)) {
+      CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(stream));
+      *bytes_read_p =
+        static_cast<ssize_t>(read(devPtr_base, *size_p, *file_offset_p, *devPtr_offset_p));
+    } else {
       CUFILE_TRY(cuFileAPI::instance().ReadAsync(
         _handle, devPtr_base, size_p, file_offset_p, devPtr_offset_p, bytes_read_p, stream));
-      return;
     }
-    CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(stream));
-    *bytes_read_p =
-      static_cast<ssize_t>(read(devPtr_base, *size_p, *file_offset_p, *devPtr_offset_p));
   }
 
   /**
@@ -561,16 +596,14 @@ class FileHandle {
                    ssize_t* bytes_written_p,
                    CUstream stream)
   {
-    // When checking for availability, we also check if cuFile's config file exist. This is because
-    // even when the stream API is available, it doesn't work if no config file exist.
-    if (kvikio::is_batch_and_stream_available() && !_compat_mode && !config_path().empty()) {
+    if (is_compat_mode_preferred_for_async(_compat_mode)) {
+      CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(stream));
+      *bytes_written_p =
+        static_cast<ssize_t>(write(devPtr_base, *size_p, *file_offset_p, *devPtr_offset_p));
+    } else {
       CUFILE_TRY(cuFileAPI::instance().WriteAsync(
         _handle, devPtr_base, size_p, file_offset_p, devPtr_offset_p, bytes_written_p, stream));
-      return;
     }
-    CUDA_DRIVER_TRY(cudaAPI::instance().StreamSynchronize(stream));
-    *bytes_written_p =
-      static_cast<ssize_t>(write(devPtr_base, *size_p, *file_offset_p, *devPtr_offset_p));
   }
 
   /**
@@ -612,14 +645,35 @@ class FileHandle {
   }
 
   /**
-   * @brief Returns `true` if the compatibility mode has been enabled for this file.
+   * @brief Returns `true` if the compatibility mode is expected to be `ON` for this file.
    *
    * Compatibility mode can be explicitly enabled in object creation. The mode is also enabled
-   * automatically, if file cannot be opened with the `O_DIRECT` flag.
+   * automatically, if file cannot be opened with the `O_DIRECT` flag, or if the system does not
+   * meet the requirements for the cuFile library under the `AUTO` compatibility mode.
+   *
+   * @return Boolean answer.
+   */
+  [[nodiscard]] bool is_compat_mode_preferred() const noexcept
+  {
+    return defaults::is_compat_mode_preferred(_compat_mode);
+  }
+
+  /**
+   * @brief Returns `true` if the compatibility mode is expected to be `ON` for the asynchronous I/O
+   * on this file.
+   *
+   * For asynchronous I/O, the compatibility mode can be automatically enabled if the cuFile batch
+   * and stream symbols are missing, or if the cuFile configuration file is missing, or if
+   * `is_compat_mode_preferred()` returns true.
    *
-   * @return compatibility mode state for the object
+   * @return Boolean answer.
    */
-  [[nodiscard]] bool is_compat_mode_on() const noexcept { return _compat_mode; }
+  [[nodiscard]] bool is_compat_mode_preferred_for_async() const noexcept
+  {
+    static bool is_extra_symbol_available = is_batch_and_stream_available();
+    static bool is_config_path_empty      = config_path().empty();
+    return is_compat_mode_preferred() || !is_extra_symbol_available || is_config_path_empty;
+  }
 };
 
 }  // namespace kvikio
diff --git a/cpp/include/kvikio/shim/cufile.hpp b/cpp/include/kvikio/shim/cufile.hpp
index c5c7a0671f..7f12c29c3d 100644
--- a/cpp/include/kvikio/shim/cufile.hpp
+++ b/cpp/include/kvikio/shim/cufile.hpp
@@ -214,7 +214,7 @@ inline bool is_cufile_available()
  * @return The boolean answer
  */
 #if defined(KVIKIO_CUFILE_STREAM_API_FOUND) && defined(KVIKIO_CUFILE_STREAM_API_FOUND)
-inline bool is_batch_and_stream_available()
+inline bool is_batch_and_stream_available() noexcept
 {
   try {
     return is_cufile_available() && cuFileAPI::instance().stream_available;
diff --git a/cpp/src/file_handle.cpp b/cpp/src/file_handle.cpp
index c5b7ada59a..2e0de2537b 100644
--- a/cpp/src/file_handle.cpp
+++ b/cpp/src/file_handle.cpp
@@ -23,6 +23,7 @@
 #include <stdexcept>
 #include <system_error>
 
+#include <kvikio/defaults.hpp>
 #include <kvikio/file_handle.hpp>
 
 namespace kvikio {
@@ -118,12 +119,12 @@ int open_fd(const std::string& file_path, const std::string& flags, bool o_direc
 FileHandle::FileHandle(const std::string& file_path,
                        const std::string& flags,
                        mode_t mode,
-                       bool compat_mode)
+                       CompatMode compat_mode)
   : _fd_direct_off{open_fd(file_path, flags, false, mode)},
     _initialized{true},
     _compat_mode{compat_mode}
 {
-  if (_compat_mode) {
+  if (is_compat_mode_preferred()) {
     return;  // Nothing to do in compatibility mode
   }
 
@@ -131,13 +132,13 @@ FileHandle::FileHandle(const std::string& file_path,
   try {
     _fd_direct_on = open_fd(file_path, flags, true, mode);
   } catch (const std::system_error&) {
-    _compat_mode = true;
+    _compat_mode = CompatMode::ON;
   } catch (const std::invalid_argument&) {
-    _compat_mode = true;
+    _compat_mode = CompatMode::ON;
   }
 
   // Create a cuFile handle, if not in compatibility mode
-  if (!_compat_mode) {
+  if (!is_compat_mode_preferred()) {
     CUfileDescr_t desc{};  // It is important to set to zero!
     desc.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD;
     // NOLINTNEXTLINE(cppcoreguidelines-pro-type-union-access)
diff --git a/cpp/tests/test_basic_io.cpp b/cpp/tests/test_basic_io.cpp
index 12ccb6d428..c884ec6230 100644
--- a/cpp/tests/test_basic_io.cpp
+++ b/cpp/tests/test_basic_io.cpp
@@ -15,29 +15,84 @@
  */
 
 #include <kvikio/file_handle.hpp>
-
+#include "kvikio/defaults.hpp"
 #include "utils.hpp"
 
 using namespace kvikio::test;
 
-TEST(BasicIO, write_read)
+class BasicIOTest : public testing::Test {
+ protected:
+  void SetUp() override
+  {
+    TempDir tmp_dir{false};
+    _filepath = tmp_dir.path() / "test";
+
+    _dev_a = std::move(DevBuffer::arange(100));
+    _dev_b = std::move(DevBuffer::zero_like(_dev_a));
+  }
+
+  void TearDown() override {}
+
+  std::filesystem::path _filepath;
+  DevBuffer _dev_a;
+  DevBuffer _dev_b;
+};
+
+TEST_F(BasicIOTest, write_read)
 {
-  TempDir tmp_dir{false};
-  auto filepath = tmp_dir.path() / "test";
+  {
+    kvikio::FileHandle f(_filepath, "w");
+    auto nbytes = f.write(_dev_a.ptr, _dev_a.nbytes, 0, 0);
+    EXPECT_EQ(nbytes, _dev_a.nbytes);
+  }
 
-  auto dev_a = DevBuffer::arange(100);
-  auto dev_b = DevBuffer::zero_like(dev_a);
+  {
+    kvikio::FileHandle f(_filepath, "r");
+    auto nbytes = f.read(_dev_b.ptr, _dev_b.nbytes, 0, 0);
+    EXPECT_EQ(nbytes, _dev_b.nbytes);
+    expect_equal(_dev_a, _dev_b);
+  }
+}
+
+TEST_F(BasicIOTest, write_read_async)
+{
+  CUstream stream{};
+  CUDA_DRIVER_TRY(kvikio::cudaAPI::instance().StreamCreate(&stream, CU_STREAM_NON_BLOCKING));
 
+  // Default compatibility mode (AUTO)
   {
-    kvikio::FileHandle f(filepath, "w");
-    auto nbytes = f.write(dev_a.ptr, dev_a.nbytes, 0, 0);
-    EXPECT_EQ(nbytes, dev_a.nbytes);
+    kvikio::FileHandle f(_filepath, "w");
+    auto stream_future = f.write_async(_dev_a.ptr, _dev_a.nbytes, 0, 0, stream);
+    auto nbytes        = stream_future.check_bytes_done();
+    EXPECT_EQ(nbytes, _dev_a.nbytes);
   }
 
   {
-    kvikio::FileHandle f(filepath, "r");
-    auto nbytes = f.read(dev_b.ptr, dev_b.nbytes, 0, 0);
-    EXPECT_EQ(nbytes, dev_b.nbytes);
-    expect_equal(dev_a, dev_b);
+    kvikio::FileHandle f(_filepath, "r");
+    auto stream_future = f.read_async(_dev_b.ptr, _dev_b.nbytes, 0, 0, stream);
+    auto nbytes        = stream_future.check_bytes_done();
+    EXPECT_EQ(nbytes, _dev_b.nbytes);
+    expect_equal(_dev_a, _dev_b);
   }
+
+  // Explicitly set compatibility mode
+  std::array<kvikio::CompatMode, 2> compat_modes{kvikio::CompatMode::AUTO, kvikio::CompatMode::ON};
+  for (const auto& compat_mode : compat_modes) {
+    {
+      kvikio::FileHandle f(_filepath, "w", kvikio::FileHandle::m644, compat_mode);
+      auto stream_future = f.write_async(_dev_a.ptr, _dev_a.nbytes, 0, 0, stream);
+      auto nbytes        = stream_future.check_bytes_done();
+      EXPECT_EQ(nbytes, _dev_a.nbytes);
+    }
+
+    {
+      kvikio::FileHandle f(_filepath, "r", kvikio::FileHandle::m644, compat_mode);
+      auto stream_future = f.read_async(_dev_b.ptr, _dev_b.nbytes, 0, 0, stream);
+      auto nbytes        = stream_future.check_bytes_done();
+      EXPECT_EQ(nbytes, _dev_b.nbytes);
+      expect_equal(_dev_a, _dev_b);
+    }
+  }
+
+  CUDA_DRIVER_TRY(kvikio::cudaAPI::instance().StreamDestroy(stream));
 }
diff --git a/cpp/tests/test_defaults.cpp b/cpp/tests/test_defaults.cpp
new file mode 100644
index 0000000000..c4a88775e4
--- /dev/null
+++ b/cpp/tests/test_defaults.cpp
@@ -0,0 +1,53 @@
+/*
+ * Copyright (c) 2024, NVIDIA CORPORATION.
+ *
+ * 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 <stdexcept>
+
+#include <gtest/gtest.h>
+#include <kvikio/defaults.hpp>
+
+TEST(Defaults, parse_compat_mode_str)
+{
+  {
+    std::vector<std::string> inputs{
+      "ON", "on", "On", "TRUE", "true", "True", "YES", "yes", "Yes", "1"};
+    for (const auto& input : inputs) {
+      EXPECT_EQ(kvikio::detail::parse_compat_mode_str(input), kvikio::CompatMode::ON);
+    }
+  }
+
+  {
+    std::vector<std::string> inputs{
+      "OFF", "off", "oFf", "FALSE", "false", "False", "NO", "no", "No", "0"};
+    for (const auto& input : inputs) {
+      EXPECT_EQ(kvikio::detail::parse_compat_mode_str(input), kvikio::CompatMode::OFF);
+    }
+  }
+
+  {
+    std::vector<std::string> inputs{"AUTO", "auto", "aUtO"};
+    for (const auto& input : inputs) {
+      EXPECT_EQ(kvikio::detail::parse_compat_mode_str(input), kvikio::CompatMode::AUTO);
+    }
+  }
+
+  {
+    std::vector<std::string> inputs{"", "invalidOption", "11", "*&^Yes"};
+    for (const auto& input : inputs) {
+      EXPECT_THROW(kvikio::detail::parse_compat_mode_str(input), std::invalid_argument);
+    }
+  }
+}
diff --git a/cpp/tests/utils.hpp b/cpp/tests/utils.hpp
index 56a2cd5c45..1c671a82bc 100644
--- a/cpp/tests/utils.hpp
+++ b/cpp/tests/utils.hpp
@@ -110,10 +110,12 @@ class TempDir {
  */
 class DevBuffer {
  public:
-  const std::size_t nelem;
-  const std::size_t nbytes;
+  std::size_t nelem;
+  std::size_t nbytes;
   void* ptr{nullptr};
 
+  DevBuffer() : nelem{0}, nbytes{0} {};
+
   DevBuffer(std::size_t nelem) : nelem{nelem}, nbytes{nelem * sizeof(std::int64_t)}
   {
     KVIKIO_CHECK_CUDA(cudaMalloc(&ptr, nbytes));
@@ -123,6 +125,21 @@ class DevBuffer {
     KVIKIO_CHECK_CUDA(cudaMemcpy(ptr, host_buffer.data(), nbytes, cudaMemcpyHostToDevice));
   }
 
+  DevBuffer(DevBuffer&& dev_buffer) noexcept
+    : nelem{std::exchange(dev_buffer.nelem, 0)},
+      nbytes{std::exchange(dev_buffer.nbytes, 0)},
+      ptr{std::exchange(dev_buffer.ptr, nullptr)}
+  {
+  }
+
+  DevBuffer& operator=(DevBuffer&& dev_buffer) noexcept
+  {
+    nelem  = std::exchange(dev_buffer.nelem, 0);
+    nbytes = std::exchange(dev_buffer.nbytes, 0);
+    ptr    = std::exchange(dev_buffer.ptr, nullptr);
+    return *this;
+  }
+
   ~DevBuffer() noexcept { cudaFree(ptr); }
 
   [[nodiscard]] static DevBuffer arange(std::size_t nelem, std::int64_t start = 0)
diff --git a/docs/source/runtime_settings.rst b/docs/source/runtime_settings.rst
index 631ba0c937..be5508741e 100644
--- a/docs/source/runtime_settings.rst
+++ b/docs/source/runtime_settings.rst
@@ -3,15 +3,21 @@ Runtime Settings
 
 Compatibility Mode ``KVIKIO_COMPAT_MODE``
 -----------------------------------------
-When KvikIO is running in compatibility mode, it doesn't load ``libcufile.so``. Instead, reads and writes are done using POSIX. Notice, this is not the same as the compatibility mode in cuFile. That is cuFile can run in compatibility mode while KvikIO is not.
-Set the environment variable ``KVIKIO_COMPAT_MODE`` to enable/disable compatibility mode. By default, compatibility mode is enabled:
+When KvikIO is running in compatibility mode, it doesn't load ``libcufile.so``. Instead, reads and writes are done using POSIX. Notice, this is not the same as the compatibility mode in cuFile. It is possible that KvikIO performs I/O in the non-compatibility mode by using the cuFile library, but the cuFile library itself is configured to operate in its own compatibility mode. For more details, refer to `cuFile compatibility mode <https://docs.nvidia.com/gpudirect-storage/api-reference-guide/index.html#cufile-compatibility-mode>`_ and `cuFile environment variables <https://docs.nvidia.com/gpudirect-storage/troubleshooting-guide/index.html#environment-variables>`_ .
+
+The environment variable ``KVIKIO_COMPAT_MODE`` has three options (case-insensitive):
+
+  * ``ON`` (aliases: ``TRUE``, ``YES``, ``1``): Enable the compatibility mode.
+  * ``OFF`` (aliases: ``FALSE``, ``NO``, ``0``): Disable the compatibility mode, and enforce cuFile I/O. GDS will be activated if the system requirements for cuFile are met and cuFile is properly configured. However, if the system is not suited for cuFile, I/O operations under the ``OFF`` option may error out, crash or hang.
+  * ``AUTO``: Try cuFile I/O first, and fall back to POSIX I/O if the system requirements for cuFile are not met.
+
+Under ``AUTO``, KvikIO falls back to the compatibility mode:
 
   * when ``libcufile.so`` cannot be found.
   * when running in Windows Subsystem for Linux (WSL).
   * when ``/run/udev`` isn't readable, which typically happens when running inside a docker image not launched with ``--volume /run/udev:/run/udev:ro``.
 
-This setting can also be controlled by :py:func:`kvikio.defaults.compat_mode`, :py:func:`kvikio.defaults.compat_mode_reset`, and :py:func:`kvikio.defaults.set_compat_mode`.
-
+This setting can also be programmatically controlled by :py:func:`kvikio.defaults.set_compat_mode` and :py:func:`kvikio.defaults.compat_mode_reset`.
 
 Thread Pool ``KVIKIO_NTHREADS``
 -------------------------------
diff --git a/python/kvikio/kvikio/__init__.py b/python/kvikio/kvikio/__init__.py
index f4db6d1d05..64aa95df5c 100644
--- a/python/kvikio/kvikio/__init__.py
+++ b/python/kvikio/kvikio/__init__.py
@@ -12,6 +12,7 @@
     del libkvikio
 
 
+from kvikio._lib.defaults import CompatMode  # noqa: F401
 from kvikio._version import __git_commit__, __version__
 from kvikio.cufile import CuFile
 from kvikio.remote_file import RemoteFile, is_remote_file_available
diff --git a/python/kvikio/kvikio/_lib/defaults.pyx b/python/kvikio/kvikio/_lib/defaults.pyx
index f59cad5cb4..9042069b74 100644
--- a/python/kvikio/kvikio/_lib/defaults.pyx
+++ b/python/kvikio/kvikio/_lib/defaults.pyx
@@ -4,13 +4,18 @@
 # distutils: language = c++
 # cython: language_level=3
 
+from libc.stdint cimport uint8_t
 from libcpp cimport bool
 
 
-cdef extern from "<kvikio/defaults.hpp>" nogil:
-    bool cpp_compat_mode "kvikio::defaults::compat_mode"() except +
+cdef extern from "<kvikio/defaults.hpp>" namespace "kvikio" nogil:
+    cpdef enum class CompatMode(uint8_t):
+        OFF = 0
+        ON = 1
+        AUTO = 2
+    CompatMode cpp_compat_mode "kvikio::defaults::compat_mode"() except +
     void cpp_compat_mode_reset \
-        "kvikio::defaults::compat_mode_reset"(bool enable) except +
+        "kvikio::defaults::compat_mode_reset"(CompatMode compat_mode) except +
     unsigned int cpp_thread_pool_nthreads \
         "kvikio::defaults::thread_pool_nthreads"() except +
     void cpp_thread_pool_nthreads_reset \
@@ -25,12 +30,12 @@ cdef extern from "<kvikio/defaults.hpp>" nogil:
         "kvikio::defaults::bounce_buffer_size_reset"(size_t nbytes) except +
 
 
-def compat_mode() -> bool:
+def compat_mode() -> CompatMode:
     return cpp_compat_mode()
 
 
-def compat_mode_reset(enable: bool) -> None:
-    cpp_compat_mode_reset(enable)
+def compat_mode_reset(compat_mode: CompatMode) -> None:
+    cpp_compat_mode_reset(compat_mode)
 
 
 def thread_pool_nthreads() -> int:
diff --git a/python/kvikio/kvikio/defaults.py b/python/kvikio/kvikio/defaults.py
index a0ff265873..9e959c1f74 100644
--- a/python/kvikio/kvikio/defaults.py
+++ b/python/kvikio/kvikio/defaults.py
@@ -7,7 +7,7 @@
 import kvikio._lib.defaults
 
 
-def compat_mode() -> bool:
+def compat_mode() -> kvikio.CompatMode:
     """Check if KvikIO is running in compatibility mode.
 
     Notice, this is not the same as the compatibility mode in cuFile. That is,
@@ -18,10 +18,11 @@ def compat_mode() -> bool:
 
     Set the environment variable `KVIKIO_COMPAT_MODE` to enable/disable compatibility
     mode. By default, compatibility mode is enabled:
+
     - when `libcufile` cannot be found
     - when running in Windows Subsystem for Linux (WSL)
     - when `/run/udev` isn't readable, which typically happens when running inside
-    a docker image not launched with `--volume /run/udev:/run/udev:ro`
+      a docker image not launched with `--volume /run/udev:/run/udev:ro`
 
     Returns
     -------
@@ -31,32 +32,36 @@ def compat_mode() -> bool:
     return kvikio._lib.defaults.compat_mode()
 
 
-def compat_mode_reset(enable: bool) -> None:
+def compat_mode_reset(compatmode: kvikio.CompatMode) -> None:
     """Reset the compatibility mode.
 
     Use this function to enable/disable compatibility mode explicitly.
 
     Parameters
     ----------
-    enable : bool
-        Set to True to enable and False to disable compatibility mode
+    compatmode : kvikio.CompatMode
+        Set to kvikio.CompatMode.ON to enable and kvikio.CompatMode.OFF to disable
+        compatibility mode, or kvikio.CompatMode.AUTO to let KvikIO determine: try
+        OFF first, and upon failure, fall back to ON.
     """
-    kvikio._lib.defaults.compat_mode_reset(enable)
+    kvikio._lib.defaults.compat_mode_reset(compatmode)
 
 
 @contextlib.contextmanager
-def set_compat_mode(enable: bool):
+def set_compat_mode(compatmode: kvikio.CompatMode):
     """Context for resetting the compatibility mode.
 
     Parameters
     ----------
-    enable : bool
-        Set to True to enable and False to disable compatibility mode
+    compatmode : kvikio.CompatMode
+        Set to kvikio.CompatMode.ON to enable and kvikio.CompatMode.OFF to disable
+        compatibility mode, or kvikio.CompatMode.AUTO to let KvikIO determine: try
+        OFF first, and upon failure, fall back to ON.
     """
     num_threads_reset(get_num_threads())  # Sync all running threads
     old_value = compat_mode()
     try:
-        compat_mode_reset(enable)
+        compat_mode_reset(compatmode)
         yield
     finally:
         compat_mode_reset(old_value)
diff --git a/python/kvikio/tests/test_defaults.py b/python/kvikio/tests/test_defaults.py
index 39892a784d..d7048c418d 100644
--- a/python/kvikio/tests/test_defaults.py
+++ b/python/kvikio/tests/test_defaults.py
@@ -8,17 +8,19 @@
 
 
 @pytest.mark.skipif(
-    kvikio.defaults.compat_mode(),
+    kvikio.defaults.compat_mode() == kvikio.CompatMode.ON,
     reason="cannot test `compat_mode` when already running in compatibility mode",
 )
 def test_compat_mode():
     """Test changing `compat_mode`"""
 
     before = kvikio.defaults.compat_mode()
-    with kvikio.defaults.set_compat_mode(True):
-        assert kvikio.defaults.compat_mode()
-        kvikio.defaults.compat_mode_reset(False)
-        assert not kvikio.defaults.compat_mode()
+    with kvikio.defaults.set_compat_mode(kvikio.CompatMode.ON):
+        assert kvikio.defaults.compat_mode() == kvikio.CompatMode.ON
+        kvikio.defaults.compat_mode_reset(kvikio.CompatMode.OFF)
+        assert kvikio.defaults.compat_mode() == kvikio.CompatMode.OFF
+        kvikio.defaults.compat_mode_reset(kvikio.CompatMode.AUTO)
+        assert kvikio.defaults.compat_mode() == kvikio.CompatMode.AUTO
     assert before == kvikio.defaults.compat_mode()