24687578389

Committed 20 Apr 2026 08:03PM UTC coverage: 78.429% (-3.6%) from 82.034%

Build # 24687578389

Build Type

push

github

Committed by

web-flow

Commit Message

Migrate `dpctl.tensor` into `dpnp.tensor` (#2856)

This PR migrates the tensor implementation from `dpctl.tensor` into
`dpnp.tensor` making dpnp the primary owner of the Array API-compliant
tensor layer

Major changes:

- Move compiled C++/SYCL extensions (`_tensor_impl,
_tensor_elementwise_impl, _tensor_reductions_impl, _tensor_sorting_impl,
_tensor_accumulation_impl, tensor linalg`) into `dpnp.tensor`
- Move `usm_ndarray`, `compute-follows-data utilities` and tensor
`tests` from dpctl
- Replace all `dpctl.tensor` references with `dpnp.tensor` in
docstrings, error messages and comments
- Remove redundant dpctl.tensor C-API  interface
- Add `tensor.rst` documentation page describing the module, its
relationship to `dpnp.ndarray` and `dpctl` and linking to the `dpctl
0.21.1 API` reference

This simplifies maintenance, reduces cross-project dependencies and
enables independent development and release cycles

Coverage Stats

1573 of 2908 branches covered (54.09%)

Branch coverage included in aggregate %.

6973 of 9803 new or added lines in 203 files covered. (71.13%)

1 existing line in 1 file now uncovered.

26259 of 32579 relevant lines covered (80.6%)

7622.15 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

95.97

/dpnp/tensor/libtensor/include/utils/type_dispatch_building.hpp

//*****************************************************************************
// Copyright (c) 2026, Intel 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 copyright holder 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 THE COPYRIGHT HOLDER OR 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.
//*****************************************************************************
///
/// \file
/// This file defines class to implement dispatch tables for pair of types
//===----------------------------------------------------------------------===//

#pragma once

#include <cassert>
#include <complex>
#include <cstdint>
#include <type_traits>
#include <vector>

#include <sycl/sycl.hpp>

namespace dpnp::tensor::type_dispatch
{
enum class typenum_t : int
{
    BOOL = 0,
    INT8, // 1
    UINT8,
    INT16,
    UINT16,
    INT32, // 5
    UINT32,
    INT64,
    UINT64,
    HALF,
    FLOAT, // 10
    DOUBLE,
    CFLOAT,
    CDOUBLE, // 13
};
inline constexpr int num_types = 14; // number of elements in typenum_t

template <typename funcPtrT,
          template <typename fnT, typename D, typename S> typename factory,
          int _num_types>
class DispatchTableBuilder
{
private:
    template <typename dstTy>
    const std::vector<funcPtrT> row_per_dst_type() const
    {
        std::vector<funcPtrT> per_dstTy = {
            factory<funcPtrT, dstTy, bool>{}.get(),
            factory<funcPtrT, dstTy, std::int8_t>{}.get(),
            factory<funcPtrT, dstTy, std::uint8_t>{}.get(),
            factory<funcPtrT, dstTy, std::int16_t>{}.get(),
            factory<funcPtrT, dstTy, std::uint16_t>{}.get(),
            factory<funcPtrT, dstTy, std::int32_t>{}.get(),
            factory<funcPtrT, dstTy, std::uint32_t>{}.get(),
            factory<funcPtrT, dstTy, std::int64_t>{}.get(),
            factory<funcPtrT, dstTy, std::uint64_t>{}.get(),
            factory<funcPtrT, dstTy, sycl::half>{}.get(),
            factory<funcPtrT, dstTy, float>{}.get(),
            factory<funcPtrT, dstTy, double>{}.get(),
            factory<funcPtrT, dstTy, std::complex<float>>{}.get(),
            factory<funcPtrT, dstTy, std::complex<double>>{}.get()};
        assert(per_dstTy.size() == _num_types);
        return per_dstTy;
    }

public:
    DispatchTableBuilder() = default;
    ~DispatchTableBuilder() = default;

    void populate_dispatch_table(funcPtrT table[][_num_types]) const
    {
        const auto map_by_dst_type = {row_per_dst_type<bool>(),
                                      row_per_dst_type<std::int8_t>(),
                                      row_per_dst_type<std::uint8_t>(),
                                      row_per_dst_type<std::int16_t>(),
                                      row_per_dst_type<std::uint16_t>(),
                                      row_per_dst_type<std::int32_t>(),
                                      row_per_dst_type<std::uint32_t>(),
                                      row_per_dst_type<std::int64_t>(),
                                      row_per_dst_type<std::uint64_t>(),
                                      row_per_dst_type<sycl::half>(),
                                      row_per_dst_type<float>(),
                                      row_per_dst_type<double>(),
                                      row_per_dst_type<std::complex<float>>(),
                                      row_per_dst_type<std::complex<double>>()};
        assert(map_by_dst_type.size() == _num_types);
        int dst_id = 0;
        for (const auto &row : map_by_dst_type) {
            int src_id = 0;
            for (const auto &fn_ptr : row) {
                table[dst_id][src_id] = fn_ptr;
                ++src_id;
            }
            ++dst_id;
        }
    }
};

template <typename funcPtrT,
          template <typename fnT, typename T> typename factory,
          int _num_types>
class DispatchVectorBuilder
{
private:
    template <typename Ty>
    const funcPtrT func_per_type() const
    {
        funcPtrT f = factory<funcPtrT, Ty>{}.get();
        return f;
    }

public:
    DispatchVectorBuilder() = default;
    ~DispatchVectorBuilder() = default;

    void populate_dispatch_vector(funcPtrT vector[]) const
    {
        const auto fn_map_by_type = {func_per_type<bool>(),
                                     func_per_type<std::int8_t>(),
                                     func_per_type<std::uint8_t>(),
                                     func_per_type<std::int16_t>(),
                                     func_per_type<std::uint16_t>(),
                                     func_per_type<std::int32_t>(),
                                     func_per_type<std::uint32_t>(),
                                     func_per_type<std::int64_t>(),
                                     func_per_type<std::uint64_t>(),
                                     func_per_type<sycl::half>(),
                                     func_per_type<float>(),
                                     func_per_type<double>(),
                                     func_per_type<std::complex<float>>(),
                                     func_per_type<std::complex<double>>()};
        assert(fn_map_by_type.size() == _num_types);
        int ty_id = 0;
        for (const auto &fn : fn_map_by_type) {
            vector[ty_id] = fn;
            ++ty_id;
        }
    }
};

/*! @brief struct to define result_type typename for Ty == ArgTy */
template <typename Ty, typename ArgTy, typename ResTy = ArgTy>
struct TypeMapResultEntry : std::is_same<Ty, ArgTy>
{
    using result_type = ResTy;
};

/*! @brief struct to define result_type typename for Ty1 == ArgTy1 && Ty2 ==
 * ArgTy2 */
template <typename Ty1,
          typename ArgTy1,
          typename Ty2,
          typename ArgTy2,
          typename ResTy>
struct BinaryTypeMapResultEntry
    : std::conjunction<std::is_same<Ty1, ArgTy1>, std::is_same<Ty2, ArgTy2>>
{
    using result_type = ResTy;
};

/*! @brief fall-through struct with specified result_type, usually void */
template <typename Ty = void>
struct DefaultResultEntry : std::true_type
{
    using result_type = Ty;
};

/*! @brief Utility struct to convert C++ type into typeid integer */
template <typename T>
struct GetTypeid
{
    int get()
    {
        if constexpr (std::is_same_v<T, bool>) {
            return static_cast<int>(typenum_t::BOOL);
        }
        else if constexpr (std::is_same_v<T, std::int8_t>) {
            return static_cast<int>(typenum_t::INT8);
        }
        else if constexpr (std::is_same_v<T, std::uint8_t>) {
            return static_cast<int>(typenum_t::UINT8);
        }
        else if constexpr (std::is_same_v<T, std::int16_t>) {
            return static_cast<int>(typenum_t::INT16);
        }
        else if constexpr (std::is_same_v<T, std::uint16_t>) {
            return static_cast<int>(typenum_t::UINT16);
        }
        else if constexpr (std::is_same_v<T, std::int32_t>) {
            return static_cast<int>(typenum_t::INT32);
        }
        else if constexpr (std::is_same_v<T, std::uint32_t>) {
            return static_cast<int>(typenum_t::UINT32);
        }
        else if constexpr (std::is_same_v<T, std::int64_t>) {
            return static_cast<int>(typenum_t::INT64);
        }
        else if constexpr (std::is_same_v<T, std::uint64_t>) {
            return static_cast<int>(typenum_t::UINT64);
        }
        else if constexpr (std::is_same_v<T, sycl::half>) {
            return static_cast<int>(typenum_t::HALF);
        }
        else if constexpr (std::is_same_v<T, float>) {
            return static_cast<int>(typenum_t::FLOAT);
        }
        else if constexpr (std::is_same_v<T, double>) {
            return static_cast<int>(typenum_t::DOUBLE);
        }
        else if constexpr (std::is_same_v<T, std::complex<float>>) {
            return static_cast<int>(typenum_t::CFLOAT);
        }
        else if constexpr (std::is_same_v<T, std::complex<double>>) {
            return static_cast<int>(typenum_t::CDOUBLE);
        }
        else if constexpr (std::is_same_v<T, void>) { // special token
            return -1;
        }

        assert(("Unsupported type T", false));
        return -2;
    }
};

/*! @brief Class to generate vector of null function pointers */
template <typename FunPtrT>
struct NullPtrVector
{

    using value_type = FunPtrT;
    using const_reference = value_type const &;

    NullPtrVector() : val(nullptr) {}

    const_reference operator[](int) const { return val; }

private:
    value_type val;
};

/*! @brief Class to generate table of null function pointers */
template <typename FunPtrT>
struct NullPtrTable
{
    using value_type = NullPtrVector<FunPtrT>;
    using const_reference = value_type const &;

    NullPtrTable() : val() {}

    const_reference operator[](int) const { return val; }

private:
    value_type val;
};

template <typename Ty1, typename ArgTy, typename Ty2, typename outTy>
struct TypePairDefinedEntry
    : std::conjunction<std::is_same<Ty1, ArgTy>, std::is_same<Ty2, outTy>>
{
    static constexpr bool is_defined = true;
};

struct NotDefinedEntry : std::true_type
{
    static constexpr bool is_defined = false;
};
} // namespace dpnp::tensor::type_dispatch

1	//*****************************************************************************
2	// Copyright (c) 2026, Intel Corporation
3	// All rights reserved.
4	//
5	// Redistribution and use in source and binary forms, with or without
6	// modification, are permitted provided that the following conditions are met:
7	// - Redistributions of source code must retain the above copyright notice,
8	// this list of conditions and the following disclaimer.
9	// - Redistributions in binary form must reproduce the above copyright notice,
10	// this list of conditions and the following disclaimer in the documentation
11	// and/or other materials provided with the distribution.
12	// - Neither the name of the copyright holder nor the names of its contributors
13	// may be used to endorse or promote products derived from this software
14	// without specific prior written permission.
15	//
16	// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
17	// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18	// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
19	// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
20	// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
21	// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
22	// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
23	// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
24	// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
25	// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
26	// THE POSSIBILITY OF SUCH DAMAGE.
27	//*****************************************************************************
28	///
29	/// \file
30	/// This file defines class to implement dispatch tables for pair of types
31	//===----------------------------------------------------------------------===//
32
33	#pragma once
34
35	#include <cassert>
36	#include <complex>
37	#include <cstdint>
38	#include <type_traits>
39	#include <vector>
40
41	#include <sycl/sycl.hpp>
42
43	namespace dpnp::tensor::type_dispatch
44	{
45	enum class typenum_t : int
46	{
47	BOOL = 0,
48	INT8, // 1
49	UINT8,
50	INT16,
51	UINT16,
52	INT32, // 5
53	UINT32,
54	INT64,
55	UINT64,
56	HALF,
57	FLOAT, // 10
58	DOUBLE,
59	CFLOAT,
60	CDOUBLE, // 13
61	};
62	inline constexpr int num_types = 14; // number of elements in typenum_t
63
64	template <typename funcPtrT,
65	template <typename fnT, typename D, typename S> typename factory,
66	int _num_types>
67	class DispatchTableBuilder
68	{
69	private:
70	template <typename dstTy>
71	const std::vector<funcPtrT> row_per_dst_type() const
72	{	3,640✔
73	std::vector<funcPtrT> per_dstTy = {	3,640✔
74	factory<funcPtrT, dstTy, bool>{}.get(),	3,640✔
75	factory<funcPtrT, dstTy, std::int8_t>{}.get(),	3,640✔
76	factory<funcPtrT, dstTy, std::uint8_t>{}.get(),	3,640✔
77	factory<funcPtrT, dstTy, std::int16_t>{}.get(),	3,640✔
78	factory<funcPtrT, dstTy, std::uint16_t>{}.get(),	3,640✔
79	factory<funcPtrT, dstTy, std::int32_t>{}.get(),	3,640✔
80	factory<funcPtrT, dstTy, std::uint32_t>{}.get(),	3,640✔
81	factory<funcPtrT, dstTy, std::int64_t>{}.get(),	3,640✔
82	factory<funcPtrT, dstTy, std::uint64_t>{}.get(),	3,640✔
83	factory<funcPtrT, dstTy, sycl::half>{}.get(),	3,640✔
84	factory<funcPtrT, dstTy, float>{}.get(),	3,640✔
85	factory<funcPtrT, dstTy, double>{}.get(),	3,640✔
86	factory<funcPtrT, dstTy, std::complex<float>>{}.get(),	3,640✔
87	factory<funcPtrT, dstTy, std::complex<double>>{}.get()};	3,640✔
88	assert(per_dstTy.size() == _num_types);	3,640!
89	return per_dstTy;	3,640✔
90	}	3,640✔
91
92	public:
93	DispatchTableBuilder() = default;
94	~DispatchTableBuilder() = default;
95
96	void populate_dispatch_table(funcPtrT table[][_num_types]) const
97	{	260✔
98	const auto map_by_dst_type = {row_per_dst_type<bool>(),	260✔
99	row_per_dst_type<std::int8_t>(),	260✔
100	row_per_dst_type<std::uint8_t>(),	260✔
101	row_per_dst_type<std::int16_t>(),	260✔
102	row_per_dst_type<std::uint16_t>(),	260✔
103	row_per_dst_type<std::int32_t>(),	260✔
104	row_per_dst_type<std::uint32_t>(),	260✔
105	row_per_dst_type<std::int64_t>(),	260✔
106	row_per_dst_type<std::uint64_t>(),	260✔
107	row_per_dst_type<sycl::half>(),	260✔
108	row_per_dst_type<float>(),	260✔
109	row_per_dst_type<double>(),	260✔
110	row_per_dst_type<std::complex<float>>(),	260✔
111	row_per_dst_type<std::complex<double>>()};	260✔
112	assert(map_by_dst_type.size() == _num_types);	260!
113	int dst_id = 0;	260✔
114	for (const auto &row : map_by_dst_type) {	3,640✔
115	int src_id = 0;	3,640✔
116	for (const auto &fn_ptr : row) {	50,960✔
117	table[dst_id][src_id] = fn_ptr;	50,960✔
118	++src_id;	50,960✔
119	}	50,960✔
120	++dst_id;	3,640✔
121	}	3,640✔
122	}	260✔
123	};
124
125	template <typename funcPtrT,
126	template <typename fnT, typename T> typename factory,
127	int _num_types>
128	class DispatchVectorBuilder
129	{
130	private:
131	template <typename Ty>
132	const funcPtrT func_per_type() const
133	{	7,784✔
134	funcPtrT f = factory<funcPtrT, Ty>{}.get();	7,784✔
135	return f;	7,784✔
136	}	7,784✔
137
138	public:
139	DispatchVectorBuilder() = default;
140	~DispatchVectorBuilder() = default;
141
142	void populate_dispatch_vector(funcPtrT vector[]) const
143	{	556✔
144	const auto fn_map_by_type = {func_per_type<bool>(),	556✔
145	func_per_type<std::int8_t>(),	556✔
146	func_per_type<std::uint8_t>(),	556✔
147	func_per_type<std::int16_t>(),	556✔
148	func_per_type<std::uint16_t>(),	556✔
149	func_per_type<std::int32_t>(),	556✔
150	func_per_type<std::uint32_t>(),	556✔
151	func_per_type<std::int64_t>(),	556✔
152	func_per_type<std::uint64_t>(),	556✔
153	func_per_type<sycl::half>(),	556✔
154	func_per_type<float>(),	556✔
155	func_per_type<double>(),	556✔
156	func_per_type<std::complex<float>>(),	556✔
157	func_per_type<std::complex<double>>()};	556✔
158	assert(fn_map_by_type.size() == _num_types);	556!
159	int ty_id = 0;	556✔
160	for (const auto &fn : fn_map_by_type) {	7,784✔
161	vector[ty_id] = fn;	7,784✔
162	++ty_id;	7,784✔
163	}	7,784✔
164	}	556✔
165	};
166
167	/! @brief struct to define result_type typename for Ty == ArgTy /
168	template <typename Ty, typename ArgTy, typename ResTy = ArgTy>
169	struct TypeMapResultEntry : std::is_same<Ty, ArgTy>
170	{
171	using result_type = ResTy;
172	};
173
174	/*! @brief struct to define result_type typename for Ty1 == ArgTy1 && Ty2 ==
175	* ArgTy2 */
176	template <typename Ty1,
177	typename ArgTy1,
178	typename Ty2,
179	typename ArgTy2,
180	typename ResTy>
181	struct BinaryTypeMapResultEntry
182	: std::conjunction<std::is_same<Ty1, ArgTy1>, std::is_same<Ty2, ArgTy2>>
183	{
184	using result_type = ResTy;
185	};
186
187	/! @brief fall-through struct with specified result_type, usually void /
188	template <typename Ty = void>
189	struct DefaultResultEntry : std::true_type
190	{
191	using result_type = Ty;
192	};
193
194	/! @brief Utility struct to convert C++ type into typeid integer /
195	template <typename T>
196	struct GetTypeid
197	{
198	int get()
199	{	20,608✔
200	if constexpr (std::is_same_v<T, bool>) {	20,608✔
201	return static_cast<int>(typenum_t::BOOL);	8✔
202	}
203	else if constexpr (std::is_same_v<T, std::int8_t>) {	32✔
204	return static_cast<int>(typenum_t::INT8);	32✔
205	}
206	else if constexpr (std::is_same_v<T, std::uint8_t>) {	60✔
207	return static_cast<int>(typenum_t::UINT8);	60✔
208	}
209	else if constexpr (std::is_same_v<T, std::int16_t>) {	28✔
210	return static_cast<int>(typenum_t::INT16);	28✔
211	}
212	else if constexpr (std::is_same_v<T, std::uint16_t>) {	28✔
213	return static_cast<int>(typenum_t::UINT16);	28✔
214	}
215	else if constexpr (std::is_same_v<T, std::int32_t>) {	40✔
216	return static_cast<int>(typenum_t::INT32);	40✔
217	}
218	else if constexpr (std::is_same_v<T, std::uint32_t>) {	28✔
219	return static_cast<int>(typenum_t::UINT32);	28✔
220	}
221	else if constexpr (std::is_same_v<T, std::int64_t>) {	28✔
222	return static_cast<int>(typenum_t::INT64);	28✔
223	}
224	else if constexpr (std::is_same_v<T, std::uint64_t>) {	28✔
225	return static_cast<int>(typenum_t::UINT64);	28✔
226	}
227	else if constexpr (std::is_same_v<T, sycl::half>) {	88✔
228	return static_cast<int>(typenum_t::HALF);	88✔
229	}
230	else if constexpr (std::is_same_v<T, float>) {	276✔
231	return static_cast<int>(typenum_t::FLOAT);	276✔
232	}
233	else if constexpr (std::is_same_v<T, double>) {	272✔
234	return static_cast<int>(typenum_t::DOUBLE);	272✔
235	}
236	else if constexpr (std::is_same_v<T, std::complex<float>>) {	104✔
237	return static_cast<int>(typenum_t::CFLOAT);	104✔
238	}
239	else if constexpr (std::is_same_v<T, std::complex<double>>) {	104✔
240	return static_cast<int>(typenum_t::CDOUBLE);	104✔
241	}
242	else if constexpr (std::is_same_v<T, void>) { // special token	19,484✔
243	return -1;	19,484✔
244	}	19,484✔
245
246	assert(("Unsupported type T", false));	20,608!
NEW 247	return -2;	×
248	}	20,608✔
249	};
250
251	/! @brief Class to generate vector of null function pointers /
252	template <typename FunPtrT>
253	struct NullPtrVector
254	{
255
256	using value_type = FunPtrT;
257	using const_reference = value_type const &;
258
259	NullPtrVector() : val(nullptr) {}	65,873✔
260
261	const_reference operator[](int) const { return val; }	19✔
262
263	private:
264	value_type val;
265	};
266
267	/! @brief Class to generate table of null function pointers /
268	template <typename FunPtrT>
269	struct NullPtrTable
270	{
271	using value_type = NullPtrVector<FunPtrT>;
272	using const_reference = value_type const &;
273
274	NullPtrTable() : val() {}	59,844✔
275
276	const_reference operator[](int) const { return val; }	19✔
277
278	private:
279	value_type val;
280	};
281
282	template <typename Ty1, typename ArgTy, typename Ty2, typename outTy>
283	struct TypePairDefinedEntry
284	: std::conjunction<std::is_same<Ty1, ArgTy>, std::is_same<Ty2, outTy>>
285	{
286	static constexpr bool is_defined = true;
287	};
288
289	struct NotDefinedEntry : std::true_type
290	{
291	static constexpr bool is_defined = false;
292	};
293	} // namespace dpnp::tensor::type_dispatch

IntelPython / dpnp / 24687578389

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous