[ESIMD] Add tests for atomic operations.

SYCL/ESIMD/dword_atomic_smoke.cpp

@@ -0,0 +1,18 @@
+//==---------------- dword_atomic_smoke.cpp  - DPC++ ESIMD on-device test --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This test checks DWORD atomic operations.
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// This macro disables usage of LSC atomics in the included test.
+#define UNDEF_USE_LSC_ATOMICS
+
+#include "lsc/atomic_smoke.cpp"

SYCL/ESIMD/esimd_test_utils.hpp

@@ -522,4 +522,20 @@ inline void iterate_ops(OpSeq<OpClass, Ops...> ops, F f) {
   ConstexprForLoop<0, sizeof...(Ops)>::unroll(act);
 }
 
+struct USMDeleter {
+  queue Q;
+  void operator()(void *Ptr) {
+    if (Ptr) {
+      sycl::free(Ptr, Q);
+    }
+  }
+};
+
+template <class T>
+std::unique_ptr<T, USMDeleter> usm_malloc_shared(queue q, int n) {
+  std::unique_ptr<T, USMDeleter> res(sycl::malloc_shared<T>(n, q),
+                                     USMDeleter{q});
+  return std::move(res);
+}
+
 } // namespace esimd_test