From e2bae5131afc7ad1ad44c3ee64cea1177a4370d9 Mon Sep 17 00:00:00 2001
From: akolliasAMD <99202231+akolliasAMD@users.noreply.github.com>
Date: Wed, 25 Jun 2025 17:35:54 -0600
Subject: [PATCH] added new gfx target (#171)

---
 CMakeLists.txt   |  4 +++-
 src/assembly.hpp | 18 +++++++++---------
 2 files changed, 12 insertions(+), 10 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 5614bb68ab..b4b5faa9ad 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -109,7 +109,9 @@ if (NOT BUILD_TESTS_ONLY)
       gfx90a:xnack-;
       gfx90a:xnack+;
       gfx942:xnack-;
-      gfx942:xnack+)
+      gfx942:xnack+;
+      gfx950:xnack-;
+      gfx950:xnack+)
 
   if (BUILD_LOCAL_GPU_TARGET_ONLY)
     message(STATUS "Building only for local GPU target")
diff --git a/src/assembly.hpp b/src/assembly.hpp
index aa48099ac1..67827f3e9d 100644
--- a/src/assembly.hpp
+++ b/src/assembly.hpp
@@ -53,7 +53,7 @@ __device__ __forceinline__ int uncached_load_ubyte(uint8_t* src) {
       : "=v"(ret)
       : "v"(src));
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
   asm volatile(
       "global_load_ubyte %0 %1 off sc0 sc1 \n"
       "s_waitcnt vmcnt(0)"
@@ -76,7 +76,7 @@ __device__ __forceinline__ void refresh_volatile_sbyte(volatile int *assigned_va
     : "=v"(*assigned_value)
     : "v"(read_value));
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
   asm volatile(
     "global_load_sbyte %0 %1 off sc0 sc1\n "
     "s_waitcnt vmcnt(0)"
@@ -98,7 +98,7 @@ __device__ __forceinline__ void refresh_volatile_dwordx2(volatile uint64_t *assi
     : "=v"(*assigned_value)
     : "v"(read_value));
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
   asm volatile(
     "global_load_dwordx2 %0 %1 off sc0 sc1\n "
     "s_waitcnt vmcnt(0)"
@@ -129,7 +129,7 @@ NOWARN(-Wdeprecated-volatile,
             : "=v"(ret)
             : "v"(src));
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
         asm volatile(
             "global_load_dword %0 %1 off sc0 sc1 \n"
             "s_waitcnt vmcnt(0)"
@@ -149,7 +149,7 @@ NOWARN(-Wdeprecated-volatile,
             : "=v"(ret)
             : "v"(src));
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
         asm volatile(
             "global_load_dwordx2 %0 %1 off sc0 sc1 \n"
             "s_waitcnt vmcnt(0)"
@@ -175,7 +175,7 @@ __device__ __forceinline__ void __roc_flush() {
 //  asm volatile("s_dcache_wb;");
 //  asm volatile("buffer_wbl2;");
 #endif
-#if defined(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
 //  asm volatile("s_dcache_wb;");
 //  asm volatile("buffer_wbl2;");
 #endif
@@ -194,7 +194,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(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
       asm volatile("flat_store_short %0 %1 sc0 sc1" : : "v"(dst), "v"(val16));
 #endif
       break;
@@ -208,7 +208,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(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
       asm volatile("flat_store_dword %0 %1 sc0 sc1" : : "v"(dst), "v"(val32));
 #endif
       break;
@@ -222,7 +222,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(__gfx942__)
+#if defined(__gfx942__) || defined(__gfx950__)
       asm volatile("flat_store_dwordx2 %0 %1 sc0 sc1" : : "v"(dst), "v"(val64));
 #endif
       break;