1 /**
2 * Copyright 2020 Huawei Technologies Co., Ltd
3 *
4 * Licensed under the Apache License, Version 2.0 (the "License");
5 * you may not use this file except in compliance with the License.
6 * You may obtain a copy of the License at
7 *
8 * http://www.apache.org/licenses/LICENSE-2.0
9 *
10 * Unless required by applicable law or agreed to in writing, software
11 * distributed under the License is distributed on an "AS IS" BASIS,
12 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 * See the License for the specific language governing permissions and
14 * limitations under the License.
15 */
16
17 #include "hash_impl.cuh"
18 #include "runtime/device/gpu/cuda_common.h"
19
20 template <typename T>
HashSwapOut(const T * hash_table,T * swap_out_value,const int * swap_out_index,const int index_size,const int hash_dim)21 __global__ void HashSwapOut(const T *hash_table, T *swap_out_value, const int *swap_out_index, const int index_size,
22 const int hash_dim) {
23 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < index_size; i += blockDim.x * gridDim.x) {
24 int hash_index = swap_out_index[i];
25 for (int j = 0; j < hash_dim; j++) {
26 swap_out_value[i * hash_dim + j] = hash_table[hash_index * hash_dim + j];
27 }
28 }
29 return;
30 }
31
32 template <typename T>
HashSwapIn(T * hash_table,const T * swap_in_value,const int * swap_in_index,const int index_size,const int hash_dim)33 __global__ void HashSwapIn(T *hash_table, const T *swap_in_value, const int *swap_in_index, const int index_size,
34 const int hash_dim) {
35 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < index_size; i += blockDim.x * gridDim.x) {
36 int hash_index = swap_in_index[i];
37 for (int j = 0; j < hash_dim; j++) {
38 hash_table[hash_index * hash_dim + j] = swap_in_value[i * hash_dim + j];
39 }
40 }
41 return;
42 }
43
44 template <typename T>
DoHashSwapOut(const T * hash_table,T * swap_out_value,const int * swap_out_index,const int index_size,const int hash_dim,cudaStream_t cuda_stream)45 void DoHashSwapOut(const T *hash_table, T *swap_out_value, const int *swap_out_index, const int index_size,
46 const int hash_dim, cudaStream_t cuda_stream) {
47 HashSwapOut<<<GET_BLOCKS(index_size), GET_THREADS, 0, cuda_stream>>>(hash_table, swap_out_value, swap_out_index,
48 index_size, hash_dim);
49 return;
50 }
51
52 template <typename T>
DoHashSwapIn(T * hash_table,const T * swap_in_value,const int * swap_in_index,const int index_size,const int hash_dim,cudaStream_t cuda_stream)53 void DoHashSwapIn(T *hash_table, const T *swap_in_value, const int *swap_in_index, const int index_size,
54 const int hash_dim, cudaStream_t cuda_stream) {
55 HashSwapIn<<<GET_BLOCKS(index_size), GET_THREADS, 0, cuda_stream>>>(hash_table, swap_in_value, swap_in_index,
56 index_size, hash_dim);
57 return;
58 }
59
60 template void DoHashSwapOut<float>(const float *hash_table, float *swap_out_value, const int *swap_out_index,
61 const int index_size, const int hash_dim, cudaStream_t cuda_stream);
62
63 template void DoHashSwapIn<float>(float *hash_table, const float *swap_in_value, const int *swap_in_index,
64 const int index_size, const int hash_dim, cudaStream_t cuda_stream);
65