| 306 | template<typename OffsetType, int32_t BlockSize> |
| 307 | __launch_bounds__(BlockSize) __global__ |
| 308 | void naive_kernel(void** in_ptr, void** out_ptr, const OffsetType* sizes) |
| 309 | { |
| 310 | using underlying_type = unsigned char; |
| 311 | constexpr int32_t items_per_thread = 4; |
| 312 | constexpr int32_t tile_size = items_per_thread * BlockSize; |
| 313 | |
| 314 | const int32_t buffer_id = rocprim::flat_block_id(); |
| 315 | auto in = reinterpret_cast<underlying_type*>(in_ptr[buffer_id]); |
| 316 | auto out = reinterpret_cast<underlying_type*>(out_ptr[buffer_id]); |
| 317 | |
| 318 | const auto size = sizes[buffer_id]; |
| 319 | const auto size_in_elements = size / sizeof(underlying_type); |
| 320 | const auto tiles = size_in_elements / tile_size; |
| 321 | |
| 322 | auto num_items_to_copy = size; |
| 323 | |
| 324 | for(size_t i = 0; i < tiles; ++i) |
| 325 | { |
| 326 | underlying_type data[items_per_thread]; |
| 327 | rocprim::block_load_direct_blocked(rocprim::flat_block_thread_id(), |
| 328 | in, |
| 329 | data, |
| 330 | num_items_to_copy); |
| 331 | rocprim::block_store_direct_blocked(rocprim::flat_block_thread_id(), |
| 332 | out, |
| 333 | data, |
| 334 | num_items_to_copy); |
| 335 | |
| 336 | in += tile_size; |
| 337 | out += tile_size; |
| 338 | num_items_to_copy -= tile_size; |
| 339 | } |
| 340 | } |
| 341 | |
| 342 | template<typename ValueType, |
| 343 | typename BufferSizeType, |
nothing calls this directly
no test coverage detected