rocPRIM
device_adjacent_difference_config.hpp
1 // Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
2 //
3 // Permission is hereby granted, free of charge, to any person obtaining a copy
4 // of this software and associated documentation files (the "Software"), to deal
5 // in the Software without restriction, including without limitation the rights
6 // to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
7 // copies of the Software, and to permit persons to whom the Software is
8 // furnished to do so, subject to the following conditions:
9 //
10 // The above copyright notice and this permission notice shall be included in
11 // all copies or substantial portions of the Software.
12 //
13 // THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
14 // IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
15 // FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
16 // AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
17 // LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
18 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
19 // THE SOFTWARE.
20 
21 #ifndef ROCPRIM_DEVICE_DEVICE_ADJACENT_DIFFERENCE_CONFIG_HPP_
22 #define ROCPRIM_DEVICE_DEVICE_ADJACENT_DIFFERENCE_CONFIG_HPP_
23 
24 #include <type_traits>
25 
26 #include "../config.hpp"
27 #include "../detail/various.hpp"
28 #include "../functional.hpp"
29 
30 #include "config_types.hpp"
31 
32 #include "../block/block_load.hpp"
33 #include "../block/block_store.hpp"
34 
37 
38 BEGIN_ROCPRIM_NAMESPACE
39 
48 template <unsigned int BlockSize,
49  unsigned int ItemsPerThread,
52  unsigned int SizeLimit = ROCPRIM_GRID_SIZE_LIMIT>
53 struct adjacent_difference_config : kernel_config<BlockSize, ItemsPerThread, SizeLimit>
54 {
55  static constexpr block_load_method load_method = LoadMethod;
56  static constexpr block_store_method store_method = StoreMethod;
57 };
58 
59 namespace detail
60 {
61 
62 template <class Value>
64 {
65  static constexpr unsigned int item_scale
66  = ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Value), sizeof(int));
67 
68  using type = adjacent_difference_config<256, ::rocprim::max(1u, 16u / item_scale)>;
69 };
70 
71 template <unsigned int TargetArch, class Value>
73  : select_arch<TargetArch, adjacent_difference_config_fallback<Value>>
74 {
75 };
76 
77 } // end namespace detail
78 
79 END_ROCPRIM_NAMESPACE
80 
82 // end of group primitivesmodule_deviceconfigs
83 
84 #endif // ROCPRIM_DEVICE_DEVICE_ADJACENT_DIFFERENCE_CONFIG_HPP_
ROCPRIM_HOST_DEVICE constexpr T max(const T &a, const T &b)
Returns the maximum of its arguments.
Definition: functional.hpp:55
block_store_method
block_store_method enumerates the methods available to store a striped arrangement of items into a bl...
Definition: block_store.hpp:41
Configuration of device-level adjacent_difference primitives.
Definition: device_adjacent_difference_config.hpp:53
Deprecated: Configuration of device-level scan primitives.
Definition: block_histogram.hpp:62
Definition: device_adjacent_difference_config.hpp:63
block_load_method
block_load_method enumerates the methods available to load data from continuous memory into a blocked...
Definition: block_load.hpp:41
A striped arrangement of data from continuous memory is locally transposed into a blocked arrangement...
Definition: config_types.hpp:140
static constexpr block_load_method load_method
input values are loaded using this method
Definition: device_adjacent_difference_config.hpp:55
static constexpr block_store_method store_method
input values are stored using this method
Definition: device_adjacent_difference_config.hpp:56
A blocked arrangement of items is locally transposed and stored as a striped arrangement of data on c...
Configuration of particular kernels launched by device-level operation.
Definition: config_types.hpp:84
Definition: device_adjacent_difference_config.hpp:72