[SYCLomatic][ASM] Support migration of instruction st.volatile.global.u32 and ld.volatile.global.u32 (#2873)

Signed-off-by: chenwei.sun <chenwei.sun@intel.com>

Author: tomflinda

clang/lib/DPCT/RulesAsm/Parser/AsmTokenKinds.def

@@ -450,6 +450,7 @@ MODIFIER(cs, ".cs")
 MODIFIER(to, ".to")
 MODIFIER(aligned, ".aligned")
 MODIFIER(trans, ".trans")
+MODIFIER(ptx_volatile, ".volatile")
 
 #undef LINKAGE
 #undef TARGET

clang/test/dpct/asm/ld.cu

@@ -70,4 +70,27 @@ __device__ inline void load_global_short4(short4 &a, const short4 *addr) {
   a.w = w;
 }
 
+// CHECK: __dpct_inline__ int ld_flag_volatile(int* flag_addr) {
+// CHECK-NEXT:   int flag;
+// CHECK-NEXT:   flag = *((uint32_t *)(uintptr_t)flag_addr);
+// CHECK-NEXT:   sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::device);
+// CHECK-NEXT:   return flag;
+// CHECK-NEXT: }
+__device__ __forceinline__ int ld_flag_volatile(int* flag_addr) {
+  int flag;
+  asm volatile("ld.volatile.global.u32 %0, [%1]; membar.gl;" : "=r"(flag) : "l"(flag_addr));
+  return flag;
+}
+
+// CHECK: __dpct_inline__ int ld_flag_acquire(int* flag_addr) {
+// CHECK-NEXT:  int flag;
+// CHECK-NEXT:  flag = *((uint32_t *)(uintptr_t)flag_addr);
+// CHECK-NEXT:  return flag;
+// CHECK-NEXT: }
+__device__ __forceinline__ int ld_flag_acquire(int* flag_addr) {
+  int flag;
+  asm volatile("ld.volatile.global.u32 %0, [%1];" : "=r"(flag) : "l"(flag_addr));
+  return flag;
+}
+
 // clang-format on

clang/test/dpct/asm/st.cu

@@ -61,4 +61,19 @@ __device__ inline void store_streaming_short2(short2 *addr, short x, short y) {
     asm("st.cs.global.v2.s16 [%0+0], {%1, %2};" ::__PTR(addr), "h"(x), "h"(y));
 }
 
+// CHECK: __dpct_inline__ void st_flag_release(int* flag_addr, int flag) {
+// CHECK-NEXT:  sycl::atomic_fence(sycl::memory_order::seq_cst,sycl::memory_scope::system);
+// CHECK-NEXT:  *((uint32_t *)(uintptr_t)flag_addr) = flag;
+// CHECK-NEXT: }
+__device__ __forceinline__ void st_flag_release(int* flag_addr, int flag) {
+  asm volatile("membar.sys; st.volatile.global.u32 [%1], %0;" ::"r"(flag), "l"(flag_addr));
+}
+
+// CHECK: __dpct_inline__ void st_flag_volatile(int* flag_addr, int flag) {
+// CHECK-NEXT:   *((uint32_t *)(uintptr_t)flag_addr) = flag;
+// CHECK-NEXT: }
+__device__ __forceinline__ void st_flag_volatile(int* flag_addr, int flag) {
+  asm volatile("st.volatile.global.u32 [%1], %0;" ::"r"(flag), "l"(flag_addr));
+}
+
 // clang-format on