From a9b4df0ad0144a2ef17e5b5347e63ca7efaf4bc4 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 17 Jun 2021 21:21:19 -0700
Subject: [PATCH] SWDEV-285388 - Allow fallback to asm when builtin not
 available

Change-Id: Ifeee7f3ad1be9ce9b79cf26008b75102153143a5
---
 hipamd/include/hip/amd_detail/llvm_intrinsics.h | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/hipamd/include/hip/amd_detail/llvm_intrinsics.h b/hipamd/include/hip/amd_detail/llvm_intrinsics.h
index 8c626f7616..5494b69f8b 100644
--- a/hipamd/include/hip/amd_detail/llvm_intrinsics.h
+++ b/hipamd/include/hip/amd_detail/llvm_intrinsics.h
@@ -30,11 +30,17 @@ THE SOFTWARE.
 
 #include "hip/amd_detail/host_defines.h"
 
+#if __has_builtin(__builtin_amdgcn_groupstaticsize)
 __device__
 inline
 unsigned __llvm_amdgcn_groupstaticsize() {
   return __builtin_amdgcn_groupstaticsize();
 }
+#else
+// FIXME: This should be removed once the above builtin is available.
+__device__
+unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");
+#endif
 
 template<typename __T>
 __device__