| 274 | template<const unsigned int BlockSize, class T> |
| 275 | __global__ __launch_bounds__(BlockSize) |
| 276 | void example_global_memory_storage(const T* input, T* output, uint8_t* global_storage_raw) |
| 277 | { |
| 278 | // Indexing for this block |
| 279 | unsigned int index = (blockIdx.x * BlockSize) + threadIdx.x; |
| 280 | // specialize block_scan for type T and block of 256 threads |
| 281 | using block_scan_type = rocprim::block_scan<T, BlockSize>; |
| 282 | |
| 283 | T input_value = input[index]; |
| 284 | T output_value; |
| 285 | |
| 286 | // Cast the raw pointer to the correct storage type |
| 287 | using storage_type = typename block_scan_type::storage_type; |
| 288 | storage_type* global_storage = reinterpret_cast<storage_type*>(global_storage_raw); |
| 289 | |
| 290 | block_scan_type().inclusive_scan(input_value, |
| 291 | output_value, |
| 292 | global_storage[blockIdx.x], |
| 293 | rocprim::plus<T>()); |
| 294 | |
| 295 | output[index] = output_value; |
| 296 | } |
| 297 | |
| 298 | // Host function that runs example_global_memory_storage kernel |
| 299 | template<class T> |
nothing calls this directly
no test coverage detected