#ifndef max_scan
#define max_scan
template <typename T>
static __device__ inline T block_max_scan(T val, void* buffer) {
const int lane = threadIdx.x % WS;
const int warp = threadIdx.x / WS;
const int warps = TPB / WS;
T* const carry = (T*)buffer;
assert(WS >= warps);
T tmp = __shfl_up(val, 1);
if (lane >= 1) val = max(val, tmp);
tmp = __shfl_up(val, 2);
if (lane >= 2) val = max(val, tmp);
tmp = __shfl_up(val, 4);
if (lane >= 4) val = max(val, tmp);
tmp = __shfl_up(val, 8);
if (lane >= 8) val = max(val, tmp);
tmp = __shfl_up(val, 16);
if (lane >= 16) val = max(val, tmp);
#if defined(WS) && (WS == 64)
tmp = __shfl_up(val, 32);
if (lane >= 32) val = max(val, tmp);
#endif
if (lane == WS - 1) carry[warp] = val;
__syncthreads();
if constexpr (warps > 1) {
if (warp == 0) {
T res = carry[lane];
T tmp = __shfl_up(res, 1);
if (lane >= 1) res = max(res, tmp);
if constexpr (warps > 2) {
tmp = __shfl_up(res, 2);
if (lane >= 2) res = max(res, tmp);
if constexpr (warps > 4) {
tmp = __shfl_up(res, 4);
if (lane >= 4) res = max(res, tmp);
if constexpr (warps > 8) {
tmp = __shfl_up(res, 8);
if (lane >= 8) res = max(res, tmp);
if constexpr (warps > 16) {
tmp = __shfl_up(res, 16);
if (lane >= 16) res = max(res, tmp);
#if defined(WS) && (WS == 64)
if constexpr (warps > 32) {
tmp = __shfl_up(res, 32);
if (lane >= 32) res = max(res, tmp);
}
#endif
}
}
}
}
carry[lane] = res;
}
__syncthreads();
if (warp > 0) val = max(val, carry[warp - 1]);
}
return val;
}
#endif