Add CMakeLists.txt to compile examples
With the CMakeLists.txt all examples (except weld_vertices) are compiled with BUILD_EXAMPLES is turned on. Also the linking example in the cpp_integration subdir is built.
Note: not all examples run. Some examples throw runtime errors. These will be fixed later.
MODIFY
CMakeLists.txt
CMakeLists.txt
@@ -36,6 +36,7 @@ include(cmake/VerifyCompiler.cmake)
# Disable -Werror
option(DISABLE_WERROR "Disable building with Werror" ON)
option(BUILD_TEST "Build tests" OFF)
+option(BUILD_EXAMPLES "Build examples" OFF)
option(DOWNLOAD_ROCPRIM "Download rocPRIM and do not search for rocPRIM package" OFF)
# Set CXX flags
@@ -71,6 +72,12 @@ if(BUILD_TEST)
add_subdirectory(testing)
endif()
+# Tests
+if(BUILD_EXAMPLES)
+ message (STATUS " BUILD_EXAMPLES : ON")
+ add_subdirectory(examples)
+endif()
+
# Package
set(CPACK_DEBIAN_ARCHIVE_TYPE "gnutar")

MODIFY
examples/CMakeLists.txt
examples/CMakeLists.txt
-# message(STATUS "Adding \"examples\"")
+# ########################################################################
+# Copyright 2020 Advanced Micro Devices, Inc.
+# ########################################################################
-#aux_source_directory("testing" sources)
-FILE(GLOB SOURCES_CU *.cu)
-FILE(GLOB SOURCES_CPP *.cpp)
-set(SOURCES ${SOURCES_CU})
+function(add_thrust_example EXAMPLE)
+ set(EXAMPLE_SOURCE "${EXAMPLE}.cu")
+ set(EXAMPLE_TARGET "example_thrust_${EXAMPLE}")
+ set_source_files_properties(${EXAMPLE_SOURCE} PROPERTIES LANGUAGE CXX COMPILE_FLAGS "-x c++")
+ add_executable(${EXAMPLE_TARGET} ${EXAMPLE_SOURCE})
-list(LENGTH SOURCES index)
-message(STATUS "Found ${index} examples")
-
-set(targets "")
-foreach (src ${SOURCES})
- get_filename_component(exec_name ${src} NAME_WE)
- set(target example-${exec_name})
- thrust_add_executable(${target} ${src})
- set_target_properties(${target} PROPERTIES OUTPUT_NAME ${exec_name})
- install(TARGETS ${target} DESTINATION "examples/${HOST_BACKEND}_host_${DEVICE_BACKEND}_device_${THRUST_MODE}" OPTIONAL COMPONENT examples-bin)
- list(APPEND targets ${target})
+ target_link_libraries(${EXAMPLE_TARGET}
+ PRIVATE
+ rocthrust
+ )
+ foreach(amdgpu_target ${AMDGPU_TARGETS})
+ target_link_libraries(${EXAMPLE_TARGET}
+ PRIVATE
+ --amdgpu-target=${amdgpu_target}
+ )
endforeach()
+ set_target_properties(${EXAMPLE_TARGET}
+ PROPERTIES
+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/examples/"
+ )
+endfunction()
+
-add_subdirectory(cuda)
-add_subdirectory(omp)
add_subdirectory(cpp_integration)
-add_custom_target(examples-bin DEPENDS ${targets})
-add_custom_target(install-examples-bin
- COMMAND
- "${CMAKE_COMMAND}" -DCMAKE_INSTALL_COMPONENT=examples-bin
- -P "${CMAKE_BINARY_DIR}/cmake_install.cmake"
-)
+# ****************************************************************************
+# Examples
+# ****************************************************************************
+message (STATUS "Building examples")
-install(FILES ${SOURCES} DESTINATION "examples" COMPONENT examples)
+add_thrust_example("arbitrary_transformation")
+add_thrust_example("basic_vector")
+add_thrust_example("bounding_box")
+add_thrust_example("bucket_sort2d")
+add_thrust_example("constant_iterator")
+add_thrust_example("counting_iterator")
+add_thrust_example("device_ptr")
+add_thrust_example("discrete_voronoi")
+add_thrust_example("dot_products_with_zip")
+add_thrust_example("expand")
+add_thrust_example("fill_copy_sequence")
+add_thrust_example("histogram")
+add_thrust_example("lambda")
+add_thrust_example("lexicographical_sort")
+add_thrust_example("max_abs_diff")
+add_thrust_example("minimal_custom_backend")
+add_thrust_example("minmax")
+add_thrust_example("mode")
+add_thrust_example("monte_carlo_disjoint_sequences")
+add_thrust_example("monte_carlo")
+add_thrust_example("norm")
+add_thrust_example("padded_grid_reduction")
+add_thrust_example("permutation_iterator")
+add_thrust_example("raw_reference_cast")
+add_thrust_example("remove_points2d")
+add_thrust_example("repeated_range")
+add_thrust_example("run_length_decoding")
+add_thrust_example("run_length_encoding")
+add_thrust_example("saxpy")
+add_thrust_example("scan_by_key")
+add_thrust_example("set_operations")
+add_thrust_example("simple_moving_average")
+add_thrust_example("sort")
+add_thrust_example("sorting_aos_vs_soa")
+add_thrust_example("sparse_vector")
+add_thrust_example("stream_compaction")
+add_thrust_example("strided_range")
+add_thrust_example("sum_rows")
+add_thrust_example("sum")
+add_thrust_example("summary_statistics")
+add_thrust_example("summed_area_table")
+add_thrust_example("tiled_range")
+add_thrust_example("transform_iterator")
+add_thrust_example("transform_output_iterator")
+add_thrust_example("uninitialized_vector")
+add_thrust_example("version")
+# add_thrust_example("weld_vertices")
+add_thrust_example("word_count")

