| 670 | extern void reverse_and_sort_with_thrust(std::uint32_t *device_pointer, std::size_t array_length); |
| 671 | |
| 672 | static void sorting_with_thrust(benchmark::State &state) { |
| 673 | const auto count = static_cast<std::size_t>(state.range(0)); |
| 674 | |
| 675 | // Typically, the data is first allocated on the "host" CPU side, |
| 676 | // initialized, and then transferred to the "device" GPU memory. |
| 677 | // In our specific case, we could have also used `thrust::sequence`. |
| 678 | thrust::host_vector<std::uint32_t> host_array(count); |
| 679 | std::iota(host_array.begin(), host_array.end(), 1u); |
| 680 | thrust::device_vector<std::uint32_t> device_array = host_array; |
| 681 | |
| 682 | for (auto _ : state) { |
| 683 | reverse_and_sort_with_thrust(device_array.data().get(), count); |
| 684 | cudaError_t error = cudaDeviceSynchronize(); //! Block until the GPU has completed all tasks |
| 685 | if (error != cudaSuccess) state.SkipWithError("CUDA error after kernel launch: "s + cudaGetErrorString(error)); |
| 686 | benchmark::DoNotOptimize(device_array.data()); |
| 687 | } |
| 688 | |
| 689 | state.SetComplexityN(count); |
| 690 | state.SetItemsProcessed(count * state.iterations()); |
| 691 | state.SetBytesProcessed(count * state.iterations() * sizeof(std::uint32_t)); |
| 692 | } |
| 693 | |
| 694 | BENCHMARK(sorting_with_thrust) |
| 695 | ->RangeMultiplier(4) |