From ffd1601dc5771f5bf3ad1e315f34e1a52e868ba7 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Thu, 12 Nov 2020 15:48:22 -0500 Subject: [PATCH 01/11] Fix radix sort perf regression. The new tuning parameters for radix sort are causing a performance regression (2MiB input, 16-bit keys, GV100, 21.62 -> 16.01 GiB/sec). See referenced NVBug for details. Bug 200676467 --- cub/device/dispatch/dispatch_radix_sort.cuh | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/device/dispatch/dispatch_radix_sort.cuh index 04b5461f91..7d2bedf77f 100644 --- a/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/device/dispatch/dispatch_radix_sort.cuh @@ -865,10 +865,10 @@ struct DeviceRadixSortPolicy // ScanPolicy - typedef AgentScanPolicy <256, 23, OffsetT, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; + typedef AgentScanPolicy <512, 23, OffsetT, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; // Downsweep policies - typedef AgentRadixSortDownsweepPolicy <256, 23, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MATCH, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS> DownsweepPolicy; + typedef AgentRadixSortDownsweepPolicy <512, 23, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MATCH, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS> DownsweepPolicy; typedef AgentRadixSortDownsweepPolicy <(sizeof(KeyT) > 1) ? 256 : 128, 47, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS - 1> AltDownsweepPolicy; // Upsweep policies @@ -879,8 +879,8 @@ struct DeviceRadixSortPolicy typedef AgentRadixSortDownsweepPolicy <256, 19, DominantT, BLOCK_LOAD_DIRECT, LOAD_LDG, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SINGLE_TILE_RADIX_BITS> SingleTilePolicy; // Segmented policies - typedef AgentRadixSortDownsweepPolicy <128, 39, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS> SegmentedPolicy; - typedef AgentRadixSortDownsweepPolicy <256, 11, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS - 1> AltSegmentedPolicy; + typedef AgentRadixSortDownsweepPolicy <192, 39, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS> SegmentedPolicy; + typedef AgentRadixSortDownsweepPolicy <384, 11, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS - 1> AltSegmentedPolicy; }; @@ -907,10 +907,10 @@ struct DeviceRadixSortPolicy RADIX_SORT_STORE_DIRECT, ONESWEEP_RADIX_BITS> OnesweepPolicy; // ScanPolicy - typedef AgentScanPolicy <256, 23, OffsetT, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; + typedef AgentScanPolicy <512, 23, OffsetT, BLOCK_LOAD_WARP_TRANSPOSE, LOAD_DEFAULT, BLOCK_STORE_WARP_TRANSPOSE, BLOCK_SCAN_RAKING_MEMOIZE> ScanPolicy; // Downsweep policies - typedef AgentRadixSortDownsweepPolicy <256, 23, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MATCH, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS> DownsweepPolicy; + typedef AgentRadixSortDownsweepPolicy <512, 23, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MATCH, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS> DownsweepPolicy; typedef AgentRadixSortDownsweepPolicy <(sizeof(KeyT) > 1) ? 256 : 128, 47, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, PRIMARY_RADIX_BITS - 1> AltDownsweepPolicy; // Upsweep policies @@ -921,8 +921,8 @@ struct DeviceRadixSortPolicy typedef AgentRadixSortDownsweepPolicy <256, 19, DominantT, BLOCK_LOAD_DIRECT, LOAD_LDG, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SINGLE_TILE_RADIX_BITS> SingleTilePolicy; // Segmented policies - typedef AgentRadixSortDownsweepPolicy <128, 39, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS> SegmentedPolicy; - typedef AgentRadixSortDownsweepPolicy <256, 11, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS - 1> AltSegmentedPolicy; + typedef AgentRadixSortDownsweepPolicy <192, 39, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS> SegmentedPolicy; + typedef AgentRadixSortDownsweepPolicy <384, 11, DominantT, BLOCK_LOAD_TRANSPOSE, LOAD_DEFAULT, RADIX_RANK_MEMOIZE, BLOCK_SCAN_WARP_SCANS, SEGMENTED_RADIX_BITS - 1> AltSegmentedPolicy; }; From 618a46c27764f0e0b86fb3643a572ed039180ad8 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Wed, 18 Nov 2020 13:48:39 -0500 Subject: [PATCH 02/11] Update with release info for 1.11.0. --- CHANGELOG.md | 76 +++++++++++++++++++++++++++++++++++++++++++++++++--- README.md | 1 + 2 files changed, 74 insertions(+), 3 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index d0c8df0f74..0bf219e7c4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,7 +1,77 @@ -# Thrust 1.11.0 (NVIDIA HPC SDK 20.11) +# CUB 1.11.0 -CUB 1.11.0 is the major release accompanying the NVIDIA HPC SDK 20.11 release. -It introduces a new radix sort backend with improved performance. +## Summary + +CUB 1.11.0 is a major release providing bugfixes and performance enhancements. + +It includes a new `DeviceRadixSort` backend that improves performance by up to +2x on supported keys and hardware. + +Our CMake package and build system continue to see improvements +with `add_subdirectory` support, installation rules, status messages, and other +features that make CUB easier to use from CMake projects. + +The release includes several other bugfixes and modernizations, and received +updates from 11 contributors. + +## Breaking Changes + +- NVIDIA/cub#201: The intermediate accumulator type used when `DeviceScan` is + invoked with different input/output types is now consistent + with [P0571](https://wg21.link/P0571). This may produce different results for + some edge cases when compared with earlier releases of CUB. + +## New Features + +- NVIDIA/cub#204: Faster `DeviceRadixSort`, up to 2x performance increase for + 32/64-bit keys on Pascal and up (SM60+). Thanks to Andy Adinets for this + contribution. +- Unroll loops in `BlockRadixRank` to improve performance for 32-bit keys by + 1.5-2x on Clang CUDA. Thanks to Justin Lebar for this contribution. +- NVIDIA/cub#200: Allow CUB to be added to CMake projects via `add_subdirectory`. +- NVIDIA/cub#214: Optionally add install rules when included with + CMake's `add_subdirectory`. Thanks to Kai Germaschewski for this contribution. + +## Bug Fixes + +- NVIDIA/cub#215: Fix integer truncation in `AgentReduceByKey`, `AgentScan`, + and `AgentSegmentFixup`. Thanks to Rory Mitchell for this contribution. +- NVIDIA/cub#225: Fix compile-time regression when defining `CUB_NS_PREFIX` + /`CUB_NS_POSTFIX` macro. Thanks to Elias Stehle for this contribution. +- NVIDIA/cub#210: Fix some edge cases in `DeviceScan`: + - Use values from the input when padding temporary buffers. This prevents + custom functors from getting unexpected values. + - Prevent integer truncation when using large indices via the `DispatchScan` + layer. + - Use timesliced reads/writes for types > 128 bytes. +- NVIDIA/cub#217: Fix and add test for cmake package install rules. Thanks to + Keith Kraus and Kai Germaschewski for testing and discussion. +- NVIDIA/cub#170, NVIDIA/cub#233: Update CUDA version checks to behave on Clang + CUDA and `nvc++`. Thanks to Artem Belevich, Andrew Corrigan, and David Olsen + for these contributions. +- NVIDIA/cub#220, NVIDIA/cub#216: Various fixes for Clang CUDA. Thanks to Andrew + Corrigan for these contributions. +- NVIDIA/cub#231: Fix signedness mismatch warnings in unit tests. +- NVIDIA/cub#231: Suppress GPU deprecation warnings. +- NVIDIA/cub#214: Use semantic versioning rules for our CMake package's + compatibility checks. Thanks to Kai Germaschewski for this contribution. +- NVIDIA/cub#214: Use `FindPackageHandleStandardArgs` to print standard status + messages when our CMake package is found. Thanks to Kai Germaschewski for this + contribution. +- NVIDIA/cub#207: Fix `CubDebug` usage + in `CachingDeviceAllocator::DeviceAllocate`. Thanks to Andreas Hehn for this + contribution. +- Fix documentation for `DevicePartition`. Thanks to ByteHamster for this + contribution. +- Clean up unused code in `DispatchScan`. Thanks to ByteHamster for this + contribution. + +## Other Enhancements + +- NVIDIA/cub#213: Remove tuning policies for unsupported hardware ( Date: Tue, 24 Nov 2020 18:19:21 +0000 Subject: [PATCH 03/11] Fix incorrect issue number in changelog --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 0bf219e7c4..4cdc230635 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -363,7 +363,7 @@ CUB 1.7.2 is a minor release. ## Bug Fixes -- #104: Device-wide reduction is now "run-to-run" deterministic for +- #108: Device-wide reduction is now "run-to-run" deterministic for pseudo-associative reduction operators (like floating point addition). # CUB 1.7.1 From cb77f97d79a0569874e51f017d51529d922aa0a0 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 13:23:55 -0500 Subject: [PATCH 04/11] Fix `maybe-uninitialized` warnings on GCC 5. --- experimental/spmv_compare.cu | 2 +- test/test_device_scan.cu | 2 +- test/test_device_select_if.cu | 2 +- test/test_device_select_unique.cu | 2 +- test/test_grid_barrier.cu | 2 +- 5 files changed, 5 insertions(+), 5 deletions(-) diff --git a/experimental/spmv_compare.cu b/experimental/spmv_compare.cu index b64297d8db..8b623a5e56 100644 --- a/experimental/spmv_compare.cu +++ b/experimental/spmv_compare.cu @@ -200,7 +200,7 @@ float TestGpuCsrIoProxy( CubDebugExit(cudaGetDevice(&device_ordinal)); // Get device SM version - int sm_version; + int sm_version = 0; CubDebugExit(SmVersion(sm_version, device_ordinal)); void (*kernel)(SpmvParams, TexItr) = NonZeroIoKernel; diff --git a/test/test_device_scan.cu b/test/test_device_scan.cu index 30a7b26f8a..8318d01721 100644 --- a/test/test_device_scan.cu +++ b/test/test_device_scan.cu @@ -992,7 +992,7 @@ int main(int argc, char** argv) CubDebugExit(cudaGetDevice(&device_ordinal)); // Get device SM version - int sm_version; + int sm_version = 0; CubDebugExit(SmVersion(sm_version, device_ordinal)); // Compile/run quick tests diff --git a/test/test_device_select_if.cu b/test/test_device_select_if.cu index a4488276ca..a870c7746d 100644 --- a/test/test_device_select_if.cu +++ b/test/test_device_select_if.cu @@ -1003,7 +1003,7 @@ int main(int argc, char** argv) CubDebugExit(cudaGetDevice(&device_ordinal)); // Get device SM version - int sm_version; + int sm_version = 0; CubDebugExit(SmVersion(sm_version, device_ordinal)); // Compile/run quick tests diff --git a/test/test_device_select_unique.cu b/test/test_device_select_unique.cu index 1c13fba15d..ed0549edee 100644 --- a/test/test_device_select_unique.cu +++ b/test/test_device_select_unique.cu @@ -598,7 +598,7 @@ int main(int argc, char** argv) CubDebugExit(cudaGetDevice(&device_ordinal)); // Get device SM version - int sm_version; + int sm_version = 0; CubDebugExit(SmVersion(sm_version, device_ordinal)); // Compile/run quick tests diff --git a/test/test_grid_barrier.cu b/test/test_grid_barrier.cu index e6e3b8125c..9c989e5397 100644 --- a/test/test_grid_barrier.cu +++ b/test/test_grid_barrier.cu @@ -104,7 +104,7 @@ int main(int argc, char** argv) CubDebugExit(cudaGetDevice(&device_ordinal)); // Get device SM version - int sm_version; + int sm_version = 0; CubDebugExit(SmVersion(sm_version, device_ordinal)); // Get SM properties From 22e7b0f59e68aa888d21efda44715f994fde6473 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 14:02:07 -0500 Subject: [PATCH 05/11] Remove old project config files. If anyone is using these, just save a copy locally and add them to your repo's .git/info/exclude file, which uses regular .gitignore syntax. --- .cproject | 1223 -------------------- .project | 27 - .settings/.gitignore | 1 - .settings/org.eclipse.cdt.codan.core.prefs | 72 -- .settings/org.eclipse.cdt.core.prefs | 177 --- .settings/org.eclipse.cdt.ui.prefs | 3 - .settings/org.eclipse.core.runtime.prefs | 4 - 7 files changed, 1507 deletions(-) delete mode 100644 .cproject delete mode 100644 .project delete mode 100644 .settings/.gitignore delete mode 100644 .settings/org.eclipse.cdt.codan.core.prefs delete mode 100644 .settings/org.eclipse.cdt.core.prefs delete mode 100644 .settings/org.eclipse.cdt.ui.prefs delete mode 100644 .settings/org.eclipse.core.runtime.prefs diff --git a/.cproject b/.cproject deleted file mode 100644 index e76d1da67e..0000000000 --- a/.cproject +++ /dev/null @@ -1,1223 +0,0 @@ - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - diff --git a/.project b/.project deleted file mode 100644 index 7aca9e046b..0000000000 --- a/.project +++ /dev/null @@ -1,27 +0,0 @@ - - - GIT_CUB - - - - - - org.eclipse.cdt.managedbuilder.core.genmakebuilder - clean,full,incremental, - - - - - org.eclipse.cdt.managedbuilder.core.ScannerConfigBuilder - full,incremental, - - - - - - org.eclipse.cdt.core.cnature - org.eclipse.cdt.managedbuilder.core.managedBuildNature - org.eclipse.cdt.managedbuilder.core.ScannerConfigNature - org.eclipse.cdt.core.ccnature - - diff --git a/.settings/.gitignore b/.settings/.gitignore deleted file mode 100644 index d81d4c414c..0000000000 --- a/.settings/.gitignore +++ /dev/null @@ -1 +0,0 @@ -/language.settings.xml diff --git a/.settings/org.eclipse.cdt.codan.core.prefs b/.settings/org.eclipse.cdt.codan.core.prefs deleted file mode 100644 index 64da7771b6..0000000000 --- a/.settings/org.eclipse.cdt.codan.core.prefs +++ /dev/null @@ -1,72 +0,0 @@ -eclipse.preferences.version=1 -org.eclipse.cdt.codan.checkers.errnoreturn=Warning -org.eclipse.cdt.codan.checkers.errnoreturn.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},implicit\=>false} -org.eclipse.cdt.codan.checkers.errreturnvalue=Error -org.eclipse.cdt.codan.checkers.errreturnvalue.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.checkers.nocommentinside=-Error -org.eclipse.cdt.codan.checkers.nocommentinside.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.checkers.nolinecomment=-Error -org.eclipse.cdt.codan.checkers.nolinecomment.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.checkers.noreturn=Error -org.eclipse.cdt.codan.checkers.noreturn.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},implicit\=>false} -org.eclipse.cdt.codan.internal.checkers.AbstractClassCreation=Error -org.eclipse.cdt.codan.internal.checkers.AbstractClassCreation.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.AmbiguousProblem=Error -org.eclipse.cdt.codan.internal.checkers.AmbiguousProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.AssignmentInConditionProblem=Warning -org.eclipse.cdt.codan.internal.checkers.AssignmentInConditionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.AssignmentToItselfProblem=Error -org.eclipse.cdt.codan.internal.checkers.AssignmentToItselfProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.CaseBreakProblem=Warning -org.eclipse.cdt.codan.internal.checkers.CaseBreakProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},no_break_comment\=>"no break",last_case_param\=>true,empty_case_param\=>false} -org.eclipse.cdt.codan.internal.checkers.CatchByReference=Warning -org.eclipse.cdt.codan.internal.checkers.CatchByReference.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},unknown\=>false,exceptions\=>()} -org.eclipse.cdt.codan.internal.checkers.CircularReferenceProblem=Error -org.eclipse.cdt.codan.internal.checkers.CircularReferenceProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.ClassMembersInitialization=Warning -org.eclipse.cdt.codan.internal.checkers.ClassMembersInitialization.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},skip\=>true} -org.eclipse.cdt.codan.internal.checkers.FieldResolutionProblem=Error -org.eclipse.cdt.codan.internal.checkers.FieldResolutionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.FunctionResolutionProblem=Error -org.eclipse.cdt.codan.internal.checkers.FunctionResolutionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.InvalidArguments=Error -org.eclipse.cdt.codan.internal.checkers.InvalidArguments.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.InvalidTemplateArgumentsProblem=Error -org.eclipse.cdt.codan.internal.checkers.InvalidTemplateArgumentsProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.LabelStatementNotFoundProblem=Error -org.eclipse.cdt.codan.internal.checkers.LabelStatementNotFoundProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.MemberDeclarationNotFoundProblem=Error -org.eclipse.cdt.codan.internal.checkers.MemberDeclarationNotFoundProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.MethodResolutionProblem=Error -org.eclipse.cdt.codan.internal.checkers.MethodResolutionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.NamingConventionFunctionChecker=-Info -org.eclipse.cdt.codan.internal.checkers.NamingConventionFunctionChecker.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},pattern\=>"^[a-z]",macro\=>true,exceptions\=>()} -org.eclipse.cdt.codan.internal.checkers.NonVirtualDestructorProblem=Warning -org.eclipse.cdt.codan.internal.checkers.NonVirtualDestructorProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.OverloadProblem=Error -org.eclipse.cdt.codan.internal.checkers.OverloadProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.RedeclarationProblem=Error -org.eclipse.cdt.codan.internal.checkers.RedeclarationProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.RedefinitionProblem=Error -org.eclipse.cdt.codan.internal.checkers.RedefinitionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.ReturnStyleProblem=-Warning -org.eclipse.cdt.codan.internal.checkers.ReturnStyleProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.ScanfFormatStringSecurityProblem=-Warning -org.eclipse.cdt.codan.internal.checkers.ScanfFormatStringSecurityProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.StatementHasNoEffectProblem=Warning -org.eclipse.cdt.codan.internal.checkers.StatementHasNoEffectProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},macro\=>true,exceptions\=>()} -org.eclipse.cdt.codan.internal.checkers.SuggestedParenthesisProblem=Warning -org.eclipse.cdt.codan.internal.checkers.SuggestedParenthesisProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},paramNot\=>false} -org.eclipse.cdt.codan.internal.checkers.SuspiciousSemicolonProblem=Warning -org.eclipse.cdt.codan.internal.checkers.SuspiciousSemicolonProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},else\=>false,afterelse\=>false} -org.eclipse.cdt.codan.internal.checkers.TypeResolutionProblem=Error -org.eclipse.cdt.codan.internal.checkers.TypeResolutionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -org.eclipse.cdt.codan.internal.checkers.UnusedFunctionDeclarationProblem=Warning -org.eclipse.cdt.codan.internal.checkers.UnusedFunctionDeclarationProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},macro\=>true} -org.eclipse.cdt.codan.internal.checkers.UnusedStaticFunctionProblem=Warning -org.eclipse.cdt.codan.internal.checkers.UnusedStaticFunctionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},macro\=>true} -org.eclipse.cdt.codan.internal.checkers.UnusedVariableDeclarationProblem=Warning -org.eclipse.cdt.codan.internal.checkers.UnusedVariableDeclarationProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true},macro\=>true,exceptions\=>("@(\#)","$Id")} -org.eclipse.cdt.codan.internal.checkers.VariableResolutionProblem=Error -org.eclipse.cdt.codan.internal.checkers.VariableResolutionProblem.params={launchModes\=>{RUN_ON_FULL_BUILD\=>true,RUN_ON_INC_BUILD\=>true,RUN_ON_FILE_OPEN\=>false,RUN_ON_FILE_SAVE\=>false,RUN_AS_YOU_TYPE\=>true,RUN_ON_DEMAND\=>true}} -useParentScope=false diff --git a/.settings/org.eclipse.cdt.core.prefs b/.settings/org.eclipse.cdt.core.prefs deleted file mode 100644 index 80b8e65c79..0000000000 --- a/.settings/org.eclipse.cdt.core.prefs +++ /dev/null @@ -1,177 +0,0 @@ -eclipse.preferences.version=1 -indexer/indexAllFiles=true -indexer/indexAllHeaderVersions=false -indexer/indexAllVersionsSpecificHeaders= -indexer/indexOnOpen=false -indexer/indexUnusedHeadersWithAlternateLang=false -indexer/indexUnusedHeadersWithDefaultLang=true -indexer/indexerId=org.eclipse.cdt.core.fastIndexer -indexer/skipFilesLargerThanMB=8 -indexer/skipImplicitReferences=false -indexer/skipIncludedFilesLargerThanMB=16 -indexer/skipMacroReferences=false -indexer/skipReferences=false -indexer/skipTypeReferences=false -indexer/useHeuristicIncludeResolution=true -org.eclipse.cdt.core.formatter.alignment_for_arguments_in_method_invocation=16 -org.eclipse.cdt.core.formatter.alignment_for_assignment=16 -org.eclipse.cdt.core.formatter.alignment_for_base_clause_in_type_declaration=48 -org.eclipse.cdt.core.formatter.alignment_for_binary_expression=16 -org.eclipse.cdt.core.formatter.alignment_for_compact_if=0 -org.eclipse.cdt.core.formatter.alignment_for_conditional_expression=48 -org.eclipse.cdt.core.formatter.alignment_for_conditional_expression_chain=18 -org.eclipse.cdt.core.formatter.alignment_for_constructor_initializer_list=0 -org.eclipse.cdt.core.formatter.alignment_for_declarator_list=16 -org.eclipse.cdt.core.formatter.alignment_for_enumerator_list=48 -org.eclipse.cdt.core.formatter.alignment_for_expression_list=0 -org.eclipse.cdt.core.formatter.alignment_for_expressions_in_array_initializer=16 -org.eclipse.cdt.core.formatter.alignment_for_member_access=0 -org.eclipse.cdt.core.formatter.alignment_for_overloaded_left_shift_chain=16 -org.eclipse.cdt.core.formatter.alignment_for_parameters_in_method_declaration=48 -org.eclipse.cdt.core.formatter.alignment_for_throws_clause_in_method_declaration=48 -org.eclipse.cdt.core.formatter.brace_position_for_array_initializer=next_line -org.eclipse.cdt.core.formatter.brace_position_for_block=next_line -org.eclipse.cdt.core.formatter.brace_position_for_block_in_case=end_of_line -org.eclipse.cdt.core.formatter.brace_position_for_method_declaration=next_line -org.eclipse.cdt.core.formatter.brace_position_for_namespace_declaration=end_of_line -org.eclipse.cdt.core.formatter.brace_position_for_switch=end_of_line -org.eclipse.cdt.core.formatter.brace_position_for_type_declaration=next_line -org.eclipse.cdt.core.formatter.comment.min_distance_between_code_and_line_comment=1 -org.eclipse.cdt.core.formatter.comment.never_indent_line_comments_on_first_column=true -org.eclipse.cdt.core.formatter.comment.preserve_white_space_between_code_and_line_comments=true -org.eclipse.cdt.core.formatter.compact_else_if=true -org.eclipse.cdt.core.formatter.continuation_indentation=1 -org.eclipse.cdt.core.formatter.continuation_indentation_for_array_initializer=1 -org.eclipse.cdt.core.formatter.format_guardian_clause_on_one_line=false -org.eclipse.cdt.core.formatter.indent_access_specifier_compare_to_type_header=false -org.eclipse.cdt.core.formatter.indent_access_specifier_extra_spaces=0 -org.eclipse.cdt.core.formatter.indent_body_declarations_compare_to_access_specifier=true -org.eclipse.cdt.core.formatter.indent_body_declarations_compare_to_namespace_header=false -org.eclipse.cdt.core.formatter.indent_breaks_compare_to_cases=true -org.eclipse.cdt.core.formatter.indent_declaration_compare_to_template_header=false -org.eclipse.cdt.core.formatter.indent_empty_lines=false -org.eclipse.cdt.core.formatter.indent_statements_compare_to_block=true -org.eclipse.cdt.core.formatter.indent_statements_compare_to_body=true -org.eclipse.cdt.core.formatter.indent_switchstatements_compare_to_cases=true -org.eclipse.cdt.core.formatter.indent_switchstatements_compare_to_switch=false -org.eclipse.cdt.core.formatter.indentation.size=4 -org.eclipse.cdt.core.formatter.insert_new_line_after_opening_brace_in_array_initializer=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_after_template_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_at_end_of_file_if_missing=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_before_catch_in_try_statement=insert -org.eclipse.cdt.core.formatter.insert_new_line_before_closing_brace_in_array_initializer=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_before_colon_in_constructor_initializer_list=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_before_else_in_if_statement=insert -org.eclipse.cdt.core.formatter.insert_new_line_before_identifier_in_function_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_before_while_in_do_statement=do not insert -org.eclipse.cdt.core.formatter.insert_new_line_in_empty_block=insert -org.eclipse.cdt.core.formatter.insert_space_after_assignment_operator=insert -org.eclipse.cdt.core.formatter.insert_space_after_binary_operator=insert -org.eclipse.cdt.core.formatter.insert_space_after_closing_angle_bracket_in_template_arguments=insert -org.eclipse.cdt.core.formatter.insert_space_after_closing_angle_bracket_in_template_parameters=insert -org.eclipse.cdt.core.formatter.insert_space_after_closing_brace_in_block=insert -org.eclipse.cdt.core.formatter.insert_space_after_closing_paren_in_cast=insert -org.eclipse.cdt.core.formatter.insert_space_after_colon_in_base_clause=insert -org.eclipse.cdt.core.formatter.insert_space_after_colon_in_case=insert -org.eclipse.cdt.core.formatter.insert_space_after_colon_in_conditional=insert -org.eclipse.cdt.core.formatter.insert_space_after_colon_in_labeled_statement=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_array_initializer=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_base_types=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_declarator_list=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_enum_declarations=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_expression_list=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_declaration_parameters=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_declaration_throws=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_invocation_arguments=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_template_arguments=insert -org.eclipse.cdt.core.formatter.insert_space_after_comma_in_template_parameters=insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_angle_bracket_in_template_arguments=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_angle_bracket_in_template_parameters=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_brace_in_array_initializer=insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_bracket=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_cast=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_catch=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_exception_specification=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_for=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_if=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_method_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_method_invocation=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_parenthesized_expression=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_switch=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_while=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_postfix_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_prefix_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_after_question_in_conditional=insert -org.eclipse.cdt.core.formatter.insert_space_after_semicolon_in_for=insert -org.eclipse.cdt.core.formatter.insert_space_after_unary_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_assignment_operator=insert -org.eclipse.cdt.core.formatter.insert_space_before_binary_operator=insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_angle_bracket_in_template_arguments=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_angle_bracket_in_template_parameters=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_brace_in_array_initializer=insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_bracket=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_cast=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_catch=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_exception_specification=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_for=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_if=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_method_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_method_invocation=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_parenthesized_expression=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_switch=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_while=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_colon_in_base_clause=insert -org.eclipse.cdt.core.formatter.insert_space_before_colon_in_case=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_colon_in_conditional=insert -org.eclipse.cdt.core.formatter.insert_space_before_colon_in_default=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_colon_in_labeled_statement=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_array_initializer=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_base_types=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_declarator_list=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_enum_declarations=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_expression_list=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_declaration_parameters=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_declaration_throws=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_invocation_arguments=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_template_arguments=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_comma_in_template_parameters=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_angle_bracket_in_template_arguments=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_angle_bracket_in_template_parameters=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_array_initializer=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_block=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_method_declaration=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_namespace_declaration=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_switch=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_type_declaration=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_bracket=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_catch=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_exception_specification=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_for=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_if=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_method_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_method_invocation=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_parenthesized_expression=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_switch=insert -org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_while=insert -org.eclipse.cdt.core.formatter.insert_space_before_postfix_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_prefix_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_question_in_conditional=insert -org.eclipse.cdt.core.formatter.insert_space_before_semicolon=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_semicolon_in_for=do not insert -org.eclipse.cdt.core.formatter.insert_space_before_unary_operator=do not insert -org.eclipse.cdt.core.formatter.insert_space_between_empty_braces_in_array_initializer=do not insert -org.eclipse.cdt.core.formatter.insert_space_between_empty_brackets=do not insert -org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_exception_specification=do not insert -org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_method_declaration=do not insert -org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_method_invocation=do not insert -org.eclipse.cdt.core.formatter.join_wrapped_lines=true -org.eclipse.cdt.core.formatter.keep_else_statement_on_same_line=false -org.eclipse.cdt.core.formatter.keep_empty_array_initializer_on_one_line=false -org.eclipse.cdt.core.formatter.keep_imple_if_on_one_line=true -org.eclipse.cdt.core.formatter.keep_then_statement_on_same_line=false -org.eclipse.cdt.core.formatter.lineSplit=80 -org.eclipse.cdt.core.formatter.number_of_empty_lines_to_preserve=1 -org.eclipse.cdt.core.formatter.put_empty_statement_on_new_line=true -org.eclipse.cdt.core.formatter.tabulation.char=space -org.eclipse.cdt.core.formatter.tabulation.size=4 -org.eclipse.cdt.core.formatter.use_tabs_only_for_leading_indentations=false diff --git a/.settings/org.eclipse.cdt.ui.prefs b/.settings/org.eclipse.cdt.ui.prefs deleted file mode 100644 index ca73f82def..0000000000 --- a/.settings/org.eclipse.cdt.ui.prefs +++ /dev/null @@ -1,3 +0,0 @@ -eclipse.preferences.version=1 -formatter_profile=_B40C -formatter_settings_version=1 diff --git a/.settings/org.eclipse.core.runtime.prefs b/.settings/org.eclipse.core.runtime.prefs deleted file mode 100644 index 2e6330e751..0000000000 --- a/.settings/org.eclipse.core.runtime.prefs +++ /dev/null @@ -1,4 +0,0 @@ -content-types/enabled=true -content-types/org.eclipse.cdt.core.cxxHeader/file-extensions=cuh -content-types/org.eclipse.cdt.core.cxxSource/file-extensions=cu -eclipse.preferences.version=1 From b61c65b3e6fb1980b84882dd4923ddc315fdee0c Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 14:04:55 -0500 Subject: [PATCH 06/11] Fix some files that used CRLF dos line endings. --- LICENSE.TXT | 48 ++--- test/half.h | 612 ++++++++++++++++++++++++++-------------------------- 2 files changed, 330 insertions(+), 330 deletions(-) diff --git a/LICENSE.TXT b/LICENSE.TXT index a678e64f8c..6aeea8da6f 100644 --- a/LICENSE.TXT +++ b/LICENSE.TXT @@ -1,24 +1,24 @@ -Copyright (c) 2010-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. \ No newline at end of file +Copyright (c) 2010-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. diff --git a/test/half.h b/test/half.h index c5255d61fe..842f9f730d 100644 --- a/test/half.h +++ b/test/half.h @@ -1,43 +1,43 @@ -/****************************************************************************** - * Copyright (c) 2011, Duane Merrill. All rights reserved. - * Copyright (c) 2011-2019, 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. - * - ******************************************************************************/ - -#pragma once - -/** - * \file - * Utilities for interacting with the opaque CUDA __half type - */ - -#include -#include -#include - -#include +/****************************************************************************** + * Copyright (c) 2011, Duane Merrill. All rights reserved. + * Copyright (c) 2011-2019, 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. + * + ******************************************************************************/ + +#pragma once + +/** + * \file + * Utilities for interacting with the opaque CUDA __half type + */ + +#include +#include +#include + +#include #ifdef __GNUC__ // There's a ton of type-punning going on in this file. @@ -45,272 +45,272 @@ #pragma GCC diagnostic ignored "-Wstrict-aliasing" #endif - -/****************************************************************************** - * half_t - ******************************************************************************/ - -/** - * Host-based fp16 data type compatible and convertible with __half - */ -struct half_t -{ - uint16_t __x; - - /// Constructor from __half - __host__ __device__ __forceinline__ - half_t(const __half &other) - { - __x = reinterpret_cast(other); - } - - /// Constructor from integer - __host__ __device__ __forceinline__ - half_t(int a) - { - *this = half_t(float(a)); - } - - /// Default constructor - __host__ __device__ __forceinline__ - half_t() : __x(0) - {} - - /// Constructor from float - __host__ __device__ __forceinline__ - half_t(float a) - { - // Stolen from Norbert Juffa - uint32_t ia = *reinterpret_cast(&a); - uint16_t ir; - - ir = (ia >> 16) & 0x8000; - - if ((ia & 0x7f800000) == 0x7f800000) - { - if ((ia & 0x7fffffff) == 0x7f800000) - { - ir |= 0x7c00; /* infinity */ - } - else - { - ir = 0x7fff; /* canonical NaN */ - } - } - else if ((ia & 0x7f800000) >= 0x33000000) - { - int32_t shift = (int32_t) ((ia >> 23) & 0xff) - 127; - if (shift > 15) - { - ir |= 0x7c00; /* infinity */ - } - else - { - ia = (ia & 0x007fffff) | 0x00800000; /* extract mantissa */ - if (shift < -14) - { /* denormal */ - ir |= ia >> (-1 - shift); - ia = ia << (32 - (-1 - shift)); - } - else - { /* normal */ - ir |= ia >> (24 - 11); - ia = ia << (32 - (24 - 11)); - ir = ir + ((14 + shift) << 10); - } - /* IEEE-754 round to nearest of even */ - if ((ia > 0x80000000) || ((ia == 0x80000000) && (ir & 1))) - { - ir++; - } - } - } - - this->__x = ir; - } - - /// Cast to __half - __host__ __device__ __forceinline__ - operator __half() const - { - return reinterpret_cast(__x); - } - - /// Cast to float - __host__ __device__ __forceinline__ - operator float() const - { - // Stolen from Andrew Kerr - - int sign = ((this->__x >> 15) & 1); - int exp = ((this->__x >> 10) & 0x1f); - int mantissa = (this->__x & 0x3ff); - uint32_t f = 0; - - if (exp > 0 && exp < 31) - { - // normal - exp += 112; - f = (sign << 31) | (exp << 23) | (mantissa << 13); - } - else if (exp == 0) - { - if (mantissa) - { - // subnormal - exp += 113; - while ((mantissa & (1 << 10)) == 0) - { - mantissa <<= 1; - exp--; - } - mantissa &= 0x3ff; - f = (sign << 31) | (exp << 23) | (mantissa << 13); - } - else if (sign) - { - f = 0x80000000; // negative zero - } - else - { - f = 0x0; // zero - } - } - else if (exp == 31) - { - if (mantissa) - { - f = 0x7fffffff; // not a number - } - else - { - f = (0xff << 23) | (sign << 31); // inf - } - } - return *reinterpret_cast(&f); - } - - - /// Get raw storage - __host__ __device__ __forceinline__ - uint16_t raw() - { - return this->__x; - } - - /// Equality - __host__ __device__ __forceinline__ - bool operator ==(const half_t &other) - { - return (this->__x == other.__x); - } - - /// Inequality - __host__ __device__ __forceinline__ - bool operator !=(const half_t &other) - { - return (this->__x != other.__x); - } - - /// Assignment by sum - __host__ __device__ __forceinline__ - half_t& operator +=(const half_t &rhs) - { - *this = half_t(float(*this) + float(rhs)); - return *this; - } - - /// Multiply - __host__ __device__ __forceinline__ - half_t operator*(const half_t &other) - { - return half_t(float(*this) * float(other)); - } - - /// Add - __host__ __device__ __forceinline__ - half_t operator+(const half_t &other) - { - return half_t(float(*this) + float(other)); - } - - /// Less-than - __host__ __device__ __forceinline__ - bool operator<(const half_t &other) const - { - return float(*this) < float(other); - } - - /// Less-than-equal - __host__ __device__ __forceinline__ - bool operator<=(const half_t &other) const - { - return float(*this) <= float(other); - } - - /// Greater-than - __host__ __device__ __forceinline__ - bool operator>(const half_t &other) const - { - return float(*this) > float(other); - } - - /// Greater-than-equal - __host__ __device__ __forceinline__ - bool operator>=(const half_t &other) const - { - return float(*this) >= float(other); - } - - /// numeric_traits::max - __host__ __device__ __forceinline__ - static half_t max() { - uint16_t max_word = 0x7BFF; - return reinterpret_cast(max_word); - } - - /// numeric_traits::lowest - __host__ __device__ __forceinline__ - static half_t lowest() { - uint16_t lowest_word = 0xFBFF; - return reinterpret_cast(lowest_word); - } -}; - - -/****************************************************************************** - * I/O stream overloads - ******************************************************************************/ - -/// Insert formatted \p half_t into the output stream -std::ostream& operator<<(std::ostream &out, const half_t &x) -{ - out << (float)x; - return out; -} - - -/// Insert formatted \p __half into the output stream -std::ostream& operator<<(std::ostream &out, const __half &x) -{ - return out << half_t(x); -} - - -/****************************************************************************** - * Traits overloads - ******************************************************************************/ - -template <> -struct cub::FpLimits -{ - static __host__ __device__ __forceinline__ half_t Max() { return half_t::max(); } - - static __host__ __device__ __forceinline__ half_t Lowest() { return half_t::lowest(); } -}; - -template <> struct cub::NumericTraits : cub::BaseTraits {}; - + +/****************************************************************************** + * half_t + ******************************************************************************/ + +/** + * Host-based fp16 data type compatible and convertible with __half + */ +struct half_t +{ + uint16_t __x; + + /// Constructor from __half + __host__ __device__ __forceinline__ + half_t(const __half &other) + { + __x = reinterpret_cast(other); + } + + /// Constructor from integer + __host__ __device__ __forceinline__ + half_t(int a) + { + *this = half_t(float(a)); + } + + /// Default constructor + __host__ __device__ __forceinline__ + half_t() : __x(0) + {} + + /// Constructor from float + __host__ __device__ __forceinline__ + half_t(float a) + { + // Stolen from Norbert Juffa + uint32_t ia = *reinterpret_cast(&a); + uint16_t ir; + + ir = (ia >> 16) & 0x8000; + + if ((ia & 0x7f800000) == 0x7f800000) + { + if ((ia & 0x7fffffff) == 0x7f800000) + { + ir |= 0x7c00; /* infinity */ + } + else + { + ir = 0x7fff; /* canonical NaN */ + } + } + else if ((ia & 0x7f800000) >= 0x33000000) + { + int32_t shift = (int32_t) ((ia >> 23) & 0xff) - 127; + if (shift > 15) + { + ir |= 0x7c00; /* infinity */ + } + else + { + ia = (ia & 0x007fffff) | 0x00800000; /* extract mantissa */ + if (shift < -14) + { /* denormal */ + ir |= ia >> (-1 - shift); + ia = ia << (32 - (-1 - shift)); + } + else + { /* normal */ + ir |= ia >> (24 - 11); + ia = ia << (32 - (24 - 11)); + ir = ir + ((14 + shift) << 10); + } + /* IEEE-754 round to nearest of even */ + if ((ia > 0x80000000) || ((ia == 0x80000000) && (ir & 1))) + { + ir++; + } + } + } + + this->__x = ir; + } + + /// Cast to __half + __host__ __device__ __forceinline__ + operator __half() const + { + return reinterpret_cast(__x); + } + + /// Cast to float + __host__ __device__ __forceinline__ + operator float() const + { + // Stolen from Andrew Kerr + + int sign = ((this->__x >> 15) & 1); + int exp = ((this->__x >> 10) & 0x1f); + int mantissa = (this->__x & 0x3ff); + uint32_t f = 0; + + if (exp > 0 && exp < 31) + { + // normal + exp += 112; + f = (sign << 31) | (exp << 23) | (mantissa << 13); + } + else if (exp == 0) + { + if (mantissa) + { + // subnormal + exp += 113; + while ((mantissa & (1 << 10)) == 0) + { + mantissa <<= 1; + exp--; + } + mantissa &= 0x3ff; + f = (sign << 31) | (exp << 23) | (mantissa << 13); + } + else if (sign) + { + f = 0x80000000; // negative zero + } + else + { + f = 0x0; // zero + } + } + else if (exp == 31) + { + if (mantissa) + { + f = 0x7fffffff; // not a number + } + else + { + f = (0xff << 23) | (sign << 31); // inf + } + } + return *reinterpret_cast(&f); + } + + + /// Get raw storage + __host__ __device__ __forceinline__ + uint16_t raw() + { + return this->__x; + } + + /// Equality + __host__ __device__ __forceinline__ + bool operator ==(const half_t &other) + { + return (this->__x == other.__x); + } + + /// Inequality + __host__ __device__ __forceinline__ + bool operator !=(const half_t &other) + { + return (this->__x != other.__x); + } + + /// Assignment by sum + __host__ __device__ __forceinline__ + half_t& operator +=(const half_t &rhs) + { + *this = half_t(float(*this) + float(rhs)); + return *this; + } + + /// Multiply + __host__ __device__ __forceinline__ + half_t operator*(const half_t &other) + { + return half_t(float(*this) * float(other)); + } + + /// Add + __host__ __device__ __forceinline__ + half_t operator+(const half_t &other) + { + return half_t(float(*this) + float(other)); + } + + /// Less-than + __host__ __device__ __forceinline__ + bool operator<(const half_t &other) const + { + return float(*this) < float(other); + } + + /// Less-than-equal + __host__ __device__ __forceinline__ + bool operator<=(const half_t &other) const + { + return float(*this) <= float(other); + } + + /// Greater-than + __host__ __device__ __forceinline__ + bool operator>(const half_t &other) const + { + return float(*this) > float(other); + } + + /// Greater-than-equal + __host__ __device__ __forceinline__ + bool operator>=(const half_t &other) const + { + return float(*this) >= float(other); + } + + /// numeric_traits::max + __host__ __device__ __forceinline__ + static half_t max() { + uint16_t max_word = 0x7BFF; + return reinterpret_cast(max_word); + } + + /// numeric_traits::lowest + __host__ __device__ __forceinline__ + static half_t lowest() { + uint16_t lowest_word = 0xFBFF; + return reinterpret_cast(lowest_word); + } +}; + + +/****************************************************************************** + * I/O stream overloads + ******************************************************************************/ + +/// Insert formatted \p half_t into the output stream +std::ostream& operator<<(std::ostream &out, const half_t &x) +{ + out << (float)x; + return out; +} + + +/// Insert formatted \p __half into the output stream +std::ostream& operator<<(std::ostream &out, const __half &x) +{ + return out << half_t(x); +} + + +/****************************************************************************** + * Traits overloads + ******************************************************************************/ + +template <> +struct cub::FpLimits +{ + static __host__ __device__ __forceinline__ half_t Max() { return half_t::max(); } + + static __host__ __device__ __forceinline__ half_t Lowest() { return half_t::lowest(); } +}; + +template <> struct cub::NumericTraits : cub::BaseTraits {}; + #ifdef __GNUC__ #pragma GCC diagnostic pop From 303ea5a28830bcb8b0fc20c53880cdbbb26b7663 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 14:31:59 -0500 Subject: [PATCH 07/11] Add .git-blame-ignore-revs file. See the file for more info. This will let us exclude formatting changes (clang-format, line endings, etc) from `git diff` and related tools. --- .git-blame-ignore-revs | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) create mode 100644 .git-blame-ignore-revs diff --git a/.git-blame-ignore-revs b/.git-blame-ignore-revs new file mode 100644 index 0000000000..68469e1f1f --- /dev/null +++ b/.git-blame-ignore-revs @@ -0,0 +1,16 @@ +# Exclude these commits from git-blame and similar tools. +# +# To use this file, run the following command from the repo root: +# +# ``` +# $ git config blame.ignoreRevsFile .git-blame-ignore-revs +# ``` +# +# Include a brief comment with each commit added, for example: +# +# ``` +# d92d9f8baac5ec48a8f8718dd69f415a45efe372 # Initial clang-format +# ``` +# +# Only add commits that are pure formatting changes (e.g. +# clang-format version changes, etc). From c875245650b9d3183f50e9293247fc91e6778ede Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 14:36:20 -0500 Subject: [PATCH 08/11] Fix `class-memaccess` warning on GCC 8. The testing `half_t` struct had a non-trivial default constructor and thus a non-trivial type, which caused warnings to be emitted when used with `memcpy`. --- test/half.h | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/test/half.h b/test/half.h index 842f9f730d..33ec607758 100644 --- a/test/half.h +++ b/test/half.h @@ -72,9 +72,7 @@ struct half_t } /// Default constructor - __host__ __device__ __forceinline__ - half_t() : __x(0) - {} + half_t() = default; /// Constructor from float __host__ __device__ __forceinline__ From 3f508ca8bcbb2cebda3caef075cc5b4d1ca39fb8 Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 24 Nov 2020 14:45:51 -0500 Subject: [PATCH 09/11] Fix `deprecated-copy` warning on GCC 9. Removing the copy-assignment operator allows the copy-constructor to be implicitly defined. --- cub/iterator/discard_output_iterator.cuh | 6 ------ 1 file changed, 6 deletions(-) diff --git a/cub/iterator/discard_output_iterator.cuh b/cub/iterator/discard_output_iterator.cuh index e665c784e9..2db03219ae 100644 --- a/cub/iterator/discard_output_iterator.cuh +++ b/cub/iterator/discard_output_iterator.cuh @@ -177,12 +177,6 @@ public: return; } - /// Assignment to self (no-op) - __host__ __device__ __forceinline__ void operator=(self_type const& other) - { - offset = other.offset; - } - /// Assignment to anything else (no-op) template __host__ __device__ __forceinline__ void operator=(T const&) From cb50e4582fdd9005fd10fc0b831288fe46a9890f Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Fri, 4 Dec 2020 13:45:20 -0500 Subject: [PATCH 10/11] Bump version to 1.12.0. --- cub/version.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cub/version.cuh b/cub/version.cuh index d01a4c3d6f..7157048750 100644 --- a/cub/version.cuh +++ b/cub/version.cuh @@ -43,7 +43,7 @@ * CUB_VERSION / 100 % 1000 is the minor version. * CUB_VERSION / 100000 is the major version. */ -#define CUB_VERSION 101100 +#define CUB_VERSION 101200 /*! \def CUB_MAJOR_VERSION * \brief The preprocessor macro \p CUB_MAJOR_VERSION encodes the From 0ccfb45cd9c17887ad8539c462fc1fea9025fcb1 Mon Sep 17 00:00:00 2001 From: Bryce Adelstein Lelbach aka wash Date: Fri, 30 Oct 2020 17:20:38 -0700 Subject: [PATCH 11/11] Remove TEST_ARCH, which is no longer needed. --- cmake/CubCudaConfig.cmake | 7 --- common.mk | 8 --- test/CMakeLists.txt | 7 --- test/Makefile | 101 ++++++++++++++++++-------------------- test/test_block_reduce.cu | 20 ++------ test/test_block_scan.cu | 14 +----- 6 files changed, 55 insertions(+), 102 deletions(-) diff --git a/cmake/CubCudaConfig.cmake b/cmake/CubCudaConfig.cmake index 1f3ec544ac..5db5a0bc1f 100644 --- a/cmake/CubCudaConfig.cmake +++ b/cmake/CubCudaConfig.cmake @@ -6,7 +6,6 @@ enable_language(CUDA) set(all_archs 35 37 50 52 53 60 61 62 70 72 75 80) set(arch_message "CUB: Enabled CUDA architectures:") -set(enabled_archs) # Thrust sets up the architecture flags in CMAKE_CUDA_FLAGS already. Just # reuse them if possible. After we transition to CMake 3.18 CUDA_ARCHITECTURE @@ -19,7 +18,6 @@ if (CUB_IN_THRUST) foreach (arch IN LISTS all_archs) if (THRUST_ENABLE_COMPUTE_${arch}) set(CUB_ENABLE_COMPUTE_${arch} True) - list(APPEND enabled_archs ${arch}) string(APPEND arch_message " sm_${arch}") else() set(CUB_ENABLE_COMPUTE_${arch} False) @@ -52,7 +50,6 @@ else() # NOT CUB_IN_THRUST ${option_init} ) if (CUB_ENABLE_COMPUTE_${arch}) - list(APPEND enabled_archs ${arch}) string(APPEND arch_flags " -gencode arch=compute_${arch},code=sm_${arch}") string(APPEND arch_message " sm_${arch}") endif() @@ -75,10 +72,6 @@ endif() message(STATUS ${arch_message}) -# Create a variable containing the minimal target arch for tests -list(SORT enabled_archs) -list(GET enabled_archs 0 CUB_MINIMAL_ENABLED_ARCH) - # # RDC options: # diff --git a/common.mk b/common.mk index 4010ed3098..92e16550c1 100644 --- a/common.mk +++ b/common.mk @@ -43,42 +43,34 @@ endif ifeq (700, $(findstring 700, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_70,code=\"sm_70,compute_70\" SM_DEF += -DSM700 - TEST_ARCH = 700 endif ifeq (620, $(findstring 620, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_62,code=\"sm_62,compute_62\" SM_DEF += -DSM620 - TEST_ARCH = 620 endif ifeq (610, $(findstring 610, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_61,code=\"sm_61,compute_61\" SM_DEF += -DSM610 - TEST_ARCH = 610 endif ifeq (600, $(findstring 600, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_60,code=\"sm_60,compute_60\" SM_DEF += -DSM600 - TEST_ARCH = 600 endif ifeq (520, $(findstring 520, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_52,code=\"sm_52,compute_52\" SM_DEF += -DSM520 - TEST_ARCH = 520 endif ifeq (370, $(findstring 370, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_37,code=\"sm_37,compute_37\" SM_DEF += -DSM370 - TEST_ARCH = 370 endif ifeq (350, $(findstring 350, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_35,code=\"sm_35,compute_35\" SM_DEF += -DSM350 - TEST_ARCH = 350 endif ifeq (300, $(findstring 300, $(SM_ARCH))) SM_TARGETS += -gencode=arch=compute_30,code=\"sm_30,compute_30\" SM_DEF += -DSM300 - TEST_ARCH = 300 endif diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 45a3c23b2c..23ae83af5c 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1,9 +1,3 @@ -add_subdirectory(cmake) - -# TODO investigate whether this is really needed: -math(EXPR CUB_TEST_ARCH "${CUB_MINIMAL_ENABLED_ARCH} * 10") -message(STATUS "CUB Test architecture (TEST_ARCH): ${CUB_TEST_ARCH}") - # Create meta targets that build all tests for a single configuration: foreach(cub_target IN LISTS CUB_TARGETS) cub_get_target_property(config_prefix ${cub_target} PREFIX) @@ -44,7 +38,6 @@ function(cub_add_test target_name_var test_name test_src cub_target) add_executable(${test_target} "${test_src}") target_link_libraries(${test_target} ${cub_target}) cub_clone_target_properties(${test_target} ${cub_target}) - target_compile_definitions(${test_target} PRIVATE TEST_ARCH=${CUB_TEST_ARCH}) target_include_directories(${test_target} PRIVATE "${CUB_SOURCE_DIR}/test") # Add to the active configuration's meta target diff --git a/test/Makefile b/test/Makefile index 958760a87c..c7e37d8c76 100644 --- a/test/Makefile +++ b/test/Makefile @@ -1,7 +1,7 @@ #/****************************************************************************** # * 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 @@ -12,7 +12,7 @@ # * * 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 @@ -35,8 +35,8 @@ # #------------------------------------------------------------------------------- -include ../common.mk - +include ../common.mk + #------------------------------------------------------------------------------- # Commandline Options #------------------------------------------------------------------------------- @@ -48,16 +48,16 @@ ifeq ($(quickertest), 1) else ifeq ($(quicktest), 1) NVCCFLAGS += -DQUICK_TEST TEST_SUFFIX = quick -else +else TEST_SUFFIX = thorough - NPPI = + NPPI = endif -# CUDA memcheck (enabled by default) +# CUDA memcheck (enabled by default) ifeq ($(memcheck), 0) - MEMCHECK = -else + MEMCHECK = +else MEMCHECK = cuda-memcheck endif @@ -67,14 +67,11 @@ endif #------------------------------------------------------------------------------- # Includes -INC += -I$(CUB_DIR) -I$(CUB_DIR)test +INC += -I$(CUB_DIR) -I$(CUB_DIR)test # Suffix to append to each binary SUFFIX = $(BIN_SUFFIX)_$(TEST_SUFFIX) -# Define test arch -DEFINES += -DTEST_ARCH=$(TEST_ARCH) - #------------------------------------------------------------------------------- # Dependency Lists @@ -88,20 +85,20 @@ DEPS = $(CUB_DEPS) \ $(CUB_DIR)test/mersenne.h \ BLOCK_REDUCE = test_block_reduce_raking \ - test_block_reduce_warp_reductions + test_block_reduce_warp_reductions BLOCK_SCAN = test_block_scan_raking \ test_block_scan_raking_memoize \ - test_block_scan_warp_scans + test_block_scan_warp_scans BLOCK_RADIX_SORT = test_block_radix_sort_keys \ - test_block_radix_sort_pairs + test_block_radix_sort_pairs DEVICE_RADIX_SORT = test_device_radix_sort \ - test_device_radix_sort_segmented - + test_device_radix_sort_segmented + ALL = link \ test_iterator \ test_allocator \ @@ -119,11 +116,11 @@ ALL = link \ test_device_reduce_by_key\ test_device_run_length_encode\ test_device_select_unique \ - test_device_select_if - + test_device_select_if + # test_grid_barrier \ fails on sm110 # test_device_seg_reduce - + #------------------------------------------------------------------------------- @@ -138,7 +135,7 @@ default: #------------------------------------------------------------------------------- clean : - rm -f bin/*$(CPU_ARCH_SUFFIX)* + rm -f bin/*$(CPU_ARCH_SUFFIX)* rm -f *.i* *.cubin *.cu.c *.cudafe* *.fatbin.c *.ptx *.hash *.cu.cpp *.o @@ -153,19 +150,19 @@ all : $(ALL) # make run #------------------------------------------------------------------------------- -run : +run : for i in $(ALL); do $(MEMCHECK) ./bin/$${i}_$(SUFFIX) --device=$(device) || exit 1; done -run_block_reduce : +run_block_reduce : for i in $(BLOCK_REDUCE); do $(MEMCHECK) ./bin/$${i}_$(SUFFIX) --device=$(device) || exit 1; done -run_block_scan : +run_block_scan : for i in $(BLOCK_SCAN); do $(MEMCHECK) ./bin/$${i}_$(SUFFIX) --device=$(device) || exit 1; done -run_block_radix_sort : +run_block_radix_sort : for i in $(BLOCK_RADIX_SORT); do $(MEMCHECK) ./bin/$${i}_$(SUFFIX) --device=$(device) || exit 1; done -run_device_radix_sort : +run_device_radix_sort : for i in $(DEVICE_RADIX_SORT); do $(MEMCHECK) ./bin/$${i}_$(SUFFIX) --device=$(device) || exit 1; done @@ -183,7 +180,7 @@ bin/link_$(SUFFIX) : link_a.cu link_b.cu link_main.cpp $(DEPS) #------------------------------------------------------------------------------- -# make test_iterator +# make test_iterator #------------------------------------------------------------------------------- test_iterator: bin/test_iterator_$(SUFFIX) @@ -194,7 +191,7 @@ bin/test_iterator_$(SUFFIX) : test_iterator.cu $(DEPS) #------------------------------------------------------------------------------- -# make test_allocator +# make test_allocator #------------------------------------------------------------------------------- test_allocator: bin/test_allocator_$(SUFFIX) @@ -202,39 +199,39 @@ test_allocator: bin/test_allocator_$(SUFFIX) bin/test_allocator_$(SUFFIX) : test_allocator.cu $(DEPS) mkdir -p bin $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_allocator_$(SUFFIX) test_allocator.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 - - + + #------------------------------------------------------------------------------- -# make test_grid_barrier +# make test_grid_barrier #------------------------------------------------------------------------------- test_grid_barrier: bin/test_grid_barrier_$(SUFFIX) bin/test_grid_barrier_$(SUFFIX) : test_grid_barrier.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_grid_barrier_$(SUFFIX) test_grid_barrier.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 - + $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_grid_barrier_$(SUFFIX) test_grid_barrier.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + #------------------------------------------------------------------------------- -# make test_warp_scan +# make test_warp_scan #------------------------------------------------------------------------------- test_warp_scan: bin/test_warp_scan_$(SUFFIX) bin/test_warp_scan_$(SUFFIX) : test_warp_scan.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_warp_scan_$(SUFFIX) test_warp_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_warp_scan_$(SUFFIX) test_warp_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- -# make test_warp_reduce +# make test_warp_reduce #------------------------------------------------------------------------------- test_warp_reduce: bin/test_warp_reduce_$(SUFFIX) bin/test_warp_reduce_$(SUFFIX) : test_warp_reduce.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_warp_reduce_$(SUFFIX) test_warp_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_warp_reduce_$(SUFFIX) test_warp_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- @@ -245,22 +242,22 @@ test_block_reduce_raking: bin/test_block_reduce_raking_$(SUFFIX) bin/test_block_reduce_raking_$(SUFFIX) : test_block_reduce.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) -DTEST_RAKING $(SM_TARGETS) -o bin/test_block_reduce_raking_$(SUFFIX) test_block_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) -DTEST_RAKING $(SM_TARGETS) -o bin/test_block_reduce_raking_$(SUFFIX) test_block_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- -# make test_block_reduce_warp_reductions +# make test_block_reduce_warp_reductions #------------------------------------------------------------------------------- test_block_reduce_warp_reductions: bin/test_block_reduce_warp_reductions_$(SUFFIX) bin/test_block_reduce_warp_reductions_$(SUFFIX) : test_block_reduce.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) -DTEST_WARP_REDUCTIONS $(SM_TARGETS) -o bin/test_block_reduce_warp_reductions_$(SUFFIX) test_block_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) -DTEST_WARP_REDUCTIONS $(SM_TARGETS) -o bin/test_block_reduce_warp_reductions_$(SUFFIX) test_block_reduce.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- -# make test_block_reduce +# make test_block_reduce #------------------------------------------------------------------------------- test_block_reduce: $(BLOCK_REDUCE) @@ -274,7 +271,7 @@ test_block_scan_raking: bin/test_block_scan_raking_$(SUFFIX) bin/test_block_scan_raking_$(SUFFIX) : test_block_scan.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) -DTEST_RAKING $(SM_TARGETS) -o bin/test_block_scan_raking_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) -DTEST_RAKING $(SM_TARGETS) -o bin/test_block_scan_raking_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- @@ -285,7 +282,7 @@ test_block_scan_raking_memoize: bin/test_block_scan_raking_memoize_$(SUFFIX) bin/test_block_scan_raking_memoize_$(SUFFIX) : test_block_scan.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) -DTEST_RAKING_MEMOIZE $(SM_TARGETS) -o bin/test_block_scan_raking_memoize_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) -DTEST_RAKING_MEMOIZE $(SM_TARGETS) -o bin/test_block_scan_raking_memoize_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- @@ -296,18 +293,18 @@ test_block_scan_warp_scans: bin/test_block_scan_warp_scans_$(SUFFIX) bin/test_block_scan_warp_scans_$(SUFFIX) : test_block_scan.cu $(DEPS) mkdir -p bin - $(NVCC) $(DEFINES) -DTEST_WARP_SCANS $(SM_TARGETS) -o bin/test_block_scan_warp_scans_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 + $(NVCC) $(DEFINES) -DTEST_WARP_SCANS $(SM_TARGETS) -o bin/test_block_scan_warp_scans_$(SUFFIX) test_block_scan.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- -# make test_block_scan +# make test_block_scan #------------------------------------------------------------------------------- test_block_scan: $(BLOCK_SCAN) #------------------------------------------------------------------------------- -# make test_block_load_store +# make test_block_load_store #------------------------------------------------------------------------------- test_block_load_store: bin/test_block_load_store_$(SUFFIX) @@ -315,10 +312,10 @@ test_block_load_store: bin/test_block_load_store_$(SUFFIX) bin/test_block_load_store_$(SUFFIX) : test_block_load_store.cu $(DEPS) mkdir -p bin $(NVCC) $(DEFINES) $(SM_TARGETS) -o bin/test_block_load_store_$(SUFFIX) test_block_load_store.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 - - + + #------------------------------------------------------------------------------- -# make test_block_radix_sort_keys +# make test_block_radix_sort_keys #------------------------------------------------------------------------------- test_block_radix_sort_keys: bin/test_block_radix_sort_keys_$(SUFFIX) @@ -328,7 +325,7 @@ bin/test_block_radix_sort_keys_$(SUFFIX) : test_block_radix_sort.cu $(DEPS) $(NVCC) $(DEFINES) -DTEST_KEYS_ONLY $(SM_TARGETS) -o bin/test_block_radix_sort_keys_$(SUFFIX) test_block_radix_sort.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBS) -O3 #------------------------------------------------------------------------------- -# make test_block_radix_sort_pairs +# make test_block_radix_sort_pairs #------------------------------------------------------------------------------- test_block_radix_sort_pairs: bin/test_block_radix_sort_pairs_$(SUFFIX) @@ -346,7 +343,7 @@ test_block_radix_sort : $(BLOCK_RADIX_SORT) #------------------------------------------------------------------------------- -# make test_block_histogram +# make test_block_histogram #------------------------------------------------------------------------------- test_block_histogram: bin/test_block_histogram_$(SUFFIX) diff --git a/test/test_block_reduce.cu b/test/test_block_reduce.cu index 2b439b4061..91395c4321 100644 --- a/test/test_block_reduce.cu +++ b/test/test_block_reduce.cu @@ -176,7 +176,7 @@ __global__ void FullTileReduceKernel( { // TestBarrier between thread block reductions __syncthreads(); - + // Load tile of data LoadDirectBlocked(linear_tid, d_in + block_offset, data); block_offset += TILE_SIZE; @@ -429,17 +429,12 @@ void TestFullTile( ReductionOp reduction_op) { // Check size of smem storage for the target arch to make sure it will fit - typedef BlockReduce BlockReduceT; + typedef BlockReduce BlockReduceT; - enum + enum { -#if defined(SM100) || defined(SM110) || defined(SM130) - sufficient_smem = (sizeof(typename BlockReduceT::TempStorage) <= 16 * 1024), - sufficient_threads = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 512), -#else sufficient_smem = (sizeof(typename BlockReduceT::TempStorage) <= 48 * 1024), sufficient_threads = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 1024), -#endif }; TestFullTile(gen_mode, tiles, reduction_op, Int2Type()); @@ -613,17 +608,12 @@ void TestPartialTile( ReductionOp reduction_op) { // Check size of smem storage for the target arch to make sure it will fit - typedef BlockReduce BlockReduceT; + typedef BlockReduce BlockReduceT; - enum + enum { -#if defined(SM100) || defined(SM110) || defined(SM130) - sufficient_smem = sizeof(typename BlockReduceT::TempStorage) <= 16 * 1024, - sufficient_threads = (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 512, -#else sufficient_smem = sizeof(typename BlockReduceT::TempStorage) <= 48 * 1024, sufficient_threads = (BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 1024, -#endif }; TestPartialTile(gen_mode, num_items, reduction_op, Int2Type()); diff --git a/test/test_block_scan.cu b/test/test_block_scan.cu index d3c340d2b5..ef28cf026e 100644 --- a/test/test_block_scan.cu +++ b/test/test_block_scan.cu @@ -681,21 +681,9 @@ void Test( enum { -#if defined(SM100) || defined(SM110) || defined(SM130) - sufficient_smem = (sizeof(typename BlockScanT::TempStorage) <= 16 * 1024), - sufficient_threads = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 512), -#else sufficient_smem = (sizeof(typename BlockScanT::TempStorage) <= 16 * 1024), sufficient_threads = ((BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z) <= 1024), -#endif - -#if defined(_WIN32) || defined(_WIN64) - // Accommodate ptxas crash bug (access violation) on Windows - special_skip = ((TEST_ARCH <= 130) && (Equals::VALUE) && (BLOCK_DIM_Z > 1)), -#else - special_skip = false, -#endif - sufficient_resources = (sufficient_smem && sufficient_threads && !special_skip), + sufficient_resources = (sufficient_smem && sufficient_threads), }; Test(