[SYCL] Fix tanhf to support lstm

wangzy0327 · Oct 28, 2022 · c32e75b · c32e75b
1 parent 2349d35
commit c32e75b
Show file tree

Hide file tree

Showing 2 changed files with 4 additions and 2 deletions.
diff --git a/src/runtime/sycl/sycl_module.cc b/src/runtime/sycl/sycl_module.cc
@@ -55,10 +55,10 @@ class SYCLWrappedFunc {
  void operator()(TVMArgs args, TVMRetValue* rv, void** void_args) const {
  VLOG(0) << "enter sycl wrapped func operator()";
  ICHECK(w_->devices.size() != 0) << "No SYCL device";
-
+ // get kernel
  void (*kernel_func)(sycl::queue &Q, sycl::range<3> k0_dimGrid, sycl::range<3> k0_dimBlock, void** void_args) = (void (*)(sycl::queue &Q, sycl::range<3> k0_dimGrid, sycl::range<3> k0_dimBlock, void** void_args))dlsym(so_handler_, func_name_.c_str());
  ICHECK(kernel_func != NULL) << "ERROR:"<<dlerror()<<":dlsym\n";
-
+ // get thread dimension
  ThreadWorkLoad wl = launch_param_config_.Extract(args);
  sycl::range<3> k0_dimGrid(wl.grid_dim(0), wl.grid_dim(1), wl.grid_dim(2));
  sycl::range<3> k0_dimBlock(wl.block_dim(0), wl.block_dim(1), wl.block_dim(2));

diff --git a/src/target/source/codegen_sycl.cc b/src/target/source/codegen_sycl.cc
@@ -621,6 +621,8 @@ void CodeGenSYCL::VisitExpr_(const CallNode* op, std::ostream& os) {
  func = StringImm("exp");
  }else if(func->value == "powf"){
  func = StringImm("pow");
+ }else if(func->value == "tanhf"){
+ func = StringImm("tanh");
  }
  this->PrintCallExtern(GetType(GetRef<PrimExpr>(op)), func->value, op->args, true, os);
  } else {