diff --git a/hpvm/test/dnn_benchmarks/common/include/visc.h b/hpvm/test/dnn_benchmarks/common/include/visc.h index 9bf77b6ed654a6b309a8685acd51882ea1d95a5c..a354d7c096d2c02708b51536f47389ece59771a5 100644 --- a/hpvm/test/dnn_benchmarks/common/include/visc.h +++ b/hpvm/test/dnn_benchmarks/common/include/visc.h @@ -10,110 +10,110 @@ #define DEVICE CUDNN_TARGET #endif -#include "llvm/SupportVISC/VISCHint.h" +#include "SupportHPVM/HPVMHint.h" #ifdef __cplusplus extern "C" { -void __visc__hint(visc::Target); +void __hpvm__hint(visc::Target); #else -void __visc__hint(enum Target); +void __hpvm__hint(enum Target); #endif #ifdef __cplusplus -void* __visc__node(...); -//void* __visc__createNode(...); -//void* __visc__createNode1D(...); -//void* __visc__createNode2D(...); -//void* __visc__createNode3D(...); -//void __visc__return(...); +void* __hpvm__node(...); +//void* __hpvm__createNode(...); +//void* __hpvm__createNode1D(...); +//void* __hpvm__createNode2D(...); +//void* __hpvm__createNode3D(...); +//void __hpvm__return(...); #endif -void* __visc__createNodeND(unsigned, ...); -void __visc__return(unsigned, ...); - -void __visc__attributes(unsigned, ...); -void __visc__init(); -void __visc__cleanup(); - -void __visc__bindIn(void*, unsigned, unsigned, unsigned); -void __visc__bindOut(void*, unsigned, unsigned, unsigned); -void* __visc__edge(void*, void*, unsigned, unsigned, unsigned, unsigned); -void __visc__push(void*, void*); -void* __visc__pop(void*); -void* __visc__launch(unsigned, ...); -void __visc__wait(void*); - -void* __visc__getNode(); -void* __visc__getParentNode(void*); -void __visc__barrier(); -void* __visc__malloc(long); -long __visc__getNodeInstanceID_x(void*); -long __visc__getNodeInstanceID_y(void*); -long __visc__getNodeInstanceID_z(void*); -long __visc__getNumNodeInstances_x(void*); -long __visc__getNumNodeInstances_y(void*); -long __visc__getNumNodeInstances_z(void*); +void* __hpvm__createNodeND(unsigned, ...); +void __hpvm__return(unsigned, ...); + +void __hpvm__attributes(unsigned, ...); +void __hpvm__init(); +void __hpvm__cleanup(); + +void __hpvm__bindIn(void*, unsigned, unsigned, unsigned); +void __hpvm__bindOut(void*, unsigned, unsigned, unsigned); +void* __hpvm__edge(void*, void*, unsigned, unsigned, unsigned, unsigned); +void __hpvm__push(void*, void*); +void* __hpvm__pop(void*); +void* __hpvm__launch(unsigned, ...); +void __hpvm__wait(void*); + +void* __hpvm__getNode(); +void* __hpvm__getParentNode(void*); +void __hpvm__barrier(); +void* __hpvm__malloc(long); +long __hpvm__getNodeInstanceID_x(void*); +long __hpvm__getNodeInstanceID_y(void*); +long __hpvm__getNodeInstanceID_z(void*); +long __hpvm__getNumNodeInstances_x(void*); +long __hpvm__getNumNodeInstances_y(void*); +long __hpvm__getNumNodeInstances_z(void*); // Atomic // signed int -int __visc__atomic_cmpxchg(int*, int, int); -int __visc__atomic_add(int*, int); -int __visc__atomic_sub(int*, int); -int __visc__atomic_xchg(int*, int); -int __visc__atomic_inc(int*); -int __visc__atomic_dec(int*); -int __visc__atomic_min(int*, int); -int __visc__atomic_max(int*, int); -int __visc__atomic_umax(int*, int); -int __visc__atomic_umin(int*, int); -int __visc__atomic_and(int*, int); -int __visc__atomic_or(int*, int); -int __visc__atomic_xor(int*, int); +int __hpvm__atomic_cmpxchg(int*, int, int); +int __hpvm__atomic_add(int*, int); +int __hpvm__atomic_sub(int*, int); +int __hpvm__atomic_xchg(int*, int); +int __hpvm__atomic_inc(int*); +int __hpvm__atomic_dec(int*); +int __hpvm__atomic_min(int*, int); +int __hpvm__atomic_max(int*, int); +int __hpvm__atomic_umax(int*, int); +int __hpvm__atomic_umin(int*, int); +int __hpvm__atomic_and(int*, int); +int __hpvm__atomic_or(int*, int); +int __hpvm__atomic_xor(int*, int); // Special Func -float __visc__floor(float); -float __visc__rsqrt(float); -float __visc__sqrt(float); -float __visc__sin(float); -float __visc__cos(float); +float __hpvm__floor(float); +float __hpvm__rsqrt(float); +float __hpvm__sqrt(float); +float __hpvm__sin(float); +float __hpvm__cos(float); // unsigned int -//unsigned __visc__atomic_cmpxchg(unsigned*, unsigned, unsigned); -//unsigned __visc__atomic_add(unsigned*, unsigned); -//unsigned __visc__atomic_sub(unsigned*, unsigned); -//unsigned __visc__atomic_xchg(unsigned*, unsigned); -//unsigned __visc__atomic_inc(unsigned*); -//unsigned __visc__atomic_dec(unsigned*); -//unsigned __visc__atomic_min(unsigned*, unsigned); -//unsigned __visc__atomic_max(unsigned*, unsigned); -//unsigned __visc__atomic_and(unsigned*, unsigned); -//unsigned __visc__atomic_or(unsigned*, unsigned); -//unsigned __visc__atomic_xor(unsigned*, unsigned); +//unsigned __hpvm__atomic_cmpxchg(unsigned*, unsigned, unsigned); +//unsigned __hpvm__atomic_add(unsigned*, unsigned); +//unsigned __hpvm__atomic_sub(unsigned*, unsigned); +//unsigned __hpvm__atomic_xchg(unsigned*, unsigned); +//unsigned __hpvm__atomic_inc(unsigned*); +//unsigned __hpvm__atomic_dec(unsigned*); +//unsigned __hpvm__atomic_min(unsigned*, unsigned); +//unsigned __hpvm__atomic_max(unsigned*, unsigned); +//unsigned __hpvm__atomic_and(unsigned*, unsigned); +//unsigned __hpvm__atomic_or(unsigned*, unsigned); +//unsigned __hpvm__atomic_xor(unsigned*, unsigned); /* * ApproxHPVM specific function calls */ -void* __visc__tensor_add(void*, void*); -void* __visc__tensor_mul(void*, void*); -void* __visc__tensor_convolution(void*, void*, int, int, int, int); -void* __visc__tensor_group_convolution(void*, void*, int, int, int, int, int, int); -void* __visc__tensor_batchnorm(void*, void*, void*, void*, void*, double); -void* __visc__tensor_pool_max(void*, int, int, int, int, int, int); -void* __visc__tensor_pool_mean(void*, int, int, int, int, int, int); -void* __visc__tensor_relu(void*); -void* __visc__tensor_tanh(void*); -void* __visc__tensor_softmax(void*); +void* __hpvm__tensor_add(void*, void*); +void* __hpvm__tensor_mul(void*, void*); +void* __hpvm__tensor_convolution(void*, void*, int, int, int, int); +void* __hpvm__tensor_group_convolution(void*, void*, int, int, int, int, int, int); +void* __hpvm__tensor_batchnorm(void*, void*, void*, void*, void*, double); +void* __hpvm__tensor_pool_max(void*, int, int, int, int, int, int); +void* __hpvm__tensor_pool_mean(void*, int, int, int, int, int, int); +void* __hpvm__tensor_relu(void*); +void* __hpvm__tensor_tanh(void*); +void* __hpvm__tensor_softmax(void*); // Tensor ops for image processing -void* __visc__tensor_fft(void*); -void* __visc__tensor_reduce(void*, int, void*); -void* __visc__tensor_projectiveT(void*, void*); -void* __visc__tensor_map1(void*, void*); -void* __visc__tensor_map2(void*, void*, void*); -void* __visc__tensor_map3(void*, void*, void*, void*); -void* __visc__tensor_cosineT(void*); -void* __visc__tensor_stencil(void*); +void* __hpvm__tensor_fft(void*); +void* __hpvm__tensor_reduce(void*, int, void*); +void* __hpvm__tensor_projectiveT(void*, void*); +void* __hpvm__tensor_map1(void*, void*); +void* __hpvm__tensor_map2(void*, void*, void*); +void* __hpvm__tensor_map3(void*, void*, void*, void*); +void* __hpvm__tensor_cosineT(void*); +void* __hpvm__tensor_stencil(void*); // New HPVM intrinsic for Setting Node ID -void* __visc__node_id(int); +void* __hpvm__node_id(int); #include <unistd.h> @@ -124,12 +124,12 @@ long get_local_id(int); long get_local_size(int); -void llvm_visc_track_mem(void*, size_t); -void llvm_visc_untrack_mem(void*); -void llvm_visc_request_mem(void*, size_t); +void llvm_hpvm_track_mem(void*, size_t); +void llvm_hpvm_untrack_mem(void*); +void llvm_hpvm_request_mem(void*, size_t); /*#ifdef __cplusplus -void hpvm_request_tensor(void*, visc::Target); +void hpvm_request_tensor(void*, hpvm::Target); #else void hpvm_request_tensor(void*, enum Target); #endif