diff --git a/llvm/include/llvm/IR/IntrinsicsVISC.td b/llvm/include/llvm/IR/IntrinsicsVISC.td index ab22372d80ac6ae1de89c2da6c3402fd0787e181..404903648f8656e00d283d50b27fd411cb930ac8 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 faab312087eade9f5d1e622555de270260f3259d..65e09e716c2deb0f1b3f475e4a22cf16cd25c486 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 978f29e85a9aae8e310f528bc8fa7630c85d8542..9bf77b6ed654a6b309a8685acd51882ea1d95a5c 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);