MODIFY
examples/cpp_integration/CMakeLists.txt
examples/cpp_integration/CMakeLists.txt
-FILE(GLOB SOURCES_CU *.cu)
-FILE(GLOB SOURCES_CPP *.cpp)
-FILE(GLOB SOURCES_H *.h)
-set(SOURCES_BACKEND ${SOURCES_CU} ${SOURCES_CPP} ${SOURCES_H})
-list(APPEND SOURCES_BACKEND "README")
+# ########################################################################
+# Copyright 2020 Advanced Micro Devices, Inc.
+# ########################################################################
-install(FILES ${SOURCES_BACKEND} DESTINATION "examples/cpp_integration" COMPONENT examples)
-if (NOT "x${DEVICE_BACKEND}" STREQUAL "xCUDA")
- return()
-endif()
+# link host.cpp to device.cu
-list(LENGTH SOURCES_BACKEND index)
-message(STATUS "Found ${index} examples/cpp_integration")
+set(EXAMPLE_SOURCE host.cpp device.cu device.h)
+set(EXAMPLE_TARGET example_thrust_linking)
+set_source_files_properties(device.cu PROPERTIES LANGUAGE CXX COMPILE_FLAGS "-x c++")
+add_executable(${EXAMPLE_TARGET} ${EXAMPLE_SOURCE})
-set(targets_backend "")
-set(exec_name "cpp_integration")
-set(target example-${exec_name})
-thrust_add_executable(${target} ${SOURCES_BACKEND})
-set_target_properties(${target} PROPERTIES OUTPUT_NAME ${exec_name})
-install(TARGETS ${target} DESTINATION "examples/cpp_integration/${HOST_BACKEND}_host_${DEVICE_BACKEND}_device_${THRUST_MODE}" OPTIONAL COMPONENT examples-bin)
-list(APPEND targets_backend ${target})
-
-set(targets ${targets} ${targets_backend} PARENT_SCOPE)
\ No newline at end of file
+target_link_libraries(${EXAMPLE_TARGET}
+ PRIVATE
+ rocthrust
+)
+foreach(amdgpu_target ${AMDGPU_TARGETS})
+ target_link_libraries(${EXAMPLE_TARGET}
+ PRIVATE
+ --amdgpu-target=${amdgpu_target}
+ )
+endforeach()
+set_target_properties(${EXAMPLE_TARGET}
+ PROPERTIES
+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/examples/cpp_integration"
+)
\ No newline at end of file

