| 315 | { |
| 316 | template<typename _Input, typename _Output, typename _Op, typename _Scalar> |
| 317 | __global__ void ReduceThreadKernel(dim3 virtual_size, |
| 318 | _Input input, _Output output, _Op op, int N) |
| 319 | { |
| 320 | CUMAT_KERNEL_1D_LOOP(i, virtual_size) |
| 321 | const Index O = i * N; |
| 322 | _Scalar v = input[O]; |
| 323 | for (int n = 1; n < N; ++n) v = op(v, input[O + n]); |
| 324 | output[i] = v; |
| 325 | CUMAT_KERNEL_1D_LOOP_END |
| 326 | } |
| 327 | } |
| 328 | template<typename _Input, typename _Output, int _Axis, typename _Op, typename _Scalar> |
| 329 | struct ReductionEvaluator<_Input, _Output, _Axis, _Op, _Scalar, ReductionAlg::Thread> |