MCPcopy Create free account
hub / github.com/ROCm/rocPRIM / naive_kernel

Function naive_kernel

benchmark/benchmark_device_batch_memcpy.cpp:308–340  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

306template<typename OffsetType, int32_t BlockSize>
307__launch_bounds__(BlockSize) __global__
308void 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
342template<typename ValueType,
343 typename BufferSizeType,

Callers

nothing calls this directly

Calls 3

flat_block_idFunction · 0.85
flat_block_thread_idFunction · 0.85

Tested by

no test coverage detected