DELETE
examples/cuda/CMakeLists.txt
-
-FILE(GLOB SOURCES_CU *.cu)
-FILE(GLOB SOURCES_CPP *.cpp)
-FILE(GLOB SOURCES_H *.h)
-set(SOURCES_BACKEND ${SOURCES_CU} ${SOURCES_CPP} ${SOURCES_H})
-
-install(FILES ${SOURCES_BACKEND} DESTINATION "examples/cuda" COMPONENT examples)
-
-if (NOT "x${DEVICE_BACKEND}" STREQUAL "xCUDA")
- return()
-endif()
-
-list(LENGTH SOURCES_BACKEND index)
-message(STATUS "Found ${index} examples/cuda")
-
-set(targets_backend "")
-foreach (src ${SOURCES_BACKEND})
- get_filename_component(exec_name ${src} NAME_WE)
- set(target example-${exec_name})
- thrust_add_executable(${target} ${src})
- set_target_properties(${target} PROPERTIES OUTPUT_NAME ${exec_name})
- install(TARGETS ${target} DESTINATION "examples/cuda/${HOST_BACKEND}_host_${DEVICE_BACKEND}_device_${THRUST_MODE}" OPTIONAL COMPONENT examples-bin)
- list(APPEND targets_backend ${target})
-endforeach()
-
-set(targets ${targets} ${targets_backend} PARENT_SCOPE)
-
-

DELETE
examples/cuda/async_reduce.cu
-#include <thrust/device_vector.h>
-#include <thrust/reduce.h>
-#include <thrust/system/cuda/execution_policy.h>
-#include <cassert>
-
-#if __cplusplus >= 201103L
-#include <future>
-#endif
-
-// This example demonstrates two ways to achieve algorithm invocations that are asynchronous with
-// the calling thread.
-//
-// The first method wraps a call to thrust::reduce inside a __global__ function. Since __global__ function
-// launches are asynchronous with the launching thread, this achieves asynchrony. The result of the reduction
-// is stored to a pointer to CUDA global memory. The calling thread waits for the result of the reduction to
-// be ready by synchronizing with the CUDA stream on which the __global__ function is launched.
-//
-// The second method uses the C++11 library function, std::async, to create concurrency. The lambda function
-// given to std::async returns the result of thrust::reduce to a std::future. The calling thread can use the
-// std::future to wait for the result of the reduction. This method requires a compiler which supports
-// C++11-capable language and library constructs.
-
-template<typename Iterator, typename T, typename BinaryOperation, typename Pointer>
-__global__ void reduce_kernel(Iterator first, Iterator last, T init, BinaryOperation binary_op, Pointer result)
-{
- *result = thrust::reduce(thrust::cuda::par, first, last, init, binary_op);
-}
-
-int main()
-{
- size_t n = 1 << 20;
- thrust::device_vector<unsigned int> data(n, 1);
- thrust::device_vector<unsigned int> result(1, 0);
-
- // method 1: call thrust::reduce from an asynchronous CUDA kernel launch
-
- // create a CUDA stream
- cudaStream_t s;
- cudaStreamCreate(&s);
-
- // launch a CUDA kernel with only 1 thread on our stream
- reduce_kernel<<<1,1,0,s>>>(data.begin(), data.end(), 0, thrust::plus<int>(), result.data());
-
- // wait for the stream to finish
- cudaStreamSynchronize(s);
-
- // our result should be ready
- assert(result[0] == n);
-
- cudaStreamDestroy(s);
-
- // reset the result
- result[0] = 0;
-
-#if __cplusplus >= 201103L
- // method 2: use std::async to create asynchrony
-
- // copy all the algorithm parameters
- auto begin = data.begin();
- auto end = data.end();
- unsigned int init = 0;
- auto binary_op = thrust::plus<unsigned int>();
-
- // std::async captures the algorithm parameters by value
- // use std::launch::async to ensure the creation of a new thread
- std::future<unsigned int> future_result = std::async(std::launch::async, [=]
- {
- return thrust::reduce(begin, end, init, binary_op);
- });
-
- // wait on the result and check that it is correct
- assert(future_result.get() == n);
-#endif
-
- return 0;
-}
-

