| 299 | |
| 300 | template<class Type, int32_t Length> |
| 301 | __global__ |
| 302 | void thread_reduce_kernel_array_prefix(Type* const device_input, Type* device_output, Type prefix) |
| 303 | { |
| 304 | size_t index = (blockIdx.x * blockDim.x + threadIdx.x) * Length; |
| 305 | |
| 306 | Type device_input_array[Length]; |
| 307 | for(int32_t i = 0; i < Length; i++) |
| 308 | { |
| 309 | device_input_array[i] = device_input[index + i]; |
| 310 | } |
| 311 | |
| 312 | device_output[index] = rocprim::thread_reduce(device_input_array, sum_op(), prefix); |
| 313 | } |
| 314 | |
| 315 | TYPED_TEST(RocprimThreadOperationTests, Reduction) |
| 316 | { |