Skip to content
Merged
Show file tree
Hide file tree
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion Applications/monte_carlo_pi/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -39,10 +39,12 @@ include("${CMAKE_CURRENT_LIST_DIR}/../../Common/ROCmPath.cmake")
find_package(hipcub REQUIRED)
find_package(hiprand REQUIRED)
# Workaround for hipRAND, requires manual linking with backend.
# rocThrust provides the iterators on AMD; NVIDIA gets Thrust from CUDAToolkit.
if(ROCM_EXAMPLES_HIP_PLATFORM STREQUAL "nvidia")
find_package(CUDAToolkit REQUIRED)
else()
find_package(rocrand REQUIRED)
find_package(rocthrust REQUIRED)
endif()

add_executable(${example_name} main.hip)
Expand All @@ -53,7 +55,7 @@ target_link_libraries(${example_name} PRIVATE hip::hipcub hip::hiprand)
if(ROCM_EXAMPLES_HIP_PLATFORM STREQUAL "nvidia")
target_link_libraries(${example_name} PRIVATE CUDA::curand)
else()
target_link_libraries(${example_name} PRIVATE roc::rocrand)
target_link_libraries(${example_name} PRIVATE roc::rocrand roc::rocthrust)
endif()
target_include_directories(
${example_name}
Expand Down
13 changes: 6 additions & 7 deletions Applications/monte_carlo_pi/main.hip
Original file line number Diff line number Diff line change
Expand Up @@ -25,11 +25,12 @@
#include "hiprand_utils.hpp"

#include <hipcub/device/device_reduce.hpp>
#include <hipcub/iterator/counting_input_iterator.hpp>
#include <hipcub/iterator/discard_output_iterator.hpp>
#include <hipcub/iterator/transform_input_iterator.hpp>
#include <hiprand/hiprand.h>

// CUB 3.0 removed hipcub's CUB-style iterators; use portable Thrust ones.
Comment thread
lamb-j marked this conversation as resolved.
Outdated
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>

#include <hip/hip_runtime.h>

#include <cstddef>
Expand Down Expand Up @@ -66,13 +67,11 @@ float calculate_pi(int sample_count, float* d_data)
// 4. Set up the input and output iterator for hipCUB's Sum.

// Represents the samples' index.
auto input_counting = hipcub::CountingInputIterator<int>(0);
auto input_counting = thrust::counting_iterator<int>(0);

// Converts the sample's index to a 0 or 1, indicating whether the sample lies within the disk.
conversion_op convert_op(sample_count, d_data);
auto input = hipcub::TransformInputIterator<bool, conversion_op, decltype(input_counting)>(
input_counting,
convert_op);
auto input = thrust::make_transform_iterator(input_counting, convert_op);

int* d_output{};
HIP_CHECK(hipMalloc(&d_output, sizeof(int)));
Expand Down
Loading