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