From 6f78b4300f6f29c9e71ad646d58752c0416d4bc2 Mon Sep 17 00:00:00 2001
From: Avinash Kethineedi <avinash.kethineedi@amd.com>
Date: Fri, 21 Mar 2025 14:31:49 -0500
Subject: [PATCH] Remove support code for GFX940 and GFX941 targets (#55)

[ROCm/rocshmem commit: 1210b6419fe6fa8948542016747182e292e181d3]
---
 projects/rocshmem/src/assembly.hpp | 22 +++++++++++-----------
 1 file changed, 11 insertions(+), 11 deletions(-)

diff --git a/projects/rocshmem/src/assembly.hpp b/projects/rocshmem/src/assembly.hpp
index a14907a0bd..72720ef75f 100644
--- a/projects/rocshmem/src/assembly.hpp
+++ b/projects/rocshmem/src/assembly.hpp
@@ -51,7 +51,7 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
       : "=v"(ret)
       : "v"(src));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
   asm volatile(
       "global_load_ubyte %0 %1 off sc0 sc1 \n"
       "s_waitcnt vmcnt(0)"
@@ -74,7 +74,7 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va
     : "=v"(*assigned_value)
     : "v"(read_value));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
   asm volatile(
     "global_load_sbyte %0 %1 off sc0 sc1\n "
     "s_waitcnt vmcnt(0)"
@@ -96,7 +96,7 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi
     : "=v"(*assigned_value)
     : "v"(read_value));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
   asm volatile(
     "global_load_dwordx2 %0 %1 off sc0 sc1\n "
     "s_waitcnt vmcnt(0)"
@@ -127,7 +127,7 @@ NOWARN(-Wdeprecated-volatile,
             : "=v"(ret)
             : "v"(src));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
         asm volatile(
             "global_load_dword %0 %1 off sc0 sc1 \n"
             "s_waitcnt vmcnt(0)"
@@ -147,7 +147,7 @@ NOWARN(-Wdeprecated-volatile,
             : "=v"(ret)
             : "v"(src));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
         asm volatile(
             "global_load_dwordx2 %0 %1 off sc0 sc1 \n"
             "s_waitcnt vmcnt(0)"
@@ -172,7 +172,7 @@ __device__ __forceinline__ void __roc_inv() {
 #if defined(__gfx90a__)
 //  asm volatile("buffer_wbinvl1;");
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
 //  asm volatile("buffer_inv sc0 sc1;");
 #endif
 #endif
@@ -188,7 +188,7 @@ __device__ __forceinline__ void __roc_flush() {
 //  asm volatile("s_dcache_wb;");
 //  asm volatile("buffer_wbl2;");
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
 //  asm volatile("s_dcache_wb;");
 //  asm volatile("buffer_wbl2;");
 #endif
@@ -207,7 +207,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #if defined(__gfx90a__)
       asm volatile("flat_store_short %0 %1 glc slc" : : "v"(dst), "v"(val16));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
       asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16));
 #endif
       break;
@@ -221,7 +221,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #if defined(__gfx90a__)
       asm volatile("flat_store_dword %0 %1 glc slc" : : "v"(dst), "v"(val32));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
       asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32));
 #endif
       break;
@@ -235,7 +235,7 @@ __device__ __forceinline__ void store_asm(uint8_t* val, uint8_t* dst,
 #if defined(__gfx90a__)
       asm volatile("flat_store_dwordx2 %0 %1 glc slc" : : "v"(dst), "v"(val64));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
       asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64));
 #endif
       break;
@@ -257,7 +257,7 @@ __device__ __forceinline__ uint64_t __read_clock() {
       "s_waitcnt lgkmcnt(0)\n"
       : "=s"(clock));
 #endif
-#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)
+#if defined(__gfx942__)
   asm volatile(
       "s_memrealtime %0\n"
       "s_waitcnt lgkmcnt(0)\n"