DELETE
examples/cuda/custom_temporary_allocation.cu
-#include <thrust/system/cuda/vector.h>
-#include <thrust/system/cuda/execution_policy.h>
-#include <thrust/host_vector.h>
-#include <thrust/generate.h>
-#include <thrust/sort.h>
-#include <thrust/pair.h>
-#include <cstdlib>
-#include <iostream>
-#include <sstream>
-#include <map>
-#include <cassert>
-
-// This example demonstrates how to intercept calls to get_temporary_buffer
-// and return_temporary_buffer to control how Thrust allocates temporary storage
-// during algorithms such as thrust::sort. The idea will be to create a simple
-// cache of allocations to search when temporary storage is requested. If a hit
-// is found in the cache, we quickly return the cached allocation instead of
-// resorting to the more expensive thrust::cuda::malloc.
-//
-// Note: this implementation cached_allocator is not thread-safe. If multiple
-// (host) threads use the same cached_allocator then they should gain exclusive
-// access to the allocator before accessing its methods.
-
-struct not_my_pointer
-{
- not_my_pointer(void* p)
- : message()
- {
- std::stringstream s;
- s << "Pointer `" << p << "` was not allocated by this allocator.";
- message = s.str();
- }
-
- virtual ~not_my_pointer() {}
-
- virtual const char* what() const
- {
- return message.c_str();
- }
-
-private:
- std::string message;
-};
-
-// A simple allocator for caching cudaMalloc allocations.
-struct cached_allocator
-{
- typedef char value_type;
-
- cached_allocator() {}
-
- ~cached_allocator()
- {
- free_all();
- }
-
- char *allocate(std::ptrdiff_t num_bytes)
- {
- std::cout << "cached_allocator::allocate(): num_bytes == "
- << num_bytes
- << std::endl;
-
- char *result = 0;
-
- // Search the cache for a free block.
- free_blocks_type::iterator free_block = free_blocks.find(num_bytes);
-
- if (free_block != free_blocks.end())
- {
- std::cout << "cached_allocator::allocate(): found a free block"
- << std::endl;
-
- result = free_block->second;
-
- // Erase from the `free_blocks` map.
- free_blocks.erase(free_block);
- }
- else
- {
- // No allocation of the right size exists, so create a new one with
- // `thrust::cuda::malloc`.
- try
- {
- std::cout << "cached_allocator::allocate(): allocating new block"
- << std::endl;
-
- // Allocate memory and convert the resulting `thrust::cuda::pointer` to
- // a raw pointer.
- result = thrust::cuda::malloc<char>(num_bytes).get();
- }
- catch (std::runtime_error&)
- {
- throw;
- }
- }
-
- // Insert the allocated pointer into the `allocated_blocks` map.
- allocated_blocks.insert(std::make_pair(result, num_bytes));
-
- return result;
- }
-
- void deallocate(char *ptr, size_t)
- {
- std::cout << "cached_allocator::deallocate(): ptr == "
- << reinterpret_cast<void*>(ptr) << std::endl;
-
- // Erase the allocated block from the allocated blocks map.
- allocated_blocks_type::iterator iter = allocated_blocks.find(ptr);
-
- if (iter == allocated_blocks.end())
- throw not_my_pointer(reinterpret_cast<void*>(ptr));
-
- std::ptrdiff_t num_bytes = iter->second;
- allocated_blocks.erase(iter);
-
- // Insert the block into the free blocks map.
- free_blocks.insert(std::make_pair(num_bytes, ptr));
- }
-
-private:
- typedef std::multimap<std::ptrdiff_t, char*> free_blocks_type;
- typedef std::map<char*, std::ptrdiff_t> allocated_blocks_type;
-
- free_blocks_type free_blocks;
- allocated_blocks_type allocated_blocks;
-
- void free_all()
- {
- std::cout << "cached_allocator::free_all()" << std::endl;
-
- // Deallocate all outstanding blocks in both lists.
- for ( free_blocks_type::iterator i = free_blocks.begin()
- ; i != free_blocks.end()
- ; ++i)
- {
- // Transform the pointer to cuda::pointer before calling cuda::free.
- thrust::cuda::free(thrust::cuda::pointer<char>(i->second));
- }
-
- for( allocated_blocks_type::iterator i = allocated_blocks.begin()
- ; i != allocated_blocks.end()
- ; ++i)
- {
- // Transform the pointer to cuda::pointer before calling cuda::free.
- thrust::cuda::free(thrust::cuda::pointer<char>(i->first));
- }
- }
-};
-
-int main()
-{
- std::size_t num_elements = 32768;
-
- thrust::host_vector<int> h_input(num_elements);
-
- // Generate random input.
- thrust::generate(h_input.begin(), h_input.end(), rand);
-
- thrust::cuda::vector<int> d_input = h_input;
- thrust::cuda::vector<int> d_result(num_elements);
-
- std::size_t num_trials = 5;
-
- cached_allocator alloc;
-
- for (std::size_t i = 0; i < num_trials; ++i)
- {
- d_result = d_input;
-
- // Pass alloc through cuda::par as the first parameter to sort
- // to cause allocations to be handled by alloc during sort.
- thrust::sort(thrust::cuda::par(alloc), d_result.begin(), d_result.end());
-
- // Ensure the result is sorted.
- assert(thrust::is_sorted(d_result.begin(), d_result.end()));
- }
-
- return 0;
-}
-

