| /************************************************************************************\ |
| * * |
| * Copyright � 2014 Advanced Micro Devices, Inc. * |
| * Copyright (c) 2015 Mark D. Hill and David A. Wood * |
| * Copyright (c) 2021 Gaurav Jain and Matthew D. Sinclair * |
| * All rights reserved. * |
| * * |
| * Redistribution and use in source and binary forms, with or without * |
| * modification, are permitted provided that the following are met: * |
| * * |
| * You must reproduce the above copyright notice. * |
| * * |
| * 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 from at least the copyright holder. * |
| * * |
| * You must include the following terms in your license and/or other materials * |
| * provided with the software. * |
| * * |
| * 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, NON-INFRINGEMENT, 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. * |
| * * |
| * Without limiting the foregoing, the software may implement third party * |
| * technologies for which you must obtain licenses from parties other than AMD. * |
| * You agree that AMD has not obtained or conveyed to you, and that you shall * |
| * be responsible for obtaining the rights to use and/or distribute the applicable * |
| * underlying intellectual property rights related to the third party technologies. * |
| * These third party technologies are not licensed hereunder. * |
| * * |
| * If you use the software (in whole or in part), you shall adhere to all * |
| * applicable U.S., European, and other export laws, including but not limited to * |
| * the U.S. Export Administration Regulations ("EAR") (15 C.F.R Sections 730-774), * |
| * and E.U. Council Regulation (EC) No 428/2009 of 5 May 2009. Further, pursuant * |
| * to Section 740.6 of the EAR, you hereby certify that, except pursuant to a * |
| * license granted by the United States Department of Commerce Bureau of Industry * |
| * and Security or as otherwise permitted pursuant to a License Exception under * |
| * the U.S. Export Administration Regulations ("EAR"), you will not (1) export, * |
| * re-export or release to a national of a country in Country Groups D:1, E:1 or * |
| * E:2 any restricted technology, software, or source code you receive hereunder, * |
| * or (2) export to Country Groups D:1, E:1 or E:2 the direct product of such * |
| * technology or software, if such foreign produced direct product is subject to * |
| * national security controls as identified on the Commerce Control List (currently * |
| * found in Supplement 1 to Part 774 of EAR). For the most current Country Group * |
| * listings, or for additional information about the EAR or your obligations under * |
| * those regulations, please refer to the U.S. Bureau of Industry and Security's * |
| * website at http://www.bis.doc.gov/. * |
| * * |
| \************************************************************************************/ |
| |
| #ifndef KERNEL_H_ |
| #define KERNEL_H_ |
| |
| #include "hip/hip_runtime.h" |
| |
| /** |
| * @brief Breadth-first traversal |
| * @param row CSR pointer array |
| * @param col CSR column array |
| * @param d Distance array |
| * @param rho Rho array |
| * @param p Dependency array |
| * @param cont Termination variable |
| * @param num_nodes Termination variable |
| * @param num_edges Termination variable |
| * @param dist Current traversal layer |
| */ |
| |
| __global__ void |
| bfs_kernel(int *row, int *col, int *d, float *rho, int *cont, |
| const int num_nodes, const int num_edges, const int dist) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| //navigate the current layer |
| if (tid < num_nodes && d[tid] == dist) { |
| |
| //get the starting and ending pointers |
| //of the neighbor list |
| |
| int start = row[tid]; |
| int end; |
| if (tid + 1 < num_nodes) |
| end = row[tid + 1]; |
| else |
| end = num_edges; |
| |
| //navigate through the neighbor list |
| for (int edge = start; edge < end; edge++) { |
| int w = col[edge]; |
| if (d[w] < 0) { |
| *cont = 1; |
| //traverse another layer |
| d[w] = dist + 1; |
| } |
| //transfer the rho value to the neighbor |
| if (d[w] == (dist + 1)) { |
| atomicAdd(&rho[w], rho[tid]); |
| } |
| } |
| } |
| } |
| |
| /** |
| * @brief Back traversal |
| * @param row CSR pointer array |
| * @param col CSR column array |
| * @param d Distance array |
| * @param rho Rho array |
| * @param sigma Sigma array |
| * @param p Dependency array |
| * @param cont Termination variable |
| * @param num_nodes Termination variable |
| * @param num_edges Termination variable |
| * @param dist Current traversal layer |
| * @param s Source vertex |
| * @param bc Betweeness Centrality array |
| */ |
| |
| __global__ void |
| backtrack_kernel(int *row, int *col, int *d, float *rho, float *sigma, |
| const int num_nodes, const int num_edges, const int dist, |
| const int s, float* bc) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| // Navigate the current layer |
| if (tid < num_nodes && d[tid] == dist - 1) { |
| |
| int start = row[tid]; |
| int end; |
| if (tid + 1 < num_nodes) |
| end = row[tid + 1]; |
| else |
| end = num_edges; |
| |
| // Get the starting and ending pointers |
| // of the neighbor list in the reverse graph |
| for (int edge = start; edge < end; edge++) { |
| int w = col[edge]; |
| // Update the sigma value traversing back |
| if (d[w] == dist - 2) |
| atomicAdd(&sigma[w], rho[w] / rho[tid] * (1 + sigma[tid])); |
| } |
| |
| // Update the BC value |
| if (tid != s) |
| bc[tid] = bc[tid] + sigma[tid]; |
| } |
| |
| } |
| |
| /** |
| * @brief back_sum_kernel (not used) |
| * @param s Source vertex |
| * @param dist Current traversal layer |
| * @param d Distance array |
| * @param sigma Sigma array |
| * @param bc Betweeness Centrality array |
| * @param num_nodes Termination variable |
| * @param num_edges Termination variable |
| */ |
| __global__ void |
| back_sum_kernel(const int s, const int dist, int *d, float *sigma, float *bc, |
| const int num_nodes) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| if (tid < num_nodes) { |
| // If it is not the source |
| if (s != tid && d[tid] == dist - 1) { |
| bc[tid] = bc[tid] + sigma[tid]; |
| } |
| } |
| } |
| |
| /** |
| * @brief array set 1D |
| * @param s Source vertex |
| * @param dist_array Distance array |
| * @param sigma Sigma array |
| * @param rho Rho array |
| * @param num_nodes Termination variable |
| */ |
| __global__ void |
| clean_1d_array(const int source, int *dist_array, float *sigma, float *rho, |
| const int num_nodes) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| if (tid < num_nodes) { |
| |
| sigma[tid] = 0; |
| |
| if (tid == source) { |
| // If source vertex rho = 1, dist = 0 |
| rho[tid] = 1; |
| dist_array[tid] = 0; |
| } else { |
| // If other vertices rho = 0, dist = -1 |
| rho[tid] = 0; |
| dist_array[tid] = -1; |
| } |
| } |
| } |
| |
| /** |
| * @brief array set 2D |
| * @param p Dependency array |
| * @param num_nodes Number of vertices |
| */ |
| __global__ void clean_2d_array(int *p, const int num_nodes) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| if (tid < num_nodes * num_nodes) |
| p[tid] = 0; |
| } |
| |
| /** |
| * @brief clean BC |
| * @param bc_d Betweeness Centrality array |
| * @param num_nodes Number of vertices |
| */ |
| __global__ void clean_bc(float *bc_d, const int num_nodes) |
| { |
| int tid = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x; |
| |
| if (tid < num_nodes) |
| bc_d[tid] = 0; |
| } |
| |
| #endif // KERNEL_H_ |