| 198 | template<const unsigned int BlockSize, class T> |
| 199 | __global__ __launch_bounds__(BlockSize) |
| 200 | void example_dynamic_shared_memory(const T* input, T* output) |
| 201 | { |
| 202 | // Indexing for this block |
| 203 | unsigned int index = (blockIdx.x * BlockSize) + threadIdx.x; |
| 204 | |
| 205 | // Initialize primitives |
| 206 | using block_scan_type = rocprim::block_scan<T, BlockSize>; |
| 207 | |
| 208 | // Allocation done in runtime, for more information please visit: |
| 209 | // https://github.com/ROCm-Developer-Tools/HIP/tree/master/samples/2_Cookbook/6_dynamic_shared |
| 210 | HIP_DYNAMIC_SHARED(typename block_scan_type::storage_type, primitive_storage); |
| 211 | |
| 212 | // Variables required for performing a scan |
| 213 | T input_value, output_value; |
| 214 | |
| 215 | // execute inclusive scan |
| 216 | input_value = input[index]; |
| 217 | block_scan_type().inclusive_scan(input_value, |
| 218 | output_value, |
| 219 | *primitive_storage, |
| 220 | rocprim::plus<T>()); |
| 221 | |
| 222 | output[index] = output_value; |
| 223 | } |
| 224 | |
| 225 | // Host function that runs example_dynamic_shared_memory kernel |
| 226 | template<class T> |
nothing calls this directly
no test coverage detected