diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 129eba7543..c1b83839ed 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -10,6 +10,7 @@ SYCL/AtomicRef @AGindinson
 SYCL/Assert @intel/llvm-reviewers-runtime
 SYCL/Basic @intel/llvm-reviewers-runtime
 SYCL/Config @intel/llvm-reviewers-runtime
+SYCL/DiscardEvents @intel/llvm-reviewers-runtime
 SYCL/FilterSelector @intel/llvm-reviewers-runtime
 SYCL/HostInteropTask @intel/llvm-reviewers-runtime
 SYCL/InorderQueue @intel/llvm-reviewers-runtime
diff --git a/SYCL/DiscardEvents/discard_events_accessors.cpp b/SYCL/DiscardEvents/discard_events_accessors.cpp
new file mode 100644
index 0000000000..e77645cc02
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_accessors.cpp
@@ -0,0 +1,106 @@
+// FIXME unsupported on level_zero until L0 Plugin support becomes available for
+// discard_queue_events
+// UNSUPPORTED: level_zero
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt || true
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for
+// piEnqueueKernelLaunch for USM kernel using local accessor, but
+// is not `nullptr` for kernel using buffer accessor.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: The test passed.
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+static constexpr int MAGIC_NUM = -1;
+static constexpr size_t BUFFER_SIZE = 16;
+
+void RunKernelHelper(sycl::queue Q,
+                     const std::function<void(int *Harray)> &TestFunction) {
+  int *Harray = sycl::malloc_host<int>(BUFFER_SIZE, Q);
+  assert(Harray != nullptr);
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    Harray[i] = MAGIC_NUM;
+  }
+
+  TestFunction(Harray);
+
+  // Checks result
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    size_t expected = i + 10;
+    assert(Harray[i] == expected);
+  }
+  free(Harray, Q);
+}
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(props);
+  sycl::range<1> Range(BUFFER_SIZE);
+
+  RunKernelHelper(Q, [&](int *Harray) {
+    Q.submit([&](sycl::handler &CGH) {
+      const size_t LocalMemSize = BUFFER_SIZE;
+      using LocalAccessor =
+          sycl::accessor<int, 1, sycl::access::mode::read_write,
+                         sycl::access::target::local>;
+      LocalAccessor LocalAcc(LocalMemSize, CGH);
+
+      CGH.parallel_for<class kernel_using_local_memory>(
+          Range, [=](sycl::item<1> itemID) {
+            size_t i = itemID.get_id(0);
+            int *Ptr = LocalAcc.get_pointer();
+            Ptr[i] = i + 5;
+            Harray[i] = Ptr[i] + 5;
+          });
+    });
+    Q.wait();
+  });
+
+  RunKernelHelper(Q, [&](int *Harray) {
+    sycl::buffer<int, 1> Buf(Range);
+    Q.submit([&](sycl::handler &CGH) {
+      auto Acc = Buf.get_access<sycl::access::mode::read_write>(CGH);
+      CGH.parallel_for<class kernel_using_buffer_accessor>(
+          Range, [=](sycl::item<1> itemID) {
+            size_t i = itemID.get_id(0);
+            Harray[i] = i + 10;
+            Acc[i] = i + 20;
+          });
+    });
+    Q.wait();
+
+    // Checks result
+    auto HostAcc = Buf.get_access<sycl::access::mode::read>();
+    for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+      size_t expected = i + 20;
+      assert(HostAcc[i] == expected);
+    }
+  });
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/discard_events_host_task.cpp b/SYCL/DiscardEvents/discard_events_host_task.cpp
new file mode 100644
index 0000000000..be536f9865
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_host_task.cpp
@@ -0,0 +1,57 @@
+// If necessary, the test can be removed as run_on_host_intel() is deprecated
+// and host_task() which should be used instead does not use the PI call
+// piEnqueueNativeKernel
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for
+// piEnqueueNativeKernel.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueNativeKernel(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: The test passed.
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+
+void CheckArray(sycl::queue Q, int *x, size_t buffer_size, int expected) {
+  Q.wait();
+  for (size_t i = 0; i < buffer_size; ++i)
+    assert(x[i] == expected);
+}
+
+static constexpr size_t BUFFER_SIZE = 16;
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(Props);
+
+  int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(x != nullptr);
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.run_on_host_intel([=]() {
+      for (size_t i = 0; i < BUFFER_SIZE; ++i)
+        x[i] = 8;
+    });
+  });
+  CheckArray(Q, x, BUFFER_SIZE, 8);
+
+  Q.wait();
+  free(x, Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/discard_events_kernel_using_assert.hpp b/SYCL/DiscardEvents/discard_events_kernel_using_assert.hpp
new file mode 100644
index 0000000000..7924a0195d
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_kernel_using_assert.hpp
@@ -0,0 +1,45 @@
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+static constexpr int MAGIC_NUM = -1;
+static constexpr size_t BUFFER_SIZE = 16;
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(Props);
+
+  sycl::range<1> Range(BUFFER_SIZE);
+  int *Harray = sycl::malloc_host<int>(BUFFER_SIZE, Q);
+  if (Harray == nullptr) {
+    return -1;
+  }
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    Harray[i] = MAGIC_NUM;
+  }
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.parallel_for<class kernel_using_assert>(
+        Range, [=](sycl::item<1> itemID) {
+          size_t i = itemID.get_id(0);
+          Harray[i] = i + 10;
+          assert(Harray[i] == i + 10 && "assert message");
+        });
+  });
+  Q.wait();
+
+  // Checks result
+  for (size_t i = 0; i < BUFFER_SIZE; ++i) {
+    size_t expected = i + 10;
+    if (Harray[i] != expected)
+      return -1;
+  }
+  free(Harray, Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/discard_events_test_queue_ops.hpp b/SYCL/DiscardEvents/discard_events_test_queue_ops.hpp
new file mode 100644
index 0000000000..47e82ec234
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_test_queue_ops.hpp
@@ -0,0 +1,132 @@
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+
+void CheckArray(sycl::queue Q, int *x, size_t buffer_size, int expected) {
+  Q.wait();
+  for (size_t i = 0; i < buffer_size; ++i)
+    assert(x[i] == expected);
+}
+
+static constexpr size_t BUFFER_SIZE = 16;
+
+void TestQueueOperations(sycl::queue Q) {
+  sycl::range<1> Range(BUFFER_SIZE);
+  auto Dev = Q.get_device();
+  auto Ctx = Q.get_context();
+  const int MemAdvice =
+      ((Dev.get_backend() == sycl::backend::ext_oneapi_cuda) ? 1 : 0);
+  int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(x != nullptr);
+  int *y = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(y != nullptr);
+
+  Q.memset(x, 0, BUFFER_SIZE * sizeof(int));
+  CheckArray(Q, x, BUFFER_SIZE, 0);
+
+  Q.memcpy(y, x, BUFFER_SIZE * sizeof(int));
+  CheckArray(Q, y, BUFFER_SIZE, 0);
+
+  Q.fill(y, 1, BUFFER_SIZE);
+  CheckArray(Q, y, BUFFER_SIZE, 1);
+
+  Q.copy(y, x, BUFFER_SIZE);
+  CheckArray(Q, x, BUFFER_SIZE, 1);
+
+  Q.prefetch(y, BUFFER_SIZE * sizeof(int));
+  Q.mem_advise(y, BUFFER_SIZE * sizeof(int), MemAdvice);
+  Q.ext_oneapi_submit_barrier();
+
+  Q.single_task([=] {
+    for (auto i = 0u; i < BUFFER_SIZE; ++i)
+      y[i] *= 2;
+  });
+  CheckArray(Q, y, BUFFER_SIZE, 2);
+
+  Q.parallel_for(Range,
+                 [=](sycl::item<1> itemID) { y[itemID.get_id(0)] *= 3; });
+  CheckArray(Q, y, BUFFER_SIZE, 6);
+
+  // Creates new queue with the same context/device, but without discard_events
+  // property. This queue returns a normal event, not a discarded one.
+  sycl::queue RegularQ(Ctx, Dev, sycl::property::queue::in_order{});
+  int *x1 = sycl::malloc_shared<int>(BUFFER_SIZE, RegularQ);
+  assert(x1 != nullptr);
+  auto event = RegularQ.memset(x1, 0, BUFFER_SIZE * sizeof(int));
+
+  Q.memcpy(y, x, 0, event);
+  CheckArray(Q, y, BUFFER_SIZE, 6);
+
+  Q.wait();
+  free(x, Q);
+  free(y, Q);
+  free(x1, RegularQ);
+}
+
+void TestQueueOperationsViaSubmit(sycl::queue Q) {
+  sycl::range<1> Range(BUFFER_SIZE);
+  auto Dev = Q.get_device();
+  auto Ctx = Q.get_context();
+  const int MemAdvice =
+      ((Dev.get_backend() == sycl::backend::ext_oneapi_cuda) ? 1 : 0);
+  int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(x != nullptr);
+  int *y = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(y != nullptr);
+
+  Q.submit(
+      [&](sycl::handler &CGH) { CGH.memset(x, 0, BUFFER_SIZE * sizeof(int)); });
+  CheckArray(Q, x, BUFFER_SIZE, 0);
+
+  Q.submit(
+      [&](sycl::handler &CGH) { CGH.memcpy(y, x, BUFFER_SIZE * sizeof(int)); });
+  CheckArray(Q, y, BUFFER_SIZE, 0);
+
+  Q.submit([&](sycl::handler &CGH) { CGH.fill(y, 1, BUFFER_SIZE); });
+  CheckArray(Q, y, BUFFER_SIZE, 1);
+
+  Q.submit([&](sycl::handler &CGH) { CGH.copy(y, x, BUFFER_SIZE); });
+  CheckArray(Q, x, BUFFER_SIZE, 1);
+
+  Q.submit(
+      [&](sycl::handler &CGH) { CGH.prefetch(y, BUFFER_SIZE * sizeof(int)); });
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.mem_advise(y, BUFFER_SIZE * sizeof(int), MemAdvice);
+  });
+  Q.submit([&](sycl::handler &CGH) { CGH.ext_oneapi_barrier(); });
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.single_task([=] {
+      for (auto i = 0u; i < BUFFER_SIZE; ++i)
+        y[i] *= 2;
+    });
+  });
+  CheckArray(Q, y, BUFFER_SIZE, 2);
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.parallel_for(Range,
+                     [=](sycl::item<1> itemID) { y[itemID.get_id(0)] *= 3; });
+  });
+  CheckArray(Q, y, BUFFER_SIZE, 6);
+
+  // Creates new queue with the same context/device, but without discard_events
+  // property. This queue returns a normal event, not a discarded one.
+  sycl::queue RegularQ(Ctx, Dev, sycl::property::queue::in_order{});
+  int *x1 = sycl::malloc_shared<int>(BUFFER_SIZE, RegularQ);
+  assert(x1 != nullptr);
+  auto event = RegularQ.memset(x1, 0, BUFFER_SIZE * sizeof(int));
+
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(event);
+    CGH.memcpy(y, x, 0);
+  });
+  CheckArray(Q, y, BUFFER_SIZE, 6);
+
+  Q.wait();
+  free(x, Q);
+  free(y, Q);
+  free(x1, RegularQ);
+}
diff --git a/SYCL/DiscardEvents/discard_events_using_assert.cpp b/SYCL/DiscardEvents/discard_events_using_assert.cpp
new file mode 100644
index 0000000000..bcfc24c5af
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_using_assert.cpp
@@ -0,0 +1,24 @@
+// FIXME unsupported on CUDA and HIP until fallback libdevice becomes available
+// UNSUPPORTED: cuda || hip
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is not `nullptr` for
+// piEnqueueKernelLaunch.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: The test passed.
+
+#include "discard_events_kernel_using_assert.hpp"
diff --git a/SYCL/DiscardEvents/discard_events_using_assert_ndebug.cpp b/SYCL/DiscardEvents/discard_events_using_assert_ndebug.cpp
new file mode 100644
index 0000000000..943dc6cff2
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_using_assert_ndebug.cpp
@@ -0,0 +1,24 @@
+// FIXME unsupported on level_zero until L0 Plugin support becomes available for
+// discard_queue_events
+// UNSUPPORTED: level_zero
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -DNDEBUG -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for
+// piEnqueueKernelLaunch.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: The test passed.
+
+#include "discard_events_kernel_using_assert.hpp"
diff --git a/SYCL/DiscardEvents/discard_events_usm.cpp b/SYCL/DiscardEvents/discard_events_usm.cpp
new file mode 100644
index 0000000000..91e34f9917
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_usm.cpp
@@ -0,0 +1,133 @@
+// FIXME unsupported on level_zero until L0 Plugin support becomes available for
+// discard_queue_events
+// UNSUPPORTED: level_zero
+//
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is `nullptr` for all PI calls that
+// should discard events.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// Everything that follows TestQueueOperations()
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// Q.fill don't use piEnqueueMemBufferFill
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// ---> piEnqueueMemBufferCopy(
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueuePrefetch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueueMemAdvise(
+// CHECK: ) --->  pi_result : PI_SUCCESS
+// CHECK-NEXT:         [out]pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueEventsWaitWithBarrier(
+// CHECK: ) --->  pi_result : PI_SUCCESS
+// CHECK-NEXT:         [out]pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// RegularQueue
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWait(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// Everything that follows TestQueueOperationsViaSubmit()
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// Q.fill don't use piEnqueueMemBufferFill
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// ---> piEnqueueMemBufferCopy(
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueuePrefetch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piextUSMEnqueueMemAdvise(
+// CHECK: ) --->  pi_result : PI_SUCCESS
+// CHECK-NEXT:         [out]pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueEventsWaitWithBarrier(
+// CHECK: ) --->  pi_result : PI_SUCCESS
+// CHECK-NEXT:         [out]pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// RegularQueue
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWait(
+// CHECK:        pi_event * :
+// CHECK-NEXT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+//
+// CHECK: The test passed.
+
+#include "discard_events_test_queue_ops.hpp"
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Q(Props);
+
+  TestQueueOperations(Q);
+
+  TestQueueOperationsViaSubmit(Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/discard_events_usm_ooo_queue.cpp b/SYCL/DiscardEvents/discard_events_usm_ooo_queue.cpp
new file mode 100644
index 0000000000..0338e31bf0
--- /dev/null
+++ b/SYCL/DiscardEvents/discard_events_usm_ooo_queue.cpp
@@ -0,0 +1,144 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %CPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %GPU_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+// RUN: env SYCL_PI_TRACE=2 %ACC_RUN_PLACEHOLDER %t.out &> %t.txt
+// RUN: %ACC_RUN_PLACEHOLDER FileCheck %s --input-file %t.txt
+//
+// The test checks that the last parameter is not `nullptr` for all PI calls
+// that should discard events.
+// {{0|0000000000000000}} is required for various output on Linux and Windows.
+//
+// Everything that follows TestQueueOperations()
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// Q.fill don't use piEnqueueMemBufferFill
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// ---> piEnqueueMemBufferCopy(
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueuePrefetch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueueMemAdvise(
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWaitWithBarrier(
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// RegularQueue
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWait(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// Everything that follows TestQueueOperationsViaSubmit()
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// Q.fill don't use piEnqueueMemBufferFill
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// ---> piEnqueueMemBufferCopy(
+// CHECK: ---> piextUSMEnqueueMemcpy(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueuePrefetch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piextUSMEnqueueMemAdvise(
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWaitWithBarrier(
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueKernelLaunch(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// RegularQueue
+// CHECK: ---> piextUSMEnqueueMemset(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: ---> piEnqueueEventsWait(
+// CHECK:        pi_event * :
+// CHECK-NOT:        pi_event * : {{0|0000000000000000}}[ nullptr ]
+// CHECK: --->  pi_result : PI_SUCCESS
+//
+// CHECK: The test passed.
+
+#include "discard_events_test_queue_ops.hpp"
+
+int main(int Argc, const char *Argv[]) {
+
+  sycl::property_list Props{
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue OOO_Q(Props);
+
+  TestQueueOperations(OOO_Q);
+
+  TestQueueOperationsViaSubmit(OOO_Q);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/invalid_event.cpp b/SYCL/DiscardEvents/invalid_event.cpp
new file mode 100644
index 0000000000..070625142e
--- /dev/null
+++ b/SYCL/DiscardEvents/invalid_event.cpp
@@ -0,0 +1,97 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//
+// The test checks that each PI call to the queue returns a discarded event
+// with the status "ext_oneapi_unknown"
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+static constexpr size_t BUFFER_SIZE = 16;
+
+void QueueAPIsReturnDiscardedEvent(sycl::queue Q) {
+  sycl::range<1> range(BUFFER_SIZE);
+
+  auto Dev = Q.get_device();
+  const int MemAdvice =
+      ((Dev.get_backend() == sycl::backend::ext_oneapi_cuda) ? 1 : 0);
+  int *x = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(x != nullptr);
+  int *y = sycl::malloc_shared<int>(BUFFER_SIZE, Q);
+  assert(y != nullptr);
+
+  sycl::event DiscardedEvent;
+
+  DiscardedEvent = Q.memset(x, 0, BUFFER_SIZE * sizeof(int));
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.memcpy(y, x, BUFFER_SIZE * sizeof(int));
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.fill(y, 1, BUFFER_SIZE);
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.copy(y, x, BUFFER_SIZE);
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.prefetch(y, BUFFER_SIZE * sizeof(int));
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.mem_advise(y, BUFFER_SIZE * sizeof(int), MemAdvice);
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.single_task([=] {});
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.submit([&](sycl::handler &CGH) {
+    CGH.parallel_for(range, [=](sycl::item<1> itemID) {});
+  });
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  DiscardedEvent = Q.submit_barrier();
+  assert(
+      DiscardedEvent.get_info<sycl::info::event::command_execution_status>() ==
+      sycl::info::event_command_status::ext_oneapi_unknown);
+
+  Q.wait();
+  free(x, Q);
+  free(y, Q);
+}
+
+int main(int Argc, const char *Argv[]) {
+  sycl::property_list Props1{
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue OOO_Queue(Props1);
+  QueueAPIsReturnDiscardedEvent(OOO_Queue);
+
+  sycl::property_list Props2{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Inorder_Queue(Props2);
+  QueueAPIsReturnDiscardedEvent(Inorder_Queue);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}
diff --git a/SYCL/DiscardEvents/invalid_event_exceptions.cpp b/SYCL/DiscardEvents/invalid_event_exceptions.cpp
new file mode 100644
index 0000000000..ab795e41f0
--- /dev/null
+++ b/SYCL/DiscardEvents/invalid_event_exceptions.cpp
@@ -0,0 +1,170 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+//
+// RUN: %HOST_RUN_PLACEHOLDER %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//
+// The test checks 3 things:
+// 1. An attempt to construct a queue with both properties(discard_events and
+// enable_profiling) throws an exception.
+// 2. Checks the APIs for discarded event that should throw an exception that
+// they do it.
+// 3. An attempt to pass discarded event into depends_on throws an exception.
+
+#include <CL/sycl.hpp>
+#include <cassert>
+#include <iostream>
+
+using namespace cl::sycl;
+
+void DiscardedEventWaitExceptionHelper(
+    const std::function<void()> &FunctionToTry) {
+  try {
+    FunctionToTry();
+    assert(false && "No exception was thrown.");
+  } catch (const sycl::exception &e) {
+    assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+           "sycl::exception code was not the expected sycl::errc::invalid.");
+  } catch (...) {
+    assert(false &&
+           "Unexpected exception was thrown in kernel invocation function.");
+  }
+}
+
+void DependsOnDiscardedEventException(sycl::queue Q) {
+  auto DiscardedEvent =
+      Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+
+  Q.submit([&](sycl::handler &CGH) {
+    try {
+      CGH.depends_on(DiscardedEvent);
+      assert(false && "No exception was thrown.");
+    } catch (const sycl::exception &e) {
+      assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+             "sycl::exception code was not the expected sycl::errc::invalid.");
+    } catch (...) {
+      assert(false &&
+             "Unexpected exception was thrown in kernel invocation function.");
+    }
+    CGH.single_task([] {});
+  });
+
+  sycl::event e1, e2;
+  Q.submit([&](sycl::handler &CGH) {
+    try {
+      CGH.depends_on({e1, DiscardedEvent, e2});
+      assert(false && "No exception was thrown.");
+    } catch (const sycl::exception &e) {
+      assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+             "sycl::exception code was not the expected sycl::errc::invalid.");
+    } catch (...) {
+      assert(false &&
+             "Unexpected exception was thrown in kernel invocation function.");
+    }
+    CGH.single_task([] {});
+  });
+
+  sycl::queue RegularQ;
+  RegularQ.submit([&](sycl::handler &CGH) {
+    try {
+      CGH.depends_on(DiscardedEvent);
+      assert(false && "No exception was thrown.");
+    } catch (const sycl::exception &e) {
+      assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+             "sycl::exception code was not the expected sycl::errc::invalid.");
+    } catch (...) {
+      assert(false &&
+             "Unexpected exception was thrown in kernel invocation function.");
+    }
+    CGH.single_task([] {});
+  });
+
+  RegularQ.submit([&](sycl::handler &CGH) {
+    try {
+      CGH.depends_on({e1, DiscardedEvent, e2});
+      assert(false && "No exception was thrown.");
+    } catch (const sycl::exception &e) {
+      assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+             "sycl::exception code was not the expected sycl::errc::invalid.");
+    } catch (...) {
+      assert(false &&
+             "Unexpected exception was thrown in kernel invocation function.");
+    }
+    CGH.single_task([] {});
+  });
+}
+
+void CheckDiscardedEventAPIException(sycl::queue Q) {
+  DiscardedEventWaitExceptionHelper([&]() {
+    auto DiscardedEvent =
+        Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+    DiscardedEvent.wait();
+  });
+
+  DiscardedEventWaitExceptionHelper([&]() {
+    auto DiscardedEvent =
+        Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+    sycl::event::wait({DiscardedEvent});
+  });
+
+  DiscardedEventWaitExceptionHelper([&]() {
+    auto DiscardedEvent =
+        Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+    DiscardedEvent.wait_and_throw();
+  });
+
+  DiscardedEventWaitExceptionHelper([&]() {
+    auto DiscardedEvent =
+        Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+    sycl::event::wait_and_throw({DiscardedEvent});
+  });
+
+  DiscardedEventWaitExceptionHelper([&]() {
+    auto DiscardedEvent =
+        Q.submit([&](sycl::handler &CGH) { CGH.single_task([] {}); });
+    DiscardedEvent.get_wait_list();
+  });
+}
+
+void CreatingEnableProfilingQueueException(sycl::property_list Props) {
+  try {
+    sycl::queue Q{Props};
+    assert(false && "No exception was thrown.");
+  } catch (const sycl::exception &e) {
+    assert(e.code().value() == static_cast<int>(sycl::errc::invalid) &&
+           "sycl::exception code was not the expected sycl::errc::invalid.");
+  } catch (...) {
+    assert(false &&
+           "Unexpected exception was thrown in kernel invocation function.");
+  }
+}
+
+int main(int Argc, const char *Argv[]) {
+  sycl::property_list Props1{
+      sycl::property::queue::enable_profiling{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  CreatingEnableProfilingQueueException(Props1);
+
+  sycl::property_list Props2{
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue OOO_Queue(Props2);
+  DependsOnDiscardedEventException(OOO_Queue);
+  CheckDiscardedEventAPIException(OOO_Queue);
+
+  sycl::property_list Props3{
+      sycl::property::queue::in_order{},
+      sycl::property::queue::enable_profiling{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  CreatingEnableProfilingQueueException(Props3);
+
+  sycl::property_list Props4{
+      sycl::property::queue::in_order{},
+      sycl::ext::oneapi::property::queue::discard_events{}};
+  sycl::queue Inorder_Queue(Props4);
+  DependsOnDiscardedEventException(Inorder_Queue);
+  CheckDiscardedEventAPIException(Inorder_Queue);
+
+  std::cout << "The test passed." << std::endl;
+  return 0;
+}