| 431 | |
| 432 | template<typename T, typename Tw> |
| 433 | T mean_all_weighted(CParam<T> in, CParam<Tw> iwt) { |
| 434 | int in_elements = in.dims[0] * in.dims[1] * in.dims[2] * in.dims[3]; |
| 435 | |
| 436 | // FIXME: Use better heuristics to get to the optimum number |
| 437 | if (in_elements > 4096) { |
| 438 | bool in_is_linear = (in.strides[0] == 1); |
| 439 | bool wt_is_linear = (iwt.strides[0] == 1); |
| 440 | for (int k = 1; k < 4; k++) { |
| 441 | in_is_linear &= |
| 442 | (in.strides[k] == (in.strides[k - 1] * in.dims[k - 1])); |
| 443 | wt_is_linear &= |
| 444 | (iwt.strides[k] == (iwt.strides[k - 1] * iwt.dims[k - 1])); |
| 445 | } |
| 446 | |
| 447 | if (in_is_linear && wt_is_linear) { |
| 448 | in.dims[0] = in_elements; |
| 449 | for (int k = 1; k < 4; k++) { |
| 450 | in.dims[k] = 1; |
| 451 | in.strides[k] = in_elements; |
| 452 | } |
| 453 | |
| 454 | for (int k = 0; k < 4; k++) { |
| 455 | iwt.dims[k] = in.dims[k]; |
| 456 | iwt.strides[k] = in.strides[k]; |
| 457 | } |
| 458 | } |
| 459 | |
| 460 | uint threads_x = nextpow2(std::max(32u, (uint)in.dims[0])); |
| 461 | threads_x = std::min(threads_x, THREADS_PER_BLOCK); |
| 462 | uint threads_y = THREADS_PER_BLOCK / threads_x; |
| 463 | |
| 464 | uint blocks_x = divup(in.dims[0], threads_x * REPEAT); |
| 465 | uint blocks_y = divup(in.dims[1], threads_y); |
| 466 | |
| 467 | Array<T> tmpOut = |
| 468 | createEmptyArray<T>({blocks_x, in.dims[1], in.dims[2], in.dims[3]}); |
| 469 | Array<Tw> tmpWt = createEmptyArray<Tw>( |
| 470 | {blocks_x, in.dims[1], in.dims[2], in.dims[3]}); |
| 471 | |
| 472 | int tmp_elements = tmpOut.elements(); |
| 473 | |
| 474 | mean_first_launcher<T, Tw, T>(tmpOut, tmpWt, in, iwt, blocks_x, |
| 475 | blocks_y, threads_x); |
| 476 | |
| 477 | std::vector<T> h_ptr(tmp_elements); |
| 478 | std::vector<Tw> h_wptr(tmp_elements); |
| 479 | |
| 480 | CUDA_CHECK(cudaMemcpyAsync( |
| 481 | h_ptr.data(), tmpOut.get(), tmp_elements * sizeof(T), |
| 482 | cudaMemcpyDeviceToHost, getStream(getActiveDeviceId()))); |
| 483 | CUDA_CHECK(cudaMemcpyAsync( |
| 484 | h_wptr.data(), tmpWt.get(), tmp_elements * sizeof(Tw), |
| 485 | cudaMemcpyDeviceToHost, getStream(getActiveDeviceId()))); |
| 486 | CUDA_CHECK(cudaStreamSynchronize(getStream(getActiveDeviceId()))); |
| 487 | |
| 488 | compute_t<T> val = static_cast<compute_t<T>>(h_ptr[0]); |
| 489 | compute_t<Tw> weight = static_cast<compute_t<Tw>>(h_wptr[0]); |
| 490 |
nothing calls this directly
no test coverage detected