DELETE
examples/cuda/simple_cuda_streams.cu
-#include <thrust/device_vector.h>
-#include <thrust/for_each.h>
-#include <thrust/system/cuda/execution_policy.h>
-#include <cstdio>
-#include <cassert>
-
-// This example demonstrates how to achieve asynchronous, concurrent algorithm execution using
-// the CUDA backend's low-level stream-based interface. This program uses thrust::for_each to invoke
-// two functors, "ping", and "pong", which communicate via a shared variable, "ball". To encourage
-// concurrency, we execute thrust::for_each on two independent CUDA streams using the thrust::cuda::par
-// execution policy.
-//
-// Note that stream usage provides no guarantee of concurrency. If the ping and pong functions
-// do not happen to be scheduled concurrently, this program will deadlock.
-
-struct ping
-{
- // XXX nvcc issue prevents us from making ball volatile
- //__device__
- //void operator()(volatile int &ball)
- __device__
- void operator()(int &ball)
- {
- // we're not guaranteed concurrency, so only attempt this 1000 times
- unsigned int attempt = 0;
-
- ball = 1;
-
- for(unsigned int next_state = 2;
- next_state < 25 && attempt < 1000;
- next_state += 2)
- {
- while(ball != next_state && attempt < 1000)
- {
-#if __CUDA_ARCH__ >= 200
- printf("ping waiting for return\n");
-#endif
- ++attempt;
- }
-
- ball += 1;
-
-#if __CUDA_ARCH__ >= 200
- printf("ping! ball is now %d\n", next_state + 1);
-#endif
- }
- }
-};
-
-struct pong
-{
- // XXX nvcc issue prevents us from making ball volatile
- //__device__
- //void operator()(volatile int &ball)
- __device__
- void operator()(int &ball)
- {
- // we're not guaranteed concurrency, so only attempt this 1000 times
- unsigned int attempt = 0;
-
- for(unsigned int next_state = 1;
- next_state < 25 && attempt < 1000;
- next_state += 2)
- {
- while(ball != next_state && attempt < 1000)
- {
-#if __CUDA_ARCH__ >= 200
- printf("pong waiting for return\n");
-#endif
- ++attempt;
- }
-
- ball += 1;
-
-#if __CUDA_ARCH__ >= 200
- printf("pong! ball is now %d\n", next_state + 1);
-#endif
- }
- }
-};
-
-int main()
-{
- cudaStream_t s1, s2;
- cudaStreamCreate(&s1);
- cudaStreamCreate(&s2);
-
- thrust::device_vector<int> ball(1);
-
- // Invoke thrust::for_each with the thrust::cuda::par
- // execution policy. Pass the stream s1 as an argument
- // to the .on() function
- thrust::for_each(thrust::cuda::par.on(s1),
- ball.begin(),
- ball.end(),
- ping());
-
- // Invoke thrust::for_each with the thrust::cuda::par
- // execution policy. Pass the stream s2 as an argument
- // to the .on() function
- thrust::for_each(thrust::cuda::par.on(s2),
- ball.begin(),
- ball.end(),
- pong());
-
- // Wait for all algorithms executed on the streams to terminate.
- cudaStreamSynchronize(s1);
- cudaStreamSynchronize(s2);
-
- cudaStreamDestroy(s1);
- cudaStreamDestroy(s2);
-
- return 0;
-}
-
-

