Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix assertion in adjacent_difference #10

Draft
wants to merge 135 commits into
base: develop
Choose a base branch
from
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
135 commits
Select commit Hold shift + click to select a range
46354d2
Use builtin for atomic_fetch in the HIP backend
Rombur Mar 18, 2024
a2af4e0
Deprecate trailing Proxy template argument in Kokkos::Array
dalg24 Apr 11, 2024
de3a263
Merge pull request #6934 from dalg24/deprecate_kokkos_array_proxy_tem…
dalg24 Apr 15, 2024
f2d3780
Remove unnecessary header include
dalg24 Apr 15, 2024
8c7cc95
Merge pull request #6940 from dalg24/unused_limits_header_include_in_…
dalg24 Apr 16, 2024
a8115e5
Adding converting constructor in Kokkos::RandomAccessIterator (#6929)
yasahi-hpc Apr 16, 2024
f94e8d3
Prefer standard C++ feature testing to guard the C++20 requires expre…
dalg24 Apr 16, 2024
c9e21ce
Add `kokkos_swap(Array<T, N>)` sepcialization
dalg24 Apr 16, 2024
730d8d8
Deprecate specialization of Kokkos::pair for a single element
dalg24 Apr 17, 2024
906e8ce
Merge pull request #6942 from dalg24/fix_nightlies_cxx20_requires_exp…
dalg24 Apr 17, 2024
d914fe3
Fix deprecated warning from `Kokkos::Array` specialization (#6945)
dalg24 Apr 17, 2024
69c527a
[ci skip] Enable deprecated code and deprecated warnings in nightly CI
Rombur Apr 17, 2024
e7b486f
Serial: Use the provided execution space instance in TeamPolicy
masterleinad Apr 17, 2024
0859ab0
Fixed the link for P6601 (Threads backend change)
nliber Apr 17, 2024
d5fd512
Merge pull request #6947 from dalg24/deprecate_kokkos_pair_void_speci…
dalg24 Apr 17, 2024
04bc3d9
Merge pull request #6952 from nliber/changelog43
crtrott Apr 18, 2024
34d0db2
Add test
masterleinad Apr 18, 2024
44fde21
Use Kokkos::AUTO for OpenMPTarget
masterleinad Apr 18, 2024
8706b68
kokkos_swap(Array) member friend should not be templated on some othe…
dalg24 Apr 18, 2024
86f5988
Fix noexcept specification for kokkos_swap on zero-sized arrays
dalg24 Apr 18, 2024
cc60295
Merge pull request #6951 from masterleinad/fix_serial_space_team_policy
crtrott Apr 19, 2024
205fd15
Replace deprecated sycl::device_ptr/sycl::host_ptr
masterleinad Apr 22, 2024
5932685
Introduce alias based on feature macro
masterleinad Apr 22, 2024
a782773
Kokkos::Impl::SYCLTypes:: -> Kokkos::Impl::sycl_
masterleinad Apr 22, 2024
e2b7bb9
Merge pull request #6958 from masterleinad/sycl_replace_deprecated_us…
crtrott Apr 23, 2024
cf59f31
Merge pull request #6943 from dalg24/kokkos_swap_specialization_for_k…
crtrott Apr 23, 2024
ab3cae4
Fix wrong macro guards for deprecated Kokkos::pair<T1,void> specializ…
dalg24 Apr 24, 2024
fafe861
Fix support for Kokkos::Array of const-qualified element type
dalg24 Apr 24, 2024
2e82fdd
Merge pull request #6961 from dalg24/fixup_deprcated_guards_pair_void
dalg24 Apr 24, 2024
63eef46
Try to fix the CUDA 11.0 build
dalg24 Apr 24, 2024
ebb1cb3
Revert "Try to fix the CUDA 11.0 build"
dalg24 Apr 25, 2024
031f6d9
Alternate definition of Impl::is_nothrow_swappable_v for NVCC version…
dalg24 Apr 25, 2024
2391f17
Avoid introducing a 2nd definition of the Impl::swappable trait
dalg24 Apr 25, 2024
d434f87
Do not require OpenMP support for languages other than CXX
dalg24 Apr 25, 2024
19ca9ce
Update version
crtrott Apr 26, 2024
9686392
Add Linux Foundation notice and fix C++ standard
crtrott Apr 26, 2024
1864287
Merge pull request #6967 from crtrott/update-readme-kk-version
crtrott Apr 26, 2024
7e7709f
SYCL: Avoid deprecated floating-point number abs overloads (#6959)
masterleinad Apr 27, 2024
4ec8296
OpenMPTarget: Update loop order in MDRange (#6925)
rgayatri23 Apr 28, 2024
77ea52f
Threads: Don't silently allow m_instance to be a nullptr (#6969)
masterleinad May 1, 2024
4f416f3
Merge pull request #6965 from dalg24/cmake_openmp_cxx
dalg24 May 1, 2024
f699a2c
Fix enabling OpenMP with HIP and "compile as CMake language"
dalg24 May 1, 2024
2574b80
Fix OpenMP+CUDA when `Kokkos_ENABLE_COMPILE_AS_CMAKE_LANGUAGE` is `ON`
dalg24 May 1, 2024
27b3ced
Merge pull request #6949 from Rombur/nightly_deprecated
dalg24 May 1, 2024
15d13f2
Merge pull request #6882 from Rombur/hip_atomic_fetch
dalg24 May 1, 2024
ed4d254
Merge pull request #6972 from dalg24/fix_kokkos_compile_language_cuda…
crtrott May 1, 2024
dbd7f58
Merge pull request #6962 from dalg24/kokkos_array_const_qualified_ele…
crtrott May 1, 2024
ccd0126
Fix fedora CI builds with flang-new
masterleinad May 1, 2024
45a1404
Fix Copyright file
crtrott May 2, 2024
85610f4
Merge pull request #6984 from crtrott/Copyright
crtrott May 2, 2024
a75dc70
Merge pull request #6982 from masterleinad/fix_fedora
crtrott May 2, 2024
c6d8647
Also use is_nothrow_swappable workaround for Intel Classic Compilers …
masterleinad May 2, 2024
69567f3
Add thread-safety tests (#6938)
masterleinad May 3, 2024
9c79202
Fix deprecation warnings with GCC for pair<T1,void> comparison operators
dalg24 May 3, 2024
7b8e3a6
Fix TPL_LIBRARY_SUFFIXES for 32-bit build
masterleinad May 6, 2024
2826017
Avoid duplicated definition of KOKKOS_IMPL_32BIT
masterleinad May 6, 2024
ccadc7d
Disable failing parallel_scan_with_reducers test
masterleinad May 6, 2024
06e4c5b
Merge pull request #6989 from dalg24/deprecated_attribute_comparison_…
dalg24 May 6, 2024
e4cc686
Merge pull request #6990 from masterleinad/fix_32bit_tpl_library_path
dalg24 May 7, 2024
d61d75a
Fix a bug when using realloc on views of non-default constructible el…
aprokop May 8, 2024
50a862c
SYCL: Prepare Parallel* for Graphs (#6988)
masterleinad May 8, 2024
f5b3422
SYCL: Fix deprecation in custom parallel_for RangePolicy implementation
masterleinad May 8, 2024
37986fd
[ci skip] update changelog for 4.3.1 (#6995)
ndellingwood May 8, 2024
7cad3e7
OpenMPTarget: Use mutex lock for parallel scan.
May 8, 2024
a69e81a
Merge pull request #6998 from rgayatri23/ompt_scan_lock
dalg24 May 9, 2024
5a5306c
Merge pull request #6997 from masterleinad/sycl_fix_custom_parallel_f…
dalg24 May 9, 2024
00170ae
Remove cuSPARSE TPL
dalg24 May 9, 2024
506da18
Merge pull request #7002 from dalg24/rm_tpl_cusparse
dalg24 May 9, 2024
1d9d0df
SYCL: Print submission command queue property (#7004)
masterleinad May 10, 2024
cadab6c
Test DualView resize/realloc for types without default constructor
masterleinad May 7, 2024
df018d9
Suppress deprecated warnings via pragma push/pop in the tests (#6999)
dalg24 May 13, 2024
da8be22
This PR changes the default execution behavior of the parallel_for(te…
seyonglee May 15, 2024
835dbf5
Merge pull request #7012 from seyonglee/openacc_default_async_val_for…
dalg24 May 16, 2024
2b7b98a
Use parallel_for instead of parallel_reduce for check
masterleinad May 16, 2024
fc4383a
Fix unique_any_senders nvcc template deduction
G-071 May 19, 2024
226aecf
Properly guard deprecated `Kokkos_Vector.hpp` header self contained t…
dalg24 May 20, 2024
81b63c5
mdspan converting constructors (#6830)
nmm0 May 20, 2024
64fe756
SYCL: Don't use shuffles for top-level reductions (#7009)
masterleinad May 20, 2024
6aa2ad7
Add a CITATION.cff file (#7008)
dutkalex May 20, 2024
f8f0cc4
Always run Graph tests (#7011)
masterleinad May 20, 2024
ce0915b
Fix undefined behavior in is_zero_byte (#7014)
masterleinad May 21, 2024
468faaa
Merge pull request #7015 from G-071/fix_hpx_execution_space_nvcc_comp…
dalg24 May 21, 2024
fa8b501
Disable OpenMPTarget Kokkos::Graph test (does not compile)
dalg24 May 21, 2024
bfe9aa2
Fixup for disabling deprecation warnings with NVC++
dalg24 May 21, 2024
02433b6
Merge pull request #7019 from dalg24/nvhpc_suppress_deprecation_warnings
dalg24 May 21, 2024
f3bd253
Remove unused CudaInternal::cuda_{malloc,free}_async_wrapper
dalg24 May 21, 2024
42e83f1
Merge pull request #7023 from dalg24/remove_unused_cuda_api_wrappers
dalg24 May 21, 2024
068d468
Merge pull request #7018 from dalg24/disable_openmptarget_graph_test
crtrott May 21, 2024
083fb01
Improve `Impl::is_zero_byte()` (#7017)
dalg24 May 22, 2024
6f176cd
OpenMPTarget: Fix compiling Graph tests (#7020)
masterleinad May 22, 2024
cb27c99
SYCL: Skip launch_six Graph test
masterleinad May 22, 2024
c8e0a95
HIP: Use builtin atomic for compare_exchange (#7000)
Rombur May 22, 2024
a5bb0d4
Fix Kokkos README's FENL link
mhoemmen Feb 21, 2024
6a3d918
Merge pull request #6834 from mhoemmen/fix-README-FENL-link
dalg24 May 22, 2024
07a5009
Merge pull request #7024 from masterleinad/sycl_cuda_fix_graph_tests
dalg24 May 22, 2024
24b24d0
Merge pull request #7006 from masterleinad/test_no_default_constructo…
crtrott May 23, 2024
a78d4dd
Copied the deduction guides and test cases over from branch
nliber May 23, 2024
0410363
Refactor: Replace SFINAE by `if constexpr` for `create_mirror*` funct…
pzehner May 23, 2024
7c67b02
Workaround icpc warnings
ndellingwood May 24, 2024
cf791bc
Adding `Kokkos::to_array` (#6375)
nliber May 24, 2024
1c60c80
Merge pull request #7030 from nliber/ctad-teampolicy-v3
dalg24 May 24, 2024
580dba5
Merge pull request #7034 from ndellingwood/issue-7031
dalg24 May 24, 2024
bd107d8
Add assertion in adjacent_difference with overlapping source and dest…
yasahi-hpc Mar 27, 2024
82511c4
Disallow overlap not identical inputs
yasahi-hpc Apr 3, 2024
2b66300
check equality of first and last in advance
yasahi-hpc Apr 3, 2024
f9939f6
Test overlap only if iterators are convertible
yasahi-hpc Apr 5, 2024
2acd84f
add else if for the opposite case
yasahi-hpc Apr 5, 2024
66c5df1
fix check for overlapping iterators
yasahi-hpc Apr 8, 2024
a54a3fd
avoid no-op in no-overlap check
yasahi-hpc Apr 8, 2024
417b127
fix check conditions for overlapping iterators
yasahi-hpc Apr 8, 2024
629b0ee
Formatting expect_no_overlap
yasahi-hpc Apr 8, 2024
2d11144
fix check conditions for overlapping iterators
yasahi-hpc May 16, 2024
c1cc41a
fix conflicts
yasahi-hpc May 16, 2024
1126fed
fix conflicts
yasahi-hpc May 16, 2024
70cbe74
fix conflicts
yasahi-hpc May 16, 2024
cddbec5
fix conflicts
yasahi-hpc May 16, 2024
e7da00e
avoid no-op in no-overlap check
yasahi-hpc Apr 8, 2024
62a06ec
Add maybe unused
yasahi-hpc May 16, 2024
da24fdd
Check must be made on exec space
yasahi-hpc May 16, 2024
07b871f
remove unused lines
yasahi-hpc May 16, 2024
366ea74
Enable checks for debug mode only
yasahi-hpc May 16, 2024
0b3358c
suppress unnecessary iterator conversions
yasahi-hpc May 16, 2024
28d45d8
improve comments for checks
yasahi-hpc May 16, 2024
167de91
fix iterator overlapping check
yasahi-hpc May 16, 2024
027dc0d
fix iterator overlapping check
yasahi-hpc May 17, 2024
bf6301a
fix iterator overlapping check
yasahi-hpc May 17, 2024
7c38c17
Revert "fix iterator overlapping check"
yasahi-hpc May 17, 2024
594180f
Use internval views to check iterator overlaps
yasahi-hpc May 21, 2024
c626131
fix view() method in RandomAccessIterator
yasahi-hpc May 24, 2024
5529803
unit-test for expect_no_overlap
yasahi-hpc May 24, 2024
0173e7e
fix arguments for static 2D View constructor
yasahi-hpc May 24, 2024
52a814e
remove unused s_last from expect_no_overlap
yasahi-hpc May 28, 2024
70893d4
wrapping with EXPECT_NO_THROW for no overlapping cases
yasahi-hpc May 28, 2024
9da24c5
remove NDEBUG check
yasahi-hpc May 28, 2024
4a10a40
fix comment: last element is exclusive
yasahi-hpc May 28, 2024
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,7 @@ jobs:
-DKokkos_ENABLE_DEPRECATED_CODE_4=ON \
-DKokkos_ENABLE_DEPRECATION_WARNINGS=OFF \
-DKokkos_ENABLE_COMPILER_WARNINGS=ON \
-DCMAKE_CXX_FLAGS="-Werror -m32 -DKOKKOS_IMPL_32BIT" \
-DCMAKE_CXX_FLAGS="-Werror -m32" \
-DCMAKE_CXX_COMPILER=g++ \
-DCMAKE_BUILD_TYPE=RelWithDebInfo
- name: Build
Expand Down
5 changes: 3 additions & 2 deletions .jenkins_nightly
Original file line number Diff line number Diff line change
Expand Up @@ -95,7 +95,8 @@ pipeline {
-DKokkos_ENABLE_BENCHMARKS=ON \
-DKokkos_ENABLE_EXAMPLES=ON \
-DKokkos_ENABLE_TESTS=ON \
-DKokkos_ENABLE_DEPRECATION_WARNINGS=OFF \
-DKokkos_ENABLE_DEPRECATED_CODE_4=ON \
-DKokkos_ENABLE_DEPRECATION_WARNINGS=ON \
-DKokkos_ENABLE_SERIAL=ON \
.. && \
make -j8 && ctest --verbose
Expand Down Expand Up @@ -123,7 +124,7 @@ pipeline {
-DKokkos_ARCH_NATIVE=ON \
-DKokkos_ENABLE_COMPILER_WARNINGS=ON \
-DKokkos_ENABLE_DEPRECATED_CODE_4=ON \
-DKokkos_ENABLE_DEPRECATION_WARNINGS=OFF \
-DKokkos_ENABLE_DEPRECATION_WARNINGS=ON \
-DKokkos_ENABLE_TESTS=ON \
-DKokkos_ENABLE_BENCHMARKS=ON \
-DKokkos_ENABLE_HIP=ON \
Expand Down
15 changes: 14 additions & 1 deletion CHANGELOG.md
Original file line number Diff line number Diff line change
@@ -1,5 +1,18 @@
# CHANGELOG

## [4.3.01](https://github.com/kokkos/kokkos/tree/4.3.01)
[Full Changelog](https://github.com/kokkos/kokkos/compare/4.3.00...4.3.01)

### Backend and Architecture Enhancements:

#### HIP:
* MI300 support unified memory [\#6877](https://github.com/kokkos/kokkos/pull/6877)

### Bug Fixes
* Serial: Use the provided execution space instance in TeamPolicy [\#6951](https://github.com/kokkos/kokkos/pull/6951)
* `nvcc_wrapper`: bring back support for `--fmad` option [\#6931](https://github.com/kokkos/kokkos/pull/6931)
* Fix CUDA reduction overflow for `RangePolicy` [\#6578](https://github.com/kokkos/kokkos/pull/6578)

## [4.3.00](https://github.com/kokkos/kokkos/tree/4.3.00) (2024-03-19)
[Full Changelog](https://github.com/kokkos/kokkos/compare/4.2.01...4.3.00)

Expand Down Expand Up @@ -39,7 +52,7 @@
* Make the OpenACC backend asynchronous [\#6772](https://github.com/kokkos/kokkos/pull/6772)

#### Threads:
* Add missing broadcast to TeamThreadRange parallel_scan [\#6601](https://github.com/kokkos/kokkos/pull/6446)
* Add missing broadcast to TeamThreadRange parallel_scan [\#6601](https://github.com/kokkos/kokkos/pull/6601)

#### OpenMP:
* Improve performance of view initializations and filling with zeros [\#6573](https://github.com/kokkos/kokkos/pull/6573)
Expand Down
65 changes: 65 additions & 0 deletions CITATION.cff
Original file line number Diff line number Diff line change
@@ -0,0 +1,65 @@
cff-version: 1.2.0
title: Kokkos
message: >-
If you use this software, please cite the overview paper
type: software
authors:
- name: The Kokkos authors
website: https://kokkos.org/community/team/
identifiers:
- type: url
website: https://kokkos.org/kokkos-core-wiki/citation.html
repository-code: 'https://github.com/kokkos/kokkos'
url: 'https://kokkos.org/'
license: Apache-2.0
preferred-citation:
type: article
authors:
- given-names: Christian R.
family-names: Trott
- given-names: Damien
family-names: Lebrun-Grandié
- given-names: Daniel
family-names: Arndt
- family-names: Ciesko
given-names: Jan
- given-names: Vinh
family-names: Dang
- family-names: Ellingwood
given-names: Nathan
- given-names: Rahulkumar
family-names: Gayatri
- given-names: Evan
family-names: Harvey
- given-names: Daisy S.
family-names: Hollman
- given-names: Dan
family-names: Ibanez
- given-names: Nevin
family-names: Liber
- given-names: Jonathan
family-names: Madsen
- given-names: Jeff
family-names: Miles
- given-names: David
family-names: Poliakoff
- given-names: Amy
family-names: Powell
- given-names: Sivasankaran
family-names: Rajamanickam
- given-names: Mikael
family-names: Simberg
- given-names: Dan
family-names: Sunderland
- given-names: Bruno
family-names: Turcksin
- given-names: Jeremiah
family-names: Wilke
doi: 10.1109/TPDS.2021.3097283
journal: IEEE Transactions on Parallel and Distributed Systems
start: 805
end: 817
title: "Kokkos 3: Programming Model Extensions for the Exascale Era"
volume: 33
issue: 4
year: 2022
49 changes: 8 additions & 41 deletions Copyright.txt
Original file line number Diff line number Diff line change
@@ -1,41 +1,8 @@
//@HEADER
// ************************************************************************
//
// Kokkos v. 3.0
// Copyright (2020) National Technology & Engineering
// Solutions of Sandia, LLC (NTESS).
//
// Under the terms of Contract DE-NA0003525 with NTESS,
// the U.S. Government retains certain rights in this software.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are
// met:
//
// 1. Redistributions of source code must retain the above copyright
// notice, this list of conditions and the following disclaimer.
//
// 2. 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.
//
// 3. Neither the name of the Corporation nor the names of the
// contributors may be used to endorse or promote products derived from
// this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY NTESS "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 NTESS OR THE
// CONTRIBUTORS 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.
//
// Questions? Contact Christian R. Trott ([email protected])
//
// ************************************************************************
//@HEADER
************************************************************************

Kokkos v. 4.0
Copyright (2022) National Technology & Engineering
Solutions of Sandia, LLC (NTESS).

Under the terms of Contract DE-NA0003525 with NTESS,
the U.S. Government retains certain rights in this software.
10 changes: 0 additions & 10 deletions LICENSE
Original file line number Diff line number Diff line change
@@ -1,13 +1,3 @@
************************************************************************

Kokkos v. 4.0
Copyright (2022) National Technology & Engineering
Solutions of Sandia, LLC (NTESS).

Under the terms of Contract DE-NA0003525 with NTESS,
the U.S. Government retains certain rights in this software.


==============================================================================
Kokkos is under the Apache License v2.0 with LLVM Exceptions:
==============================================================================
Expand Down
10 changes: 6 additions & 4 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,8 @@ backends in development.

**Kokkos Core is part of the [Kokkos C++ Performance Portability Programming Ecosystem](https://kokkos.org/about/abstract/).**

Kokkos is a [Linux Foundation](https://linuxfoundation.org) project.

## Learning about Kokkos

To start learning about Kokkos:
Expand All @@ -28,12 +30,12 @@ To start learning about Kokkos:

The latest release of Kokkos can be obtained from the [GitHub releases page](https://github.com/kokkos/kokkos/releases/latest).

The current release is [4.2.01](https://github.com/kokkos/kokkos/releases/tag/4.2.01).
The current release is [4.3.00](https://github.com/kokkos/kokkos/releases/tag/4.3.00).

```bash
curl -OJ -L https://github.com/kokkos/kokkos/archive/refs/tags/4.2.01.tar.gz
curl -OJ -L https://github.com/kokkos/kokkos/archive/refs/tags/4.3.00.tar.gz
# Or with wget
wget https://github.com/kokkos/kokkos/archive/refs/tags/4.2.01.tar.gz
wget https://github.com/kokkos/kokkos/archive/refs/tags/4.3.00.tar.gz
```

To clone the latest development version of Kokkos from GitHub:
Expand All @@ -44,7 +46,7 @@ git clone -b develop https://github.com/kokkos/kokkos.git

### Building Kokkos

To build Kokkos, you will need to have a C++ compiler that supports C++14 or later.
To build Kokkos, you will need to have a C++ compiler that supports C++17 or later.
All requirements including minimum and primary tested compiler versions can be found [here](https://kokkos.org/kokkos-core-wiki/requirements.html).

Building and installation instructions are described [here](https://kokkos.org/kokkos-core-wiki/building.html).
Expand Down
10 changes: 10 additions & 0 deletions algorithms/src/std_algorithms/impl/Kokkos_AdjacentDifference.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,11 @@ OutputIteratorType adjacent_difference_exespace_impl(
return first_dest;
}

#ifdef KOKKOS_ENABLE_DEBUG
// check for overlapping iterators
Impl::expect_no_overlap(first_from, last_from, first_dest);
#endif

// run
const auto num_elements =
Kokkos::Experimental::distance(first_from, last_from);
Expand Down Expand Up @@ -114,6 +119,11 @@ KOKKOS_FUNCTION OutputIteratorType adjacent_difference_team_impl(
return first_dest;
}

#ifdef KOKKOS_ENABLE_DEBUG
// check for overlapping iterators
Impl::expect_no_overlap(first_from, last_from, first_dest);
#endif

// run
const auto num_elements =
Kokkos::Experimental::distance(first_from, last_from);
Expand Down
61 changes: 54 additions & 7 deletions algorithms/src/std_algorithms/impl/Kokkos_Constraints.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -24,18 +24,21 @@ namespace Kokkos {
namespace Experimental {
namespace Impl {

template <class T>
class RandomAccessIterator;

template <typename T, typename enable = void>
struct is_admissible_to_kokkos_std_algorithms : std::false_type {};

template <typename T>
struct is_admissible_to_kokkos_std_algorithms<
T, std::enable_if_t< ::Kokkos::is_view<T>::value && T::rank() == 1 &&
(std::is_same<typename T::traits::array_layout,
Kokkos::LayoutLeft>::value ||
std::is_same<typename T::traits::array_layout,
Kokkos::LayoutRight>::value ||
std::is_same<typename T::traits::array_layout,
Kokkos::LayoutStride>::value)> >
T, std::enable_if_t<::Kokkos::is_view<T>::value && T::rank() == 1 &&
(std::is_same<typename T::traits::array_layout,
Kokkos::LayoutLeft>::value ||
std::is_same<typename T::traits::array_layout,
Kokkos::LayoutRight>::value ||
std::is_same<typename T::traits::array_layout,
Kokkos::LayoutStride>::value)>>
: std::true_type {};

template <class ViewType>
Expand All @@ -58,6 +61,18 @@ using is_iterator = Kokkos::is_detected<iterator_category_t, T>;
template <class T>
inline constexpr bool is_iterator_v = is_iterator<T>::value;

template <typename ViewType>
struct is_kokkos_iterator : std::false_type {};

template <typename ViewType>
struct is_kokkos_iterator<RandomAccessIterator<ViewType>> {
static constexpr bool value =
is_admissible_to_kokkos_std_algorithms<ViewType>::value;
};

template <class T>
inline constexpr bool is_kokkos_iterator_v = is_kokkos_iterator<T>::value;

//
// are_iterators
//
Expand Down Expand Up @@ -215,6 +230,38 @@ KOKKOS_INLINE_FUNCTION void expect_valid_range(IteratorType first,
(void)last;
}

//
// Check if kokkos iterators are overlapping
//
template <typename IteratorType1, typename IteratorType2>
KOKKOS_INLINE_FUNCTION void expect_no_overlap(
[[maybe_unused]] IteratorType1 first, [[maybe_unused]] IteratorType1 last,
[[maybe_unused]] IteratorType2 s_first) {
if constexpr (is_kokkos_iterator_v<IteratorType1> &&
is_kokkos_iterator_v<IteratorType2>) {
auto const view1 = first.view();
auto const view2 = s_first.view();

std::size_t stride1 = view1.stride(0);
std::size_t stride2 = view2.stride(0);
ptrdiff_t first_diff = view1.data() - view2.data();

// FIXME If strides are not identical, checks may not be made
// with the cost of O(1)
// Currently, checks are made only if strides are identical
// If first_diff == 0, there is already an overlap
if (stride1 == stride2 || first_diff == 0) {
[[maybe_unused]] bool is_no_overlap = (first_diff % stride1);
auto* first_pointer1 = view1.data();
auto* first_pointer2 = view2.data();
[[maybe_unused]] auto* last_pointer1 = first_pointer1 + (last - first);
[[maybe_unused]] auto* last_pointer2 = first_pointer2 + (last - first);
KOKKOS_EXPECTS(first_pointer1 >= last_pointer2 ||
last_pointer1 <= first_pointer2 || is_no_overlap);
}
}
}

} // namespace Impl
} // namespace Experimental
} // namespace Kokkos
Expand Down
39 changes: 39 additions & 0 deletions algorithms/src/std_algorithms/impl/Kokkos_RandomAccessIterator.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,38 @@ class RandomAccessIterator< ::Kokkos::View<DataType, Args...> > {
ptrdiff_t current_index)
: m_view(view), m_current_index(current_index) {}

// FIXME The C++20 requires expression is not supported with Clang 9 and GCC 9
// The following guards is unsufficient until we increase our minimum CXX20
// compiler requirements.
// #ifndef KOKKOS_ENABLE_CXX17 // C++20 and beyond
// We replace the Kokkos guards with standard C++ feature testing in the
// meantime.
#if (defined(__cpp_concepts) && (__cpp_concepts >= 201907L)) && \
(defined(__cpp_conditional_explicit) && \
(__cpp_conditional_explicit >= 201806L))
template <class OtherViewType>
requires(std::is_constructible_v<view_type, OtherViewType>) KOKKOS_FUNCTION
explicit(!std::is_convertible_v<OtherViewType, view_type>)
RandomAccessIterator(const RandomAccessIterator<OtherViewType>& other)
: m_view(other.m_view), m_current_index(other.m_current_index) {}
#else
template <
class OtherViewType,
std::enable_if_t<std::is_constructible_v<view_type, OtherViewType> &&
!std::is_convertible_v<OtherViewType, view_type>,
int> = 0>
KOKKOS_FUNCTION explicit RandomAccessIterator(
const RandomAccessIterator<OtherViewType>& other)
: m_view(other.m_view), m_current_index(other.m_current_index) {}

template <class OtherViewType,
std::enable_if_t<std::is_convertible_v<OtherViewType, view_type>,
int> = 0>
KOKKOS_FUNCTION RandomAccessIterator(
const RandomAccessIterator<OtherViewType>& other)
: m_view(other.m_view), m_current_index(other.m_current_index) {}
#endif

KOKKOS_FUNCTION
iterator_type& operator++() {
++m_current_index;
Expand Down Expand Up @@ -152,9 +184,16 @@ class RandomAccessIterator< ::Kokkos::View<DataType, Args...> > {
KOKKOS_FUNCTION
reference operator*() const { return m_view(m_current_index); }

KOKKOS_FUNCTION
view_type view() const { return m_view; }

private:
view_type m_view;
ptrdiff_t m_current_index = 0;

// Needed for the converting constructor accepting another iterator
template <class>
friend class RandomAccessIterator;
};

} // namespace Impl
Expand Down
Loading