| 120 | |
| 121 | template <class Opr> |
| 122 | void exec(Opr* opr0, Opr* opr1) { |
| 123 | opr0->param().pad_d = pad_d; |
| 124 | opr0->param().pad_h = pad_h; |
| 125 | opr0->param().pad_w = pad_w; |
| 126 | opr1->param() = opr0->param(); |
| 127 | opr1->param().sparse = param::Convolution3D::Sparse::GROUP; |
| 128 | |
| 129 | TensorND a0, b0, c0, a1, b1, c1; |
| 130 | std::tie(a0, b0, c0) = shuffle( |
| 131 | std::make_tuple(src0->tensornd(), flt0->tensornd(), dst0->tensornd())); |
| 132 | std::tie(a1, b1, c1) = shuffle( |
| 133 | std::make_tuple(src1->tensornd(), flt1->tensornd(), dst1->tensornd())); |
| 134 | WorkspaceWrapper wk( |
| 135 | handle, |
| 136 | std::max( |
| 137 | opr0->get_workspace_in_bytes(a0.layout, b0.layout, c0.layout), |
| 138 | opr1->get_workspace_in_bytes(a1.layout, b1.layout, c1.layout))); |
| 139 | cudaProfilerStart(); |
| 140 | cudaEventRecord(cuda_ev[0], cuda_stream); |
| 141 | opr0->exec(a0, b0, c0, wk.workspace()); |
| 142 | cudaEventRecord(cuda_ev[1], cuda_stream); |
| 143 | opr1->exec(a1, b1, c1, wk.workspace()); |
| 144 | cudaEventRecord(cuda_ev[2], cuda_stream); |
| 145 | cudaProfilerStop(); |
| 146 | |
| 147 | if (getenv("MEGDNN_CHANWISE_CONV3D_VERBOSE") || |
| 148 | getenv("MEGDNN_CHANWISE_CONV3D_FULLBENCH")) { |
| 149 | cudaStreamSynchronize(cuda_stream); |
| 150 | float t0 = -1, t1 = -1; |
| 151 | cudaEventElapsedTime(&t0, cuda_ev[0], cuda_ev[1]); |
| 152 | cudaEventElapsedTime(&t1, cuda_ev[1], cuda_ev[2]); |
| 153 | printf("%s;%s;%s: cudnn/megdnn: %.3fms/%.3fms=%.3f\n", |
| 154 | lsrc.TensorShape::to_string().c_str(), |
| 155 | lflt1.TensorShape::to_string().c_str(), |
| 156 | ldst.TensorShape::to_string().c_str(), t0, t1, t0 / t1); |
| 157 | } |
| 158 | } |
| 159 | |
| 160 | void cmp_dst() { |
| 161 | Tensor<> dst0_cpu(handle_cpu, ldst), dst1_cpu(handle_cpu, ldst); |
no test coverage detected