DELETE
examples/cuda/unwrap_pointer.cu
-#include <thrust/device_ptr.h>
-#include <thrust/device_malloc.h>
-#include <thrust/device_free.h>
-#include <thrust/device_vector.h>
-#include <cuda.h>
-
-int main(void)
-{
- size_t N = 10;
-
- // create a device_ptr
- thrust::device_ptr<int> dev_ptr = thrust::device_malloc<int>(N);
-
- // extract raw pointer from device_ptr
- int * raw_ptr = thrust::raw_pointer_cast(dev_ptr);
-
- // use raw_ptr in CUDA API functions
- cudaMemset(raw_ptr, 0, N * sizeof(int));
-
- // free memory
- thrust::device_free(dev_ptr);
-
- // we can use the same approach for device_vector
- thrust::device_vector<int> d_vec(N);
-
- // note: d_vec.data() returns a device_ptr
- raw_ptr = thrust::raw_pointer_cast(d_vec.data());
-
- return 0;
-}

DELETE
examples/cuda/wrap_pointer.cu
-#include <thrust/device_ptr.h>
-#include <thrust/fill.h>
-#include <cuda.h>
-
-int main(void)
-{
- size_t N = 10;
-
- // obtain raw pointer to device memory
- int * raw_ptr;
- cudaMalloc((void **) &raw_ptr, N * sizeof(int));
-
- // wrap raw pointer with a device_ptr
- thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(raw_ptr);
-
- // use device_ptr in Thrust algorithms
- thrust::fill(dev_ptr, dev_ptr + N, (int) 0);
-
- // access device memory transparently through device_ptr
- dev_ptr[0] = 1;
-
- // free memory
- cudaFree(raw_ptr);
-
- return 0;
-}

DELETE
examples/omp/CMakeLists.txt
-FILE(GLOB SOURCES_CU *.cu)
-FILE(GLOB SOURCES_CPP *.cpp)
-set(SOURCES_BACKEND ${SOURCES_CU})
-
-install(FILES ${SOURCES_BACKEND} DESTINATION "examples/omp" COMPONENT examples)
-
-if (NOT "x${DEVICE_BACKEND}" STREQUAL "xOMP")
- return()
-endif()

Add CMakeLists.txt to compile examples

With the CMakeLists.txt all examples (except weld_vertices) are compiled with BUILD_EXAMPLES is turned on. Also the linking example in the cpp_integration subdir is built.

Note: not all examples run. Some examples throw runtime errors. These will be fixed later.