| /****************************************************************************** |
| * Copyright (c) 2011, Duane Merrill. All rights reserved. |
| * Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved. |
| * |
| * Redistribution and use in source and binary forms, with or without |
| * modification, are permitted provided that the following conditions are met: |
| * * Redistributions of source code must retain the above copyright |
| * notice, this list of conditions and the following disclaimer. |
| * * Redistributions in binary form must reproduce the above copyright |
| * notice, this list of conditions and the following disclaimer in the |
| * documentation and/or other materials provided with the distribution. |
| * * Neither the name of the NVIDIA CORPORATION nor the |
| * names of its contributors may be used to endorse or promote products |
| * derived from this software without specific prior written permission. |
| * |
| * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND |
| * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED |
| * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE |
| * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY |
| * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES |
| * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; |
| * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND |
| * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT |
| * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS |
| * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
| * |
| ******************************************************************************/ |
| |
| /****************************************************************************** |
| * Test of DeviceReduce::RunLengthEncode utilities |
| ******************************************************************************/ |
| |
| // Ensure printing of CUDA runtime errors to console |
| #define CUB_STDERR |
| |
| #include <stdio.h> |
| #include <typeinfo> |
| |
| #include <thrust/device_ptr.h> |
| #include <thrust/reduce.h> |
| #include <thrust/iterator/constant_iterator.h> |
| |
| #include <cub/util_allocator.cuh> |
| #include <cub/iterator/constant_input_iterator.cuh> |
| #include <cub/device/device_reduce.cuh> |
| #include <cub/device/device_run_length_encode.cuh> |
| #include <cub/thread/thread_operators.cuh> |
| |
| #include "test_util.h" |
| |
| using namespace cub; |
| |
| |
| //--------------------------------------------------------------------- |
| // Globals, constants and typedefs |
| //--------------------------------------------------------------------- |
| |
| bool g_verbose = false; |
| int g_timing_iterations = 0; |
| int g_repeat = 0; |
| CachingDeviceAllocator g_allocator(true); |
| |
| // Dispatch types |
| enum Backend |
| { |
| CUB, // CUB method |
| THRUST, // Thrust method |
| CDP, // GPU-based (dynamic parallelism) dispatch to CUB method |
| }; |
| |
| // Operation types |
| enum RleMethod |
| { |
| RLE, // Run length encode |
| NON_TRIVIAL, |
| CSR, |
| }; |
| |
| |
| //--------------------------------------------------------------------- |
| // Dispatch to different CUB entrypoints |
| //--------------------------------------------------------------------- |
| |
| |
| /** |
| * Dispatch to run-length encode entrypoint |
| */ |
| template < |
| typename InputIteratorT, |
| typename UniqueOutputIteratorT, |
| typename OffsetsOutputIteratorT, |
| typename LengthsOutputIteratorT, |
| typename NumRunsIterator, |
| typename OffsetT> |
| CUB_RUNTIME_FUNCTION __forceinline__ |
| cudaError_t Dispatch( |
| Int2Type<RLE> /*method*/, |
| Int2Type<CUB> /*dispatch_to*/, |
| int timing_timing_iterations, |
| size_t */*d_temp_storage_bytes*/, |
| cudaError_t */*d_cdp_error*/, |
| |
| void* d_temp_storage, |
| size_t &temp_storage_bytes, |
| InputIteratorT d_in, |
| UniqueOutputIteratorT d_unique_out, |
| OffsetsOutputIteratorT /*d_offsets_out*/, |
| LengthsOutputIteratorT d_lengths_out, |
| NumRunsIterator d_num_runs, |
| cub::Equality /*equality_op*/, |
| OffsetT num_items, |
| cudaStream_t stream, |
| bool debug_synchronous) |
| { |
| cudaError_t error = cudaSuccess; |
| for (int i = 0; i < timing_timing_iterations; ++i) |
| { |
| error = DeviceRunLengthEncode::Encode( |
| d_temp_storage, |
| temp_storage_bytes, |
| d_in, |
| d_unique_out, |
| d_lengths_out, |
| d_num_runs, |
| num_items, |
| stream, |
| debug_synchronous); |
| } |
| return error; |
| } |
| |
|
|
| /** |
| * Dispatch to non-trivial runs entrypoint |
| */ |
| template < |
| typename InputIteratorT, |
| typename UniqueOutputIteratorT, |
| typename OffsetsOutputIteratorT, |
| typename LengthsOutputIteratorT, |
| typename NumRunsIterator, |
| typename OffsetT> |
| CUB_RUNTIME_FUNCTION __forceinline__ |
| cudaError_t Dispatch( |
| Int2Type<NON_TRIVIAL> /*method*/, |
| Int2Type<CUB> /*dispatch_to*/, |
| int timing_timing_iterations, |
| size_t */*d_temp_storage_bytes*/, |
| cudaError_t */*d_cdp_error*/, |
| |
| void* d_temp_storage, |
| size_t &temp_storage_bytes, |
| InputIteratorT d_in, |
| UniqueOutputIteratorT /*d_unique_out*/, |
| OffsetsOutputIteratorT d_offsets_out, |
| LengthsOutputIteratorT d_lengths_out, |
| NumRunsIterator d_num_runs, |
| cub::Equality /*equality_op*/, |
| OffsetT num_items, |
| cudaStream_t stream, |
| bool debug_synchronous) |
| { |
| cudaError_t error = cudaSuccess; |
| for (int i = 0; i < timing_timing_iterations; ++i) |
| { |
| error = DeviceRunLengthEncode::NonTrivialRuns( |
| d_temp_storage, |
| temp_storage_bytes, |
| d_in, |
| d_offsets_out, |
| d_lengths_out, |
| d_num_runs, |
| num_items, |
| stream, |
| debug_synchronous); |
| } |
| return error; |
| } |
| |
|
|
|
|
| //--------------------------------------------------------------------- |
| // Dispatch to different Thrust entrypoints |
| //--------------------------------------------------------------------- |
|
|
| /** |
| * Dispatch to run-length encode entrypoint |
| */ |
| template < |
| typename InputIteratorT, |
| typename UniqueOutputIteratorT, |
| typename OffsetsOutputIteratorT, |
| typename LengthsOutputIteratorT, |
| typename NumRunsIterator, |
| typename OffsetT> |
| cudaError_t Dispatch( |
| Int2Type<RLE> /*method*/, |
| Int2Type<THRUST> /*dispatch_to*/, |
| int timing_timing_iterations, |
| size_t */*d_temp_storage_bytes*/, |
| cudaError_t */*d_cdp_error*/, |
| |
| void *d_temp_storage, |
| size_t &temp_storage_bytes, |
| InputIteratorT d_in, |
| UniqueOutputIteratorT d_unique_out, |
| OffsetsOutputIteratorT /*d_offsets_out*/, |
| LengthsOutputIteratorT d_lengths_out, |
| NumRunsIterator d_num_runs, |
| cub::Equality /*equality_op*/, |
| OffsetT num_items, |
| cudaStream_t /*stream*/, |
| bool /*debug_synchronous*/) |
| { |
| // The input value type |
| typedef typename std::iterator_traits<InputIteratorT>::value_type InputT; |
| |
| // The output value type |
| typedef typename If<(Equals<typename std::iterator_traits<UniqueOutputIteratorT>::value_type, void>::VALUE), // OutputT = (if output iterator's value type is void) ? |
| typename std::iterator_traits<InputIteratorT>::value_type, // ... then the input iterator's value type, |
| typename std::iterator_traits<UniqueOutputIteratorT>::value_type>::Type UniqueT; // ... else the output iterator's value type |
| |
| // The lengths output value type |
| typedef typename If<(Equals<typename std::iterator_traits<LengthsOutputIteratorT>::value_type, void>::VALUE), // LengthT = (if output iterator's value type is void) ? |
| OffsetT, // ... then the OffsetT type, |
| typename std::iterator_traits<LengthsOutputIteratorT>::value_type>::Type LengthT; // ... else the output iterator's value type |
| |
| if (d_temp_storage == 0) |
| { |
| temp_storage_bytes = 1; |
| } |
| else |
| { |
| thrust::device_ptr<InputT> d_in_wrapper(d_in); |
| thrust::device_ptr<UniqueT> d_unique_out_wrapper(d_unique_out); |
| thrust::device_ptr<LengthT> d_lengths_out_wrapper(d_lengths_out); |
| |
| thrust::pair<thrust::device_ptr<UniqueT>, thrust::device_ptr<LengthT> > d_out_ends; |
| |
| LengthT one_val; |
| InitValue(INTEGER_SEED, one_val, 1); |
| thrust::constant_iterator<LengthT> constant_one(one_val); |
| |
| for (int i = 0; i < timing_timing_iterations; ++i) |
| { |
| d_out_ends = thrust::reduce_by_key( |
| d_in_wrapper, |
| d_in_wrapper + num_items, |
| constant_one, |
| d_unique_out_wrapper, |
| d_lengths_out_wrapper); |
| } |
| |
| OffsetT num_runs = OffsetT(d_out_ends.first - d_unique_out_wrapper); |
| CubDebugExit(cudaMemcpy(d_num_runs, &num_runs, sizeof(OffsetT), cudaMemcpyHostToDevice)); |
| } |
| |
| return cudaSuccess; |
| } |
| |
|
|
|
|
| //--------------------------------------------------------------------- |
| // CUDA Nested Parallelism Test Kernel |
| //--------------------------------------------------------------------- |
|
|
| /** |
| * Simple wrapper kernel to invoke DeviceRunLengthEncode |
| */ |
| template < |
| int RLE_METHOD, |
| typename InputIteratorT, |
| typename UniqueOutputIteratorT, |
| typename OffsetsOutputIteratorT, |
| typename LengthsOutputIteratorT, |
| typename NumRunsIterator, |
| typename EqualityOp, |
| typename OffsetT> |
| __global__ void CnpDispatchKernel( |
| Int2Type<RLE_METHOD> method, |
| int timing_timing_iterations, |
| size_t *d_temp_storage_bytes, |
| cudaError_t *d_cdp_error, |
| |
| void* d_temp_storage, |
| size_t temp_storage_bytes, |
| InputIteratorT d_in, |
| UniqueOutputIteratorT d_unique_out, |
| OffsetsOutputIteratorT d_offsets_out, |
| LengthsOutputIteratorT d_lengths_out, |
| NumRunsIterator d_num_runs, |
| cub::Equality equality_op, |
| OffsetT num_items, |
| cudaStream_t stream, |
| bool debug_synchronous) |
| { |
| |
| #ifndef CUB_CDP |
| *d_cdp_error = cudaErrorNotSupported; |
| #else |
| *d_cdp_error = Dispatch(method, Int2Type<CUB>(), timing_timing_iterations, d_temp_storage_bytes, d_cdp_error, |
| d_temp_storage, temp_storage_bytes, d_in, d_unique_out, d_offsets_out, d_lengths_out, d_num_runs, equality_op, num_items, 0, debug_synchronous); |
| |
| *d_temp_storage_bytes = temp_storage_bytes; |
| #endif |
| } |
| |
|
|
| /** |
| * Dispatch to CDP kernel |
| */ |
| template < |
| int RLE_METHOD, |
| typename InputIteratorT, |
| typename UniqueOutputIteratorT, |
| typename OffsetsOutputIteratorT, |
| typename LengthsOutputIteratorT, |
| typename NumRunsIterator, |
| typename EqualityOp, |
| typename OffsetT> |
| CUB_RUNTIME_FUNCTION __forceinline__ |
| cudaError_t Dispatch( |
| Int2Type<RLE_METHOD> method, |
| Int2Type<CDP> dispatch_to, |
| int timing_timing_iterations, |
| size_t *d_temp_storage_bytes, |
| cudaError_t *d_cdp_error, |
| |
| void* d_temp_storage, |
| size_t &temp_storage_bytes, |
| InputIteratorT d_in, |
| UniqueOutputIteratorT d_unique_out, |
| OffsetsOutputIteratorT d_offsets_out, |
| LengthsOutputIteratorT d_lengths_out, |
| NumRunsIterator d_num_runs, |
| EqualityOp equality_op, |
| OffsetT num_items, |
| cudaStream_t stream, |
| bool debug_synchronous) |
| { |
| // Invoke kernel to invoke device-side dispatch |
| CnpDispatchKernel<<<1,1>>>(method, timing_timing_iterations, d_temp_storage_bytes, d_cdp_error, |
| d_temp_storage, temp_storage_bytes, d_in, d_unique_out, d_offsets_out, d_lengths_out, d_num_runs, equality_op, num_items, 0, debug_synchronous); |
| |
| // Copy out temp_storage_bytes |
| CubDebugExit(cudaMemcpy(&temp_storage_bytes, d_temp_storage_bytes, sizeof(size_t) * 1, cudaMemcpyDeviceToHost)); |
| |
| // Copy out error |
| cudaError_t retval; |
| CubDebugExit(cudaMemcpy(&retval, d_cdp_error, sizeof(cudaError_t) * 1, cudaMemcpyDeviceToHost)); |
| return retval; |
| } |
| |
|
|
|
|
| //--------------------------------------------------------------------- |
| // Test generation |
| //--------------------------------------------------------------------- |
|
|
|
|
| /** |
| * Initialize problem |
| */ |
| template <typename T> |
| void Initialize( |
| int entropy_reduction, |
| T *h_in, |
| int num_items, |
| int max_segment) |
| { |
| unsigned int max_int = (unsigned int) -1; |
| |
| int key = 0; |
| int i = 0; |
| while (i < num_items) |
| { |
| // Select number of repeating occurrences for the current run |
| int repeat; |
| if (max_segment < 0) |
| { |
| repeat = num_items; |
| } |
| else if (max_segment < 2) |
| { |
| repeat = 1; |
| } |
| else |
| { |
| RandomBits(repeat, entropy_reduction); |
| repeat = (int) ((double(repeat) * double(max_segment)) / double(max_int)); |
| repeat = CUB_MAX(1, repeat); |
| } |
| |
| int j = i; |
| while (j < CUB_MIN(i + repeat, num_items)) |
| { |
| InitValue(INTEGER_SEED, h_in[j], key); |
| j++; |
| } |
| |
| i = j; |
| key++; |
| } |
| |
| if (g_verbose) |
| { |
| printf("Input:\n"); |
| DisplayResults(h_in, num_items); |
| printf("\n\n"); |
| } |
| } |
| |
|
|
| /** |
| * Solve problem. Returns total number of segments identified |
| */ |
| template < |
| RleMethod RLE_METHOD, |
| typename InputIteratorT, |
| typename T, |
| typename OffsetT, |
| typename LengthT, |
| typename EqualityOp> |
| int Solve( |
| InputIteratorT h_in, |
| T *h_unique_reference, |
| OffsetT *h_offsets_reference, |
| LengthT *h_lengths_reference, |
| EqualityOp equality_op, |
| int num_items) |
| { |
| if (num_items == 0) |
| return 0; |
| |
| // First item |
| T previous = h_in[0]; |
| LengthT length = 1; |
| int num_runs = 0; |
| int run_begin = 0; |
| |
| // Subsequent items |
| for (int i = 1; i < num_items; ++i) |
| { |
| if (!equality_op(previous, h_in[i])) |
| { |
| if ((RLE_METHOD != NON_TRIVIAL) || (length > 1)) |
| { |
| h_unique_reference[num_runs] = previous; |
| h_offsets_reference[num_runs] = run_begin; |
| h_lengths_reference[num_runs] = length; |
| num_runs++; |
| } |
| length = 1; |
| run_begin = i; |
| } |
| else |
| { |
| length++; |
| } |
| previous = h_in[i]; |
| } |
| |
| if ((RLE_METHOD != NON_TRIVIAL) || (length > 1)) |
| { |
| h_unique_reference[num_runs] = previous; |
| h_offsets_reference[num_runs] = run_begin; |
| h_lengths_reference[num_runs] = length; |
| num_runs++; |
| } |
| |
| return num_runs; |
| } |
| |
|
|
|
|
| /** |
| * Test DeviceRunLengthEncode for a given problem input |
| */ |
| template < |
| RleMethod RLE_METHOD, |
| Backend BACKEND, |
| typename DeviceInputIteratorT, |
| typename T, |
| typename OffsetT, |
| typename LengthT, |
| typename EqualityOp> |
| void Test( |
| DeviceInputIteratorT d_in, |
| T *h_unique_reference, |
| OffsetT *h_offsets_reference, |
| LengthT *h_lengths_reference, |
| EqualityOp equality_op, |
| int num_runs, |
| int num_items) |
| { |
| // Allocate device output arrays and number of segments |
| T* d_unique_out = NULL; |
| LengthT* d_offsets_out = NULL; |
| OffsetT* d_lengths_out = NULL; |
| int* d_num_runs = NULL; |
| |
| if (RLE_METHOD == RLE) |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_unique_out, sizeof(T) * num_items)); |
| if (RLE_METHOD == NON_TRIVIAL) |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_offsets_out, sizeof(OffsetT) * num_items)); |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_lengths_out, sizeof(LengthT) * num_items)); |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_num_runs, sizeof(int))); |
| |
| // Allocate CDP device arrays |
| size_t* d_temp_storage_bytes = NULL; |
| cudaError_t* d_cdp_error = NULL; |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_temp_storage_bytes, sizeof(size_t) * 1)); |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_cdp_error, sizeof(cudaError_t) * 1)); |
| |
| // Allocate temporary storage |
| void* d_temp_storage = NULL; |
| size_t temp_storage_bytes = 0; |
| CubDebugExit(Dispatch(Int2Type<RLE_METHOD>(), Int2Type<BACKEND>(), 1, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_in, d_unique_out, d_offsets_out, d_lengths_out, d_num_runs, equality_op, num_items, 0, true)); |
| CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); |
| |
| // Clear device output arrays |
| if (RLE_METHOD == RLE) |
| CubDebugExit(cudaMemset(d_unique_out, 0, sizeof(T) * num_items)); |
| if (RLE_METHOD == NON_TRIVIAL) |
| CubDebugExit(cudaMemset(d_offsets_out, 0, sizeof(OffsetT) * num_items)); |
| CubDebugExit(cudaMemset(d_lengths_out, 0, sizeof(LengthT) * num_items)); |
| CubDebugExit(cudaMemset(d_num_runs, 0, sizeof(int))); |
| |
| // Run warmup/correctness iteration |
| CubDebugExit(Dispatch(Int2Type<RLE_METHOD>(), Int2Type<BACKEND>(), 1, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_in, d_unique_out, d_offsets_out, d_lengths_out, d_num_runs, equality_op, num_items, 0, true)); |
| |
| // Check for correctness (and display results, if specified) |
| int compare0 = 0; |
| int compare1 = 0; |
| int compare2 = 0; |
| int compare3 = 0; |
| |
| if (RLE_METHOD == RLE) |
| { |
| compare0 = CompareDeviceResults(h_unique_reference, d_unique_out, num_runs, true, g_verbose); |
| printf("\t Keys %s\n", compare0 ? "FAIL" : "PASS"); |
| } |
| |
| if (RLE_METHOD != RLE) |
| { |
| compare1 = CompareDeviceResults(h_offsets_reference, d_offsets_out, num_runs, true, g_verbose); |
| printf("\t Offsets %s\n", compare1 ? "FAIL" : "PASS"); |
| } |
| |
| if (RLE_METHOD != CSR) |
| { |
| compare2 = CompareDeviceResults(h_lengths_reference, d_lengths_out, num_runs, true, g_verbose); |
| printf("\t Lengths %s\n", compare2 ? "FAIL" : "PASS"); |
| } |
| |
| compare3 = CompareDeviceResults(&num_runs, d_num_runs, 1, true, g_verbose); |
| printf("\t Count %s\n", compare3 ? "FAIL" : "PASS"); |
| |
| // Flush any stdout/stderr |
| fflush(stdout); |
| fflush(stderr); |
| |
| // Performance |
| GpuTimer gpu_timer; |
| gpu_timer.Start(); |
| CubDebugExit(Dispatch(Int2Type<RLE_METHOD>(), Int2Type<BACKEND>(), g_timing_iterations, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_in, d_unique_out, d_offsets_out, d_lengths_out, d_num_runs, equality_op, num_items, 0, false)); |
| gpu_timer.Stop(); |
| float elapsed_millis = gpu_timer.ElapsedMillis(); |
| |
| // Display performance |
| if (g_timing_iterations > 0) |
| { |
| float avg_millis = elapsed_millis / g_timing_iterations; |
| float giga_rate = float(num_items) / avg_millis / 1000.0f / 1000.0f; |
| int bytes_moved = (num_items * sizeof(T)) + (num_runs * (sizeof(OffsetT) + sizeof(LengthT))); |
| float giga_bandwidth = float(bytes_moved) / avg_millis / 1000.0f / 1000.0f; |
| printf(", %.3f avg ms, %.3f billion items/s, %.3f logical GB/s", avg_millis, giga_rate, giga_bandwidth); |
| } |
| printf("\n\n"); |
| |
| // Flush any stdout/stderr |
| fflush(stdout); |
| fflush(stderr); |
| |
| // Cleanup |
| if (d_unique_out) CubDebugExit(g_allocator.DeviceFree(d_unique_out)); |
| if (d_offsets_out) CubDebugExit(g_allocator.DeviceFree(d_offsets_out)); |
| if (d_lengths_out) CubDebugExit(g_allocator.DeviceFree(d_lengths_out)); |
| if (d_num_runs) CubDebugExit(g_allocator.DeviceFree(d_num_runs)); |
| if (d_temp_storage_bytes) CubDebugExit(g_allocator.DeviceFree(d_temp_storage_bytes)); |
| if (d_cdp_error) CubDebugExit(g_allocator.DeviceFree(d_cdp_error)); |
| if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage)); |
| |
| // Correctness asserts |
| AssertEquals(0, compare0 | compare1 | compare2 | compare3); |
| } |
| |
|
|
| /** |
| * Test DeviceRunLengthEncode on pointer type |
| */ |
| template < |
| RleMethod RLE_METHOD, |
| Backend BACKEND, |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void TestPointer( |
| int num_items, |
| int entropy_reduction, |
| int max_segment) |
| { |
| // Allocate host arrays |
| T* h_in = new T[num_items]; |
| T* h_unique_reference = new T[num_items]; |
| OffsetT* h_offsets_reference = new OffsetT[num_items]; |
| LengthT* h_lengths_reference = new LengthT[num_items]; |
| |
| for (int i = 0; i < num_items; ++i) |
| InitValue(INTEGER_SEED, h_offsets_reference[i], 1); |
| |
| // Initialize problem and solution |
| Equality equality_op; |
| Initialize(entropy_reduction, h_in, num_items, max_segment); |
| |
| int num_runs = Solve<RLE_METHOD>(h_in, h_unique_reference, h_offsets_reference, h_lengths_reference, equality_op, num_items); |
| |
| printf("\nPointer %s cub::%s on %d items, %d segments (avg run length %.3f), {%s key, %s offset, %s length}, max_segment %d, entropy_reduction %d\n", |
| (RLE_METHOD == RLE) ? "DeviceReduce::RunLengthEncode" : (RLE_METHOD == NON_TRIVIAL) ? "DeviceRunLengthEncode::NonTrivialRuns" : "Other", |
| (BACKEND == CDP) ? "CDP CUB" : (BACKEND == THRUST) ? "Thrust" : "CUB", |
| num_items, num_runs, float(num_items) / num_runs, |
| typeid(T).name(), typeid(OffsetT).name(), typeid(LengthT).name(), |
| max_segment, entropy_reduction); |
| fflush(stdout); |
| |
| // Allocate problem device arrays |
| T* d_in = NULL; |
| CubDebugExit(g_allocator.DeviceAllocate((void**)&d_in, sizeof(T) * num_items)); |
| |
| // Initialize device input |
| CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(T) * num_items, cudaMemcpyHostToDevice)); |
| |
| // Run Test |
| Test<RLE_METHOD, BACKEND>(d_in, h_unique_reference, h_offsets_reference, h_lengths_reference, equality_op, num_runs, num_items); |
| |
| // Cleanup |
| if (h_in) delete[] h_in; |
| if (h_unique_reference) delete[] h_unique_reference; |
| if (h_offsets_reference) delete[] h_offsets_reference; |
| if (h_lengths_reference) delete[] h_lengths_reference; |
| if (d_in) CubDebugExit(g_allocator.DeviceFree(d_in)); |
| } |
| |
|
|
| /** |
| * Test on iterator type |
| */ |
| template < |
| RleMethod RLE_METHOD, |
| Backend BACKEND, |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void TestIterator( |
| int num_items, |
| Int2Type<true> /*is_primitive*/) |
| { |
| // Allocate host arrays |
| T* h_unique_reference = new T[num_items]; |
| OffsetT* h_offsets_reference = new OffsetT[num_items]; |
| LengthT* h_lengths_reference = new LengthT[num_items]; |
| |
| T one_val; |
| InitValue(INTEGER_SEED, one_val, 1); |
| ConstantInputIterator<T, int> h_in(one_val); |
| |
| // Initialize problem and solution |
| Equality equality_op; |
| int num_runs = Solve<RLE_METHOD>(h_in, h_unique_reference, h_offsets_reference, h_lengths_reference, equality_op, num_items); |
| |
| printf("\nIterator %s cub::%s on %d items, %d segments (avg run length %.3f), {%s key, %s offset, %s length}\n", |
| (RLE_METHOD == RLE) ? "DeviceReduce::RunLengthEncode" : (RLE_METHOD == NON_TRIVIAL) ? "DeviceRunLengthEncode::NonTrivialRuns" : "Other", |
| (BACKEND == CDP) ? "CDP CUB" : (BACKEND == THRUST) ? "Thrust" : "CUB", |
| num_items, num_runs, float(num_items) / num_runs, |
| typeid(T).name(), typeid(OffsetT).name(), typeid(LengthT).name()); |
| fflush(stdout); |
| |
| // Run Test |
| Test<RLE_METHOD, BACKEND>(h_in, h_unique_reference, h_offsets_reference, h_lengths_reference, equality_op, num_runs, num_items); |
| |
| // Cleanup |
| if (h_unique_reference) delete[] h_unique_reference; |
| if (h_offsets_reference) delete[] h_offsets_reference; |
| if (h_lengths_reference) delete[] h_lengths_reference; |
| } |
| |
|
|
| template < |
| RleMethod RLE_METHOD, |
| Backend BACKEND, |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void TestIterator( |
| int /*num_items*/, |
| Int2Type<false> /*is_primitive*/) |
| {} |
| |
|
|
| /** |
| * Test different gen modes |
| */ |
| template < |
| RleMethod RLE_METHOD, |
| Backend BACKEND, |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void Test( |
| int num_items) |
| { |
| // Test iterator (one run) |
| TestIterator<RLE_METHOD, BACKEND, T, OffsetT, LengthT>(num_items, Int2Type<Traits<T>::PRIMITIVE>()); |
| |
| // num_items runs |
| TestPointer<RLE_METHOD, BACKEND, T, OffsetT, LengthT>(num_items, 0, 1); |
| |
| // Evaluate different run lengths |
| for (int max_segment = 3; max_segment < CUB_MIN(num_items, (unsigned short) -1); max_segment *= 3) |
| { |
| // Uniform selection run length |
| TestPointer<RLE_METHOD, BACKEND, T, OffsetT, LengthT>(num_items, 0, max_segment); |
| |
| // Reduced-entropy run length |
| TestPointer<RLE_METHOD, BACKEND, T, OffsetT, LengthT>(num_items, 4, max_segment); |
| } |
| } |
| |
|
|
| /** |
| * Test different dispatch |
| */ |
| template < |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void TestDispatch( |
| int num_items) |
| { |
| Test<RLE, CUB, T, OffsetT, LengthT>(num_items); |
| Test<NON_TRIVIAL, CUB, T, OffsetT, LengthT>(num_items); |
| |
| #ifdef CUB_CDP |
| Test<RLE, CDP, T, OffsetT, LengthT>(num_items); |
| Test<NON_TRIVIAL, CDP, T, OffsetT, LengthT>(num_items); |
| #endif |
| } |
| |
| |
| /** |
| * Test different input sizes |
| */ |
| template < |
| typename T, |
| typename OffsetT, |
| typename LengthT> |
| void TestSize( |
| int num_items) |
| { |
| if (num_items < 0) |
| { |
| TestDispatch<T, OffsetT, LengthT>(0); |
| TestDispatch<T, OffsetT, LengthT>(1); |
| TestDispatch<T, OffsetT, LengthT>(100); |
| TestDispatch<T, OffsetT, LengthT>(10000); |
| TestDispatch<T, OffsetT, LengthT>(1000000); |
| |
| // Randomly select problem size between 1:10,000,000 |
| unsigned int max_int = (unsigned int) -1; |
| for (int i = 0; i < 10; ++i) |
| { |
| unsigned int num_items; |
| RandomBits(num_items); |
| num_items = (unsigned int) ((double(num_items) * double(10000000)) / double(max_int)); |
| num_items = CUB_MAX(1, num_items); |
| TestDispatch<T, OffsetT, LengthT>(num_items); |
| } |
| } |
| else |
| { |
| TestDispatch<T, OffsetT, LengthT>(num_items); |
| } |
| |
| } |
|
|
|
|
| //--------------------------------------------------------------------- |
| // Main |
| //--------------------------------------------------------------------- |
|
|
| /** |
| * Main |
| */ |
| int main(int argc, char** argv) |
| { |
| int num_items = -1; |
| int entropy_reduction = 0; |
| int max_segment = 1000; |
| |
| // Initialize command line |
| CommandLineArgs args(argc, argv); |
| g_verbose = args.CheckCmdLineFlag("v"); |
| args.GetCmdLineArgument("n", num_items); |
| args.GetCmdLineArgument("i", g_timing_iterations); |
| args.GetCmdLineArgument("repeat", g_repeat); |
| args.GetCmdLineArgument("maxseg", max_segment); |
| args.GetCmdLineArgument("entropy", entropy_reduction); |
| |
| // Print usage |
| if (args.CheckCmdLineFlag("help")) |
| { |
| printf("%s " |
| "[--n=<input items> " |
| "[--i=<timing iterations> " |
| "[--device=<device-id>] " |
| "[--maxseg=<max segment length>]" |
| "[--entropy=<segment length bit entropy reduction rounds>]" |
| "[--repeat=<repetitions of entire test suite>]" |
| "[--v] " |
| "[--cdp]" |
| "\n", argv[0]); |
| exit(0); |
| } |
| |
| // Initialize device |
| CubDebugExit(args.DeviceInit()); |
| printf("\n"); |
| |
| // Get ptx version |
| int ptx_version = 0; |
| CubDebugExit(PtxVersion(ptx_version)); |
| |
| #ifdef QUICKER_TEST |
| |
| // Compile/run basic CUB test |
| if (num_items < 0) num_items = 32000000; |
| |
| TestPointer<RLE, CUB, int, int, int>( num_items, entropy_reduction, max_segment); |
| TestPointer<NON_TRIVIAL, CUB, int, int, int>( num_items, entropy_reduction, max_segment); |
| TestIterator<RLE, CUB, float, int, int>( num_items, Int2Type<Traits<float>::PRIMITIVE>()); |
| |
|
|
| #elif defined(QUICK_TEST) |
| |
| // Compile/run quick tests |
| if (num_items < 0) num_items = 32000000; |
| |
| TestPointer<RLE, CUB, int, int, int>( num_items, entropy_reduction, max_segment); |
| TestPointer<RLE, THRUST, int, int, int>( num_items, entropy_reduction, max_segment); |
| |
| #else |
|
|
| // Compile/run thorough tests |
| for (int i = 0; i <= g_repeat; ++i) |
| { |
| // Test different input types |
| TestSize<char, int, int>(num_items); |
| TestSize<short, int, int>(num_items); |
| TestSize<int, int, int>(num_items); |
| TestSize<long, int, int>(num_items); |
| TestSize<long long, int, int>(num_items); |
| TestSize<float, int, int>(num_items); |
| TestSize<double, int, int>(num_items); |
| |
| TestSize<uchar2, int, int>(num_items); |
| TestSize<uint2, int, int>(num_items); |
| TestSize<uint3, int, int>(num_items); |
| TestSize<uint4, int, int>(num_items); |
| TestSize<ulonglong4, int, int>(num_items); |
| TestSize<TestFoo, int, int>(num_items); |
| TestSize<TestBar, int, int>(num_items); |
| } |
| |
| #endif |
|
|
| return 0; |
| } |
| |
|
|
|
|
|
|
