From 5ad42402d6cd2894d32fe949b0dd261e39552f73 Mon Sep 17 00:00:00 2001
From: Konstantin Zhuravlyov <kzhuravl_dev@outlook.com>
Date: Thu, 18 Jun 2020 12:41:25 -0400
Subject: [PATCH] Add gfx1030

Change-Id: Icb3cbd79540dac03bd729ac93958c5b512d2ee3d
---
 ockl/src/dots.cl             |  1 +
 ockl/src/mtime.cl            |  2 +-
 oclc/src/isa_version_1030.cl | 11 +++++++++++
 3 files changed, 13 insertions(+), 1 deletion(-)
 create mode 100644 oclc/src/isa_version_1030.cl

diff --git a/ockl/src/dots.cl b/ockl/src/dots.cl
index da0d19ed..c2cb3c0f 100644
--- a/ockl/src/dots.cl
+++ b/ockl/src/dots.cl
@@ -33,6 +33,7 @@ extern __attribute__((const)) int __llvm_amdgcn_sdot8(int a, int b, int c, bool
 extern __attribute__((const)) uint __llvm_amdgcn_udot8(uint a, uint b, uint c, bool s) __asm("llvm.amdgcn.udot8");
 
 #define SWDOT __oclc_ISA_version < 9006 || __oclc_ISA_version == 9009 || __oclc_ISA_version == 10100
+
 #define AS_INT(X) __builtin_astype(X, int)
 #define AS_UINT(X) __builtin_astype(X, uint)
 #define ATTR __attribute__((const))
diff --git a/ockl/src/mtime.cl b/ockl/src/mtime.cl
index 2b77cdb9..543aaa30 100644
--- a/ockl/src/mtime.cl
+++ b/ockl/src/mtime.cl
@@ -7,7 +7,7 @@
 
 #include "ockl.h"
 
-ulong
+__attribute__((target("s-memtime-inst"))) ulong
 OCKL_MANGLE_U64(memtime)(void)
 {
     return __builtin_amdgcn_s_memtime();
diff --git a/oclc/src/isa_version_1030.cl b/oclc/src/isa_version_1030.cl
new file mode 100644
index 00000000..117645f6
--- /dev/null
+++ b/oclc/src/isa_version_1030.cl
@@ -0,0 +1,11 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ISA_version = 10300;
+