From 9a30780cca2b1e9a910d7e188cb73f353eae325e Mon Sep 17 00:00:00 2001
From: Hashim Sharif <hsharif3@tyler.cs.illinois.edu>
Date: Sat, 21 Nov 2020 00:33:21 -0600
Subject: [PATCH] Adding __visc_node_id intrinsic - untested
---
llvm/include/llvm/IR/IntrinsicsVISC.td | 3 +++
llvm/lib/Transforms/GenVISC/GenVISC.cpp | 9 +++++++++
llvm/test/VISC/DNN_Benchmarks/common/include/visc.h | 4 ++++
3 files changed, 16 insertions(+)
diff --git a/llvm/include/llvm/IR/IntrinsicsVISC.td b/llvm/include/llvm/IR/IntrinsicsVISC.td
index ab22372d80..404903648f 100644
--- a/llvm/include/llvm/IR/IntrinsicsVISC.td
+++ b/llvm/include/llvm/IR/IntrinsicsVISC.td
@@ -325,4 +325,7 @@ let TargetPrefix = "visc" in {
llvm_i32_ty,
llvm_i32_ty], []>;
+ def int_visc_node_id : Intrinsic<[llvm_ptr_ty], [llvm_i32_ty], []>;
+
+
}
diff --git a/llvm/lib/Transforms/GenVISC/GenVISC.cpp b/llvm/lib/Transforms/GenVISC/GenVISC.cpp
index faab312087..65e09e716c 100644
--- a/llvm/lib/Transforms/GenVISC/GenVISC.cpp
+++ b/llvm/lib/Transforms/GenVISC/GenVISC.cpp
@@ -178,6 +178,9 @@ IS_VISC_CALL(tensor_tanh)
IS_VISC_CALL(tensor_sigmoid)
IS_VISC_CALL(tensor_softmax)
+IS_VISC_CALL(node_id)
+
+
// Return the constant integer represented by value V
static unsigned getNumericValue(Value* V) {
assert(isa<ConstantInt>(V)
@@ -1308,6 +1311,12 @@ bool GenVISC::runOnModule(Module &M) {
if (isVISCCall_tensor_softmax(I)) {
ReplaceCallWithIntrinsic(I, Intrinsic::visc_tensor_softmax, &toBeErased);
}
+
+ // New Intrinsic to set Node ID
+ if (isVISCCall_node_id(I)) {
+ ReplaceCallWithIntrinsic(I, Intrinsic::visc_node_id, &toBeErased);
+ }
+
}
// Erase the __visc__node calls
diff --git a/llvm/test/VISC/DNN_Benchmarks/common/include/visc.h b/llvm/test/VISC/DNN_Benchmarks/common/include/visc.h
index 978f29e85a..9bf77b6ed6 100644
--- a/llvm/test/VISC/DNN_Benchmarks/common/include/visc.h
+++ b/llvm/test/VISC/DNN_Benchmarks/common/include/visc.h
@@ -112,6 +112,10 @@ void* __visc__tensor_map3(void*, void*, void*, void*);
void* __visc__tensor_cosineT(void*);
void* __visc__tensor_stencil(void*);
+// New HPVM intrinsic for Setting Node ID
+void* __visc__node_id(int);
+
+
#include <unistd.h>
long get_global_id(int);
--
GitLab