| 519 | |
| 520 | template<typename Ti, typename Tw, typename To> |
| 521 | To mean_all(CParam<Ti> in) { |
| 522 | using std::unique_ptr; |
| 523 | int in_elements = in.dims[0] * in.dims[1] * in.dims[2] * in.dims[3]; |
| 524 | bool is_linear = (in.strides[0] == 1); |
| 525 | for (int k = 1; k < 4; k++) { |
| 526 | is_linear &= (in.strides[k] == (in.strides[k - 1] * in.dims[k - 1])); |
| 527 | } |
| 528 | |
| 529 | // FIXME: Use better heuristics to get to the optimum number |
| 530 | if (in_elements > 4096 || !is_linear) { |
| 531 | if (is_linear) { |
| 532 | in.dims[0] = in_elements; |
| 533 | for (int k = 1; k < 4; k++) { |
| 534 | in.dims[k] = 1; |
| 535 | in.strides[k] = in_elements; |
| 536 | } |
| 537 | } |
| 538 | |
| 539 | uint threads_x = nextpow2(std::max(32u, (uint)in.dims[0])); |
| 540 | threads_x = std::min(threads_x, THREADS_PER_BLOCK); |
| 541 | uint threads_y = THREADS_PER_BLOCK / threads_x; |
| 542 | |
| 543 | uint blocks_x = divup(in.dims[0], threads_x * REPEAT); |
| 544 | uint blocks_y = divup(in.dims[1], threads_y); |
| 545 | |
| 546 | dim4 outDims(blocks_x, in.dims[1], in.dims[2], in.dims[3]); |
| 547 | |
| 548 | Array<To> tmpOut = createEmptyArray<To>(outDims); |
| 549 | Array<Tw> tmpCt = createEmptyArray<Tw>(outDims); |
| 550 | |
| 551 | Param<Tw> iwt; |
| 552 | mean_first_launcher<Ti, Tw, To>(tmpOut, tmpCt, in, iwt, blocks_x, |
| 553 | blocks_y, threads_x); |
| 554 | |
| 555 | int tmp_elements = tmpOut.elements(); |
| 556 | std::vector<To> h_ptr(tmp_elements); |
| 557 | std::vector<Tw> h_cptr(tmp_elements); |
| 558 | |
| 559 | CUDA_CHECK(cudaMemcpyAsync( |
| 560 | h_ptr.data(), tmpOut.get(), tmp_elements * sizeof(To), |
| 561 | cudaMemcpyDeviceToHost, getStream(getActiveDeviceId()))); |
| 562 | CUDA_CHECK(cudaMemcpyAsync( |
| 563 | h_cptr.data(), tmpCt.get(), tmp_elements * sizeof(Tw), |
| 564 | cudaMemcpyDeviceToHost, getStream(getActiveDeviceId()))); |
| 565 | CUDA_CHECK(cudaStreamSynchronize(getStream(getActiveDeviceId()))); |
| 566 | |
| 567 | compute_t<To> val = static_cast<compute_t<To>>(h_ptr[0]); |
| 568 | compute_t<Tw> weight = static_cast<compute_t<Tw>>(h_cptr[0]); |
| 569 | |
| 570 | for (int i = 1; i < tmp_elements; i++) { |
| 571 | stable_mean(&val, &weight, compute_t<To>(h_ptr[i]), |
| 572 | compute_t<Tw>(h_cptr[i])); |
| 573 | } |
| 574 | |
| 575 | return static_cast<To>(val); |
| 576 | } else { |
| 577 | std::vector<Ti> h_ptr(in_elements); |
| 578 |
nothing calls this directly
no test coverage detected