| 218 | typename std::enable_if<MemOp == block_primitive_direct, int>::type = 0> |
| 219 | __global__ __launch_bounds__(BlockSize) |
| 220 | void operation_kernel(T* input, T* output, CustomOp op) |
| 221 | { |
| 222 | constexpr unsigned int items_per_block = BlockSize * ItemsPerThread; |
| 223 | |
| 224 | using block_load_type = typename rocprim:: |
| 225 | block_load<T, BlockSize, ItemsPerThread, rocprim::block_load_method::block_load_direct>; |
| 226 | using block_store_type = typename rocprim:: |
| 227 | block_store<T, BlockSize, ItemsPerThread, rocprim::block_store_method::block_store_direct>; |
| 228 | |
| 229 | block_load_type load; |
| 230 | block_store_type store; |
| 231 | |
| 232 | __shared__ union |
| 233 | { |
| 234 | typename block_load_type::storage_type load; |
| 235 | typename block_store_type::storage_type store; |
| 236 | } storage; |
| 237 | |
| 238 | int offset = blockIdx.x * items_per_block; |
| 239 | |
| 240 | T items[ItemsPerThread]; |
| 241 | load.load(input + offset, items, storage.load); |
| 242 | __syncthreads(); |
| 243 | op(items, &storage, sizeof(storage), output); |
| 244 | store.store(output + offset, items, storage.store); |
| 245 | } |
| 246 | |
| 247 | // vectorized method base kernel |
| 248 | template<typename T, |