| 290 | |
| 291 | template<class Type, int32_t Length> |
| 292 | __global__ |
| 293 | void thread_reduce_kernel(Type* const device_input, Type* device_output) |
| 294 | { |
| 295 | size_t index = (blockIdx.x * blockDim.x + threadIdx.x) * Length; |
| 296 | |
| 297 | device_output[index] = rocprim::thread_reduce<Length>(&device_input[index], sum_op()); |
| 298 | } |
| 299 | |
| 300 | template<class Type, int32_t Length> |
| 301 | __global__ |