diff --git a/adoc/extensions/index.adoc b/adoc/extensions/index.adoc
index d16bc954b..84f46f924 100644
--- a/adoc/extensions/index.adoc
+++ b/adoc/extensions/index.adoc
@@ -18,3 +18,4 @@ include::sycl_khr_queue_flush.adoc[leveloffset=2]
 include::sycl_khr_work_item_queries.adoc[leveloffset=2]
 include::sycl_khr_static_addrspace_cast.adoc[leveloffset=2]
 include::sycl_khr_dynamic_addrspace_cast.adoc[leveloffset=2]
+include::sycl_khr_prefetch_host.adoc[leveloffset=2]
diff --git a/adoc/extensions/sycl_khr_prefetch_host.adoc b/adoc/extensions/sycl_khr_prefetch_host.adoc
new file mode 100644
index 000000000..4fe50ddd1
--- /dev/null
+++ b/adoc/extensions/sycl_khr_prefetch_host.adoc
@@ -0,0 +1,103 @@
+[[sec:khr-prefetch-host]]
+= sycl_khr_prefetch_host
+
+This extension allows developers to prefetch data from the device to the host.
+
+[[sec:khr-prefetch-host-dependencies]]
+== Dependencies
+
+This extension has no dependencies on other extensions.
+
+[[sec:khr-prefetch-host-feature-test]]
+== Feature test macro
+
+An implementation supporting this extension must predefine the macro
+[code]#SYCL_KHR_PREFETCH_HOST# to one of the values defined in the table below.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+[[sec:khr-prefetch-host-handler]]
+== Extensions to the handler class
+
+This extension adds the following function to the [code]#sycl::handler# class.
+
+[[sec:khr-prefetch-host-handler-member-funcs]]
+=== Member functions
+
+.[apidef]#handler::khr_prefetch_host#
+[source,role=synopsis,id=api:handler-khr-prefetch-host]
+----
+void khr_prefetch_host(void* ptr, size_t numBytes)
+----
+
+_Effects:_ Enqueues a prefetch host of [code]#num_bytes# of data starting at
+address [code]#ptr#.
+The [code]#ptr# must point within a USM allocation from the same context as the
+handler's queue, and the pointer must be accessible from the queue's device.
+For more detail on USM, please see <<sec:usm>>.
+
+{note} [code]#prefetch_host# inform the SYCL runtime that the specified shared
+allocation is likely to be accessed on the host in the future, and that it is
+free to migrate the allocation to the host.
+{endnote}
+
+'''
+
+[[sec:khr-prefetch-host-queue]]
+== Extensions to the queue class
+
+This extension adds the following function to the [code]#sycl::queue# class.
+
+[[sec:khr-prefetch-host-queue-member-funcs]]
+=== Member functions
+
+.[apidef]#queue::khr_prefetch_host#
+[source,role=synopsis,id=api:queue-khr-prefetch-host]
+----
+event khr_prefetch_host(void* ptr, size_t numBytes)                                      (1)
+
+event khr_prefetch_host(void* ptr, size_t numBytes, event depEvent)                      (2)
+
+event khr_prefetch_host(void* ptr, size_t numBytes, const std::vector<event>& depEvents) (3)
+
+----
+
+_Effects (1)_: Equivalent to calling [code]#queue::submit# with a command group
+function that calls [code]#handler::khr_prefetch_host(ptr, numBytes)#.
+
+_Effects (2)_: Equivalent to calling [code]#queue::submit# with a command group
+function that calls [code]#handler::depends_on(depEvent)# and
+[code]#handler::khr_prefetch_host(ptr, numBytes)#.
+
+_Effects (3)_: Equivalent to calling [code]#queue::submit# with a command group
+function that calls [code]#handler::depends_on(depEvents)# and
+[code]#handler::khr_prefetch_host(ptr, numBytes)#.
+
+_Returns_: An event which represents the command which is submitted to the
+queue.
+
+'''
+
+[[sec:khr-prefetch-host-example]]
+== Example
+
+The example below demonstrates the usage of this extension.
+
+[source,,linenums]
+----
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  int *A = Q.malloc_shared<int>(1,Q);
+  Q.prefetch_host(A, sizeof(int)).wait();
+  A[0] = 10; // Maybe faster than without the prefetch_host
+}
+----