From 31e1307c3131d03384ef83c33472309d22d2e79d Mon Sep 17 00:00:00 2001 From: Wei Kang Date: Fri, 19 Nov 2021 14:47:09 +0800 Subject: [PATCH 1/4] Fix ctc graph (make aux_labels of final arcs -1) (#877) --- k2/csrc/fsa_algo.cu | 6 +++--- k2/csrc/fsa_algo_test.cu | 8 ++++---- k2/python/tests/ctc_graph_test.py | 14 +++++++------- 3 files changed, 14 insertions(+), 14 deletions(-) diff --git a/k2/csrc/fsa_algo.cu b/k2/csrc/fsa_algo.cu index 9cb3ae023..c7106bc0c 100644 --- a/k2/csrc/fsa_algo.cu +++ b/k2/csrc/fsa_algo.cu @@ -709,7 +709,7 @@ FsaVec CtcGraphs(const Ragged &symbols, bool modified /*= false*/, case 2: // the arc pointing to the next symbol state arc.label = next_symbol; aux_labels_value = sym_state_idx01 + 1 == sym_final_state ? - 0 : next_symbol; + -1 : next_symbol; arc.dest_state = state_idx1 + 2; break; default: @@ -720,8 +720,8 @@ FsaVec CtcGraphs(const Ragged &symbols, bool modified /*= false*/, K2_CHECK_LT(arc_idx2, 2); arc.label = arc_idx2 == 0 ? 0 : current_symbol; arc.dest_state = arc_idx2 == 0 ? state_idx1 : state_idx1 + 1; - aux_labels_value = (arc_idx2 == 0 || final_state) ? - 0 : current_symbol; + aux_labels_value = arc_idx2 == 0 ? 0 : current_symbol; + if (final_state && arc_idx2 != 0) aux_labels_value = -1; } arcs_data[arc_idx012] = arc; if (aux_labels) aux_labels_data[arc_idx012] = aux_labels_value; diff --git a/k2/csrc/fsa_algo_test.cu b/k2/csrc/fsa_algo_test.cu index 38970edf2..903514318 100644 --- a/k2/csrc/fsa_algo_test.cu +++ b/k2/csrc/fsa_algo_test.cu @@ -1293,8 +1293,8 @@ TEST(FsaAlgo, TestCtcGraph) { " [ 4 4 0 0 4 5 3 0 ] [ 5 6 0 0 5 5 3 0 5 7 -1 0 ] " " [ 6 6 0 0 6 7 -1 0 ] [ ] ] ]"); Array1 aux_labels_ref(c, "[ 0 1 0 0 2 0 2 0 0 0 2 0 0 3 " - " 0 3 0 0 0 0 0 0 1 0 0 2 0 2 " - " 0 0 3 0 3 0 0 0 0 0 ]"); + " 0 3 0 0 -1 0 -1 0 1 0 0 2 0 2 " + " 0 0 3 0 3 0 0 -1 0 -1 ]"); K2_CHECK(Equal(graph, graph_ref)); K2_CHECK(Equal(aux_labels, aux_labels_ref)); } @@ -1315,8 +1315,8 @@ TEST(FsaAlgo, TestCtcGraphSimplified) { " [ 4 4 0 0 4 5 3 0 ] [ 5 6 0 0 5 5 3 0 5 7 -1 0 ] " " [ 6 6 0 0 6 7 -1 0 ] [ ] ] ]"); Array1 aux_labels_ref(c, "[ 0 1 0 0 2 0 2 0 0 2 0 2 0 " - " 0 3 0 3 0 0 0 0 0 0 1 0 0 2 " - " 0 2 0 0 3 0 3 0 0 0 0 0 ]"); + " 0 3 0 3 0 0 -1 0 -1 0 1 0 0 2 " + " 0 2 0 0 3 0 3 0 0 -1 0 -1 ]"); K2_CHECK(Equal(graph, graph_ref)); K2_CHECK(Equal(aux_labels, aux_labels_ref)); } diff --git a/k2/python/tests/ctc_graph_test.py b/k2/python/tests/ctc_graph_test.py index 569c1d8da..285a55989 100644 --- a/k2/python/tests/ctc_graph_test.py +++ b/k2/python/tests/ctc_graph_test.py @@ -50,14 +50,14 @@ def test(self): '0 0 0 0 0', '0 1 1 1 0', '1 2 0 0 0', '1 1 1 0 0', '1 3 2 2 0', '2 2 0 0 0', '2 3 2 2 0', '3 4 0 0 0', '3 3 2 0 0', '4 4 0 0 0', '4 5 2 2 0', '5 6 0 0 0', - '5 5 2 0 0', '5 7 -1 0 0', '6 6 0 0 0', '6 7 -1 0 0', '7' + '5 5 2 0 0', '5 7 -1 -1 0', '6 6 0 0 0', '6 7 -1 -1 0', '7' ]) expected_str1 = '\n'.join([ '0 0 0 0 0', '0 1 1 1 0', '1 2 0 0 0', '1 1 1 0 0', '1 3 2 2 0', '2 2 0 0 0', '2 3 2 2 0', '3 4 0 0 0', '3 3 2 0 0', '3 5 3 3 0', '4 4 0 0 0', '4 5 3 3 0', - '5 6 0 0 0', '5 5 3 0 0', '5 7 -1 0 0', '6 6 0 0 0', - '6 7 -1 0 0', '7' + '5 6 0 0 0', '5 5 3 0 0', '5 7 -1 -1 0', '6 6 0 0 0', + '6 7 -1 -1 0', '7' ]) actual_str_ragged0 = k2.to_str_simple(fsa_vec_ragged[0].to('cpu')) actual_str_ragged1 = k2.to_str_simple(fsa_vec_ragged[1].to('cpu')) @@ -81,15 +81,15 @@ def test_simplified(self): '0 0 0 0 0', '0 1 1 1 0', '1 2 0 0 0', '1 1 1 0 0', '1 3 2 2 0', '2 2 0 0 0', '2 3 2 2 0', '3 4 0 0 0', '3 3 2 0 0', '3 5 2 2 0', '4 4 0 0 0', '4 5 2 2 0', - '5 6 0 0 0', '5 5 2 0 0', '5 7 -1 0 0', '6 6 0 0 0', - '6 7 -1 0 0', '7' + '5 6 0 0 0', '5 5 2 0 0', '5 7 -1 -1 0', '6 6 0 0 0', + '6 7 -1 -1 0', '7' ]) expected_str1 = '\n'.join([ '0 0 0 0 0', '0 1 1 1 0', '1 2 0 0 0', '1 1 1 0 0', '1 3 2 2 0', '2 2 0 0 0', '2 3 2 2 0', '3 4 0 0 0', '3 3 2 0 0', '3 5 3 3 0', '4 4 0 0 0', '4 5 3 3 0', - '5 6 0 0 0', '5 5 3 0 0', '5 7 -1 0 0', '6 6 0 0 0', - '6 7 -1 0 0', '7' + '5 6 0 0 0', '5 5 3 0 0', '5 7 -1 -1 0', '6 6 0 0 0', + '6 7 -1 -1 0', '7' ]) actual_str_ragged0 = k2.to_str_simple(fsa_vec_ragged[0].to('cpu')) actual_str_ragged1 = k2.to_str_simple(fsa_vec_ragged[1].to('cpu')) From 12f591526a7afa3a3d55cc37c67576328cf926e5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ludwig=20K=C3=BCrzinger?= Date: Thu, 25 Nov 2021 00:54:45 +0100 Subject: [PATCH 2/4] Fix LICENSE location to k2 folder (#880) --- MANIFEST.in | 2 ++ setup.py | 1 - 2 files changed, 2 insertions(+), 1 deletion(-) create mode 100644 MANIFEST.in diff --git a/MANIFEST.in b/MANIFEST.in new file mode 100644 index 000000000..19ba5c47a --- /dev/null +++ b/MANIFEST.in @@ -0,0 +1,2 @@ +include LICENSE + diff --git a/setup.py b/setup.py index c45995b27..c0c4f52e0 100644 --- a/setup.py +++ b/setup.py @@ -215,7 +215,6 @@ def get_short_description(): packages=['k2', 'k2.ragged', 'k2.sparse', 'k2.version'], install_requires=install_requires, extras_require={'dev': dev_requirements}, - data_files=[('', ['LICENSE'])], ext_modules=[cmake_extension('_k2')], cmdclass={ 'build_ext': BuildExtension, From a0d75c8222a768adc1c68164f7c0730b621fff24 Mon Sep 17 00:00:00 2001 From: Fangjun Kuang Date: Mon, 29 Nov 2021 19:00:22 +0800 Subject: [PATCH 3/4] Release v1.11. (#881) It contains bugfixes. --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 47e0355c0..9f90d1d10 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -45,7 +45,7 @@ message(STATUS "Enabled languages: ${languages}") project(k2 ${languages}) -set(K2_VERSION "1.10") +set(K2_VERSION "1.11") # ----------------- Supported build types for K2 project ----------------- set(ALLOWABLE_BUILD_TYPES Debug Release RelWithDebInfo MinSizeRel) From 2cb3eeaa49b61d23f2917e5cb23ca03ef8ca709d Mon Sep 17 00:00:00 2001 From: Daniel Povey Date: Sun, 5 Dec 2021 18:00:48 +0800 Subject: [PATCH 4/4] Update documentation for hash.h (#887) * Update documentation for hash.h * Typo fix --- k2/csrc/hash.h | 63 +++++++++++++++++++++++++++++++++++--------------- 1 file changed, 44 insertions(+), 19 deletions(-) diff --git a/k2/csrc/hash.h b/k2/csrc/hash.h index 48459accb..ae651c6de 100644 --- a/k2/csrc/hash.h +++ b/k2/csrc/hash.h @@ -55,14 +55,43 @@ unsigned long long int __forceinline__ __host__ __device__ AtomicCAS( How class Hash works: - It can function as a map from key=uint32_t to value=uint32_t, or from - key=uint64_t to value=uint64_t where you choose NUM_KEY_BITS and - `key` must have only up to NUM_KEY_BITS set and `value` must have - only up to (64-NUM_KEY_BITS) set. You decide NUM_KEY_BITS when - you call Hash::Accessor() - - You can store any (key,value) pair except the pair where all the bits of + key=uint64_t to value=uint64_t, but you cannot use all 64 bits in the + key and value because we compress both of them into a single 64-bit + integer. There are several different modes of using this hash, + depending which accessor objects you use. The modes are: + + - Use Accessor with num_key_bits known at compile time; + the number of values bits will be 64 - NUM_KEY_BITS. + - Use GenericAccessor, which is like Accessor but the number of + key bits is not known at compile time; and they both must still + sum to 64. + - Use PackedAccessor, which allows you to have the number of key + plus value bits greater than 64; the rest of the bits are + implicit in groups of buckets (the number of buckets must + be >= 32 * 1 << (num_key_bits + num_value_bits - 64). + + - You must decide the number of key and value bits, and the number of + buckets, when you create the hash, but you can resize it (manually) + and when you resize it you can change the number of key and value bits. + + Some constraints: + - You can store any (key,value) pair allowed by the number of key and value + bits, except the pair where all the bits of both and key and value are set [that is used to mean "nothing here"] - - The number of buckets is a power of 2 provided by the user to the constructor; - currently no resizing is supported. + - The number of buckets must always be a power of 2. + - When deleting values from the hash you must delete them all at + once (necessary because there is no concept of a "tombstone". + + Some notes on usage: + + You use it by: constructing it, obtaining its Accessor with GetAccessor() + with appropriate template args depending on your chosen accessor type; and + inside kernels (or host code), calling functions Insert(), Find() or Delete() + of the Accessor object. Resizing is not automatic; it is the user's + responsibility to make sure the hash does not get too full (which could cause + assertion failures in kernels, and will be very slow). + + Some implementation notes: - When accessing hash[key], we use bucket_index == key % num_buckets, bucket_inc = 1 | (((key * 2) / num_buckets) ^ key). - If the bucket at `bucket_index` is occupied, we look in locations @@ -72,15 +101,7 @@ unsigned long long int __forceinline__ __host__ __device__ AtomicCAS( being odd ensures we eventually try all locations (of course for reasonable hash occupancy levels, we shouldn't ever have to try more than two or three). - - When deleting values from the hash you must delete them all at - once (necessary because there is no concept of a "tombstone". - You use it by: constructing it, obtaining its Accessor with - GetAccessor(), and inside kernels (or host code), calling - functions Insert(), Find() or Delete() of the Accessor object. There is no - resizing; sizing it correctly is the caller's responsibility and if the hash - gets full the code will just loop forever (of course it will get extremely - slow before it reaches that point). */ class Hash { public: @@ -94,10 +115,14 @@ class Hash { @param [in] num_key_bits Number of bits in the key of the hash; must satisfy 0 < num_key_bits < 64, and keys used must be less than (1< or GenericAccessor, + we require that num_key_bits + num_value_bits == 64. + For PackedAccessor we allow that num_key_bits + num_value_bits > 64, + but with the constraint that + (num_buckets >> (64 - num_key_bits - num_value_bits)) >= 32 */ Hash(ContextPtr c, int32_t num_buckets,