| 259 | template <int BlockSize, int NumPerThread, typename Self, |
| 260 | typename Reducer, typename Index> |
| 261 | __global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self input, Index num_coeffs, |
| 262 | half* output, packet_traits<Eigen::half>::type* scratch) { |
| 263 | typedef typename packet_traits<Eigen::half>::type PacketType; |
| 264 | const int packet_width = unpacket_traits<PacketType>::size; |
| 265 | eigen_assert(NumPerThread % packet_width == 0); |
| 266 | const Index first_index = |
| 267 | blockIdx.x * BlockSize * NumPerThread + packet_width * threadIdx.x; |
| 268 | |
| 269 | // Initialize the output value if it wasn't initialized by the ReductionInitKernel |
| 270 | |
| 271 | if (gridDim.x == 1) { |
| 272 | if (first_index == 0) { |
| 273 | int rem = num_coeffs % packet_width; |
| 274 | if (rem != 0) { |
| 275 | half2* p_scratch = reinterpret_cast<half2*>(scratch); |
| 276 | *scratch = reducer.template initializePacket<PacketType>(); |
| 277 | for (int i = 0; i < rem / 2; i++) { |
| 278 | *p_scratch = __halves2half2( |
| 279 | input.m_impl.coeff(num_coeffs - packet_width + 2 * i), |
| 280 | input.m_impl.coeff(num_coeffs - packet_width + 2 * i + 1)); |
| 281 | p_scratch++; |
| 282 | } |
| 283 | if ((num_coeffs & 1) != 0) { |
| 284 | half last = input.m_impl.coeff(num_coeffs - 1); |
| 285 | *p_scratch = __halves2half2(last, reducer.initialize()); |
| 286 | } |
| 287 | } else { |
| 288 | *scratch = reducer.template initializePacket<PacketType>(); |
| 289 | } |
| 290 | } |
| 291 | __syncthreads(); |
| 292 | } |
| 293 | |
| 294 | PacketType accum = reducer.template initializePacket<PacketType>(); |
| 295 | const Index max_iter = |
| 296 | numext::mini<Index>((num_coeffs - first_index) / packet_width, |
| 297 | NumPerThread * BlockSize / packet_width); |
| 298 | for (Index i = 0; i < max_iter; i += BlockSize) { |
| 299 | const Index index = first_index + packet_width * i; |
| 300 | eigen_assert(index + packet_width < num_coeffs); |
| 301 | PacketType val = input.m_impl.template packet<Unaligned>(index); |
| 302 | reducer.reducePacket(val, &accum); |
| 303 | } |
| 304 | |
| 305 | #pragma unroll |
| 306 | for (int offset = warpSize/2; offset > 0; offset /= 2) { |
| 307 | #if defined(EIGEN_HIPCC) |
| 308 | PacketType r1; |
| 309 | half2* hr = reinterpret_cast<half2*>(&r1); |
| 310 | half2* hacc = reinterpret_cast<half2*>(&accum); |
| 311 | for (int i = 0; i < packet_width / 2; i++) { |
| 312 | // FIXME : remove this workaround once we have native half/half2 support for __shfl_down |
| 313 | union { int i; half2 h; } wka_in, wka_out; |
| 314 | wka_in.h = hacc[i]; |
| 315 | wka_out.i = __shfl_down(wka_in.i, offset, warpSize); |
| 316 | hr[i] = wka_out.h; |
| 317 | } |
| 318 | reducer.reducePacket(r1, &accum); |
nothing calls this directly
no test coverage detected