special for weight preprocess
| 162 | |
| 163 | //! special for weight preprocess |
| 164 | void exec_convolution(ConvolutionForward* opr0, ConvolutionForward* opr1) { |
| 165 | opr0->param().pad_h = pad_h; |
| 166 | opr0->param().pad_w = pad_w; |
| 167 | opr1->param() = opr0->param(); |
| 168 | opr1->param().sparse = param::Convolution::Sparse::GROUP; |
| 169 | |
| 170 | TensorND a0, b0, c0, a1, b1, c1; |
| 171 | std::tie(a0, b0, c0) = shuffle( |
| 172 | std::make_tuple(src0->tensornd(), flt0->tensornd(), dst0->tensornd())); |
| 173 | std::tie(a1, b1, c1) = shuffle( |
| 174 | std::make_tuple(src1->tensornd(), flt1->tensornd(), dst1->tensornd())); |
| 175 | WorkspaceWrapper wk( |
| 176 | handle, std::max( |
| 177 | opr0->get_workspace_in_bytes( |
| 178 | a0.layout, b0.layout, c0.layout, nullptr), |
| 179 | opr1->get_workspace_in_bytes( |
| 180 | a1.layout, b1.layout, c1.layout, nullptr))); |
| 181 | cudaProfilerStart(); |
| 182 | cudaEventRecord(cuda_ev[0], cuda_stream); |
| 183 | opr0->exec(a0, b0, c0, nullptr, wk.workspace()); |
| 184 | cudaEventRecord(cuda_ev[1], cuda_stream); |
| 185 | opr1->exec(a1, b1, c1, nullptr, wk.workspace()); |
| 186 | cudaEventRecord(cuda_ev[2], cuda_stream); |
| 187 | cudaProfilerStop(); |
| 188 | |
| 189 | if (getenv("MEGDNN_CHANWISE_CONV_VERBOSE") || |
| 190 | getenv("MEGDNN_CHANWISE_CONV_FULLBENCH")) { |
| 191 | cudaStreamSynchronize(cuda_stream); |
| 192 | float t0 = -1, t1 = -1; |
| 193 | cudaEventElapsedTime(&t0, cuda_ev[0], cuda_ev[1]); |
| 194 | cudaEventElapsedTime(&t1, cuda_ev[1], cuda_ev[2]); |
| 195 | printf("%s;%s;%s: cudnn/megdnn: %.3fms/%.3fms=%.3f\n", |
| 196 | lsrc.TensorShape::to_string().c_str(), |
| 197 | lflt1.TensorShape::to_string().c_str(), |
| 198 | ldst.TensorShape::to_string().c_str(), t0, t1, t0 / t1); |
| 199 | } |
| 200 | } |
| 201 | |
| 202 | void cmp_dst() { |
| 203 | Tensor<> dst0_cpu(handle_cpu, ldst), dst1_cpu(handle_cpu, ldst); |