| 66 | } |
| 67 | |
| 68 | void LocalShareForwardImpl::AlgoCHWNBatchSizeAware::exec(const ExecArgs& args) const { |
| 69 | local_share::Param kern_param; |
| 70 | auto&& param = args.opr->param(); |
| 71 | unpack_local_share_params( |
| 72 | args.src_layout, args.filter_layout, args.dst_layout, param); |
| 73 | kern_param.n = n, kern_param.co = co, kern_param.ci = ci, kern_param.hi = hi, |
| 74 | kern_param.wi = wi, kern_param.ph = ph, kern_param.pw = pw, |
| 75 | kern_param.grp_ho = ho / sgh, kern_param.grp_wo = wo / sgw, kern_param.sgh = sgh, |
| 76 | kern_param.sgw = sgw; |
| 77 | auto&& handle = concrete_handle(args.opr->handle()); |
| 78 | auto&& cublas_hdl = cublas_handle(args.opr->handle()); |
| 79 | auto&& stream = cuda_stream(args.opr->handle()); |
| 80 | |
| 81 | auto one = handle->one_device(); |
| 82 | auto zero = handle->zero_device(); |
| 83 | |
| 84 | local_share::_do_local_share_convolution_large_batch_size( |
| 85 | args.src_tensor->ptr<dt_float32>(), args.filter_tensor->ptr<dt_float32>(), |
| 86 | args.dst_tensor->ptr<dt_float32>(), |
| 87 | reinterpret_cast<float*>(args.workspace.raw_ptr), fh, fw, sh, sw, |
| 88 | kern_param, cublas_hdl, stream, one, zero); |
| 89 | } |
| 90 | |
| 91 | // vim: syntax=cpp.doxygen |
nothing calls this directly
no test coverage detected