chore: import upstream snapshot with attribution
This commit is contained in:
wehub-resource-sync
@@ -0,0 +1,24 @@
|
||||
cmake_minimum_required(VERSION 3.23)
|
||||
project(gpu_easygraph LANGUAGES CXX CUDA)
|
||||
|
||||
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
|
||||
|
||||
file(GLOB SOURCES
|
||||
functions/*/*.c
|
||||
functions/*/*.cpp
|
||||
functions/*/*.cu
|
||||
common/*.c
|
||||
common/*.cpp
|
||||
)
|
||||
|
||||
add_library(gpu_easygraph OBJECT ${SOURCES})
|
||||
|
||||
target_include_directories(gpu_easygraph
|
||||
PRIVATE common
|
||||
PRIVATE functions
|
||||
)
|
||||
|
||||
set_target_properties(gpu_easygraph PROPERTIES
|
||||
LINK_SEARCH_START_STATIC ON
|
||||
LINK_SEARCH_END_STATIC ON
|
||||
)
|
||||
@@ -0,0 +1,23 @@
|
||||
#pragma once
|
||||
|
||||
#include "err.h"
|
||||
|
||||
#define EG_DOUBLE_INF 1e100
|
||||
|
||||
#define EXIT_IF_CUDA_FAILED(condition) \
|
||||
cuda_ret = condition; \
|
||||
if (cuda_ret != cudaSuccess) { \
|
||||
goto exit; \
|
||||
} \
|
||||
|
||||
#ifndef _IN_
|
||||
#define _IN_
|
||||
#endif
|
||||
|
||||
#ifndef _OUT_
|
||||
#define _OUT_
|
||||
#endif
|
||||
|
||||
#ifndef _BUFFER_
|
||||
#define _BUFFER_
|
||||
#endif
|
||||
@@ -0,0 +1,31 @@
|
||||
#include <string>
|
||||
|
||||
#include "err.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::string;
|
||||
|
||||
std::string err_code_detail(
|
||||
int status
|
||||
) {
|
||||
switch (status) {
|
||||
case EG_GPU_SUCC:
|
||||
return "EasyGraph GPU: success";
|
||||
case EG_GPU_FAILED_TO_ALLOCATE_HOST_MEM:
|
||||
return "EasyGraph GPU: failed to allocate host mem";
|
||||
case EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM:
|
||||
return "EasyGraph GPU: failed to allocate gpu mem";
|
||||
case EG_GPU_DEVICE_ERR:
|
||||
return "EasyGraph GPU: gpu error occurred";
|
||||
case EG_GPU_UNKNOW_ERROR:
|
||||
return "EasyGraph GPU: gpu unkonw error";
|
||||
case EG_UNSUPPORTED_GRAPH:
|
||||
return "EasyGraph GPU: unsupported graph type";
|
||||
default:
|
||||
break;
|
||||
}
|
||||
return "EasyGraph GPU: not a valid err_code";
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,20 @@
|
||||
#pragma once
|
||||
|
||||
#include <string>
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
typedef enum {
|
||||
EG_GPU_SUCC = 0,
|
||||
EG_GPU_FAILED_TO_ALLOCATE_HOST_MEM,
|
||||
EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM,
|
||||
EG_GPU_DEVICE_ERR,
|
||||
EG_GPU_UNKNOW_ERROR,
|
||||
EG_UNSUPPORTED_GRAPH
|
||||
} EG_GPU_STATUS_CODE;
|
||||
|
||||
std::string err_code_detail(
|
||||
int status
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,423 @@
|
||||
//
|
||||
// This file is sourcing from here: https://peerj.com/articles/cs-140/
|
||||
// Something such as vars' name, graph format, etc were changed
|
||||
// for adapting easygraph's GPU framework
|
||||
//
|
||||
|
||||
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __device__ double atomicAddDouble (
|
||||
_OUT_ double* address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long int* address_as_ull =
|
||||
(unsigned long long int*)address;
|
||||
unsigned long long int old = *address_as_ull, assumed;
|
||||
do {
|
||||
assumed = old;
|
||||
old = atomicCAS(address_as_ull, assumed,
|
||||
__double_as_longlong(val +
|
||||
__longlong_as_double(assumed)));
|
||||
} while (assumed != old);
|
||||
return __longlong_as_double(old);
|
||||
}
|
||||
|
||||
static __device__ double atomicMinDouble (
|
||||
_OUT_ double *address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long ret = __double_as_longlong(*address);
|
||||
while (val < __longlong_as_double(ret))
|
||||
{
|
||||
unsigned long long old = ret;
|
||||
if ((ret = atomicCAS((unsigned long long *)address, old, __double_as_longlong(val))) == old)
|
||||
break;
|
||||
}
|
||||
return __longlong_as_double(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_calc_min_edge (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ double* d_min_edge
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid < len_V) {
|
||||
double curr_min = EG_DOUBLE_INF;
|
||||
int edge_start = d_V[tid];
|
||||
int edge_end = d_V[tid + 1];
|
||||
for(int i = edge_start; i < edge_end; ++i) {
|
||||
curr_min = min(curr_min, d_W[i]);
|
||||
}
|
||||
d_min_edge[tid] = curr_min;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
static __global__ void d_dijkstra_bc (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ double* d_min_edge,
|
||||
_IN_ int* d_sources,
|
||||
_BUFFER_ double* d_dist_2D,
|
||||
_BUFFER_ double* d_sigma_2D,
|
||||
_BUFFER_ double* d_delta_2D,
|
||||
_BUFFER_ int* d_U_2D,
|
||||
_BUFFER_ int* d_F_2D,
|
||||
_BUFFER_ int* d_lock_flag_2D,
|
||||
_BUFFER_ int* d_st_2D,
|
||||
_BUFFER_ int* d_st_idx_2D,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
for (int s_idx = blockIdx.x; s_idx < len_sources; s_idx += gridDim.x) {
|
||||
int s = d_sources[s_idx];
|
||||
|
||||
double* d_dist = d_dist_2D + blockIdx.x * len_V;
|
||||
double* d_sigma = d_sigma_2D + blockIdx.x * len_V;
|
||||
double* d_delta = d_delta_2D + blockIdx.x * len_V;
|
||||
|
||||
int* d_U = d_U_2D + blockIdx.x * len_V;
|
||||
int* d_F = d_F_2D + blockIdx.x * len_V;
|
||||
int* d_lock_flag = d_lock_flag_2D + blockIdx.x * len_V;
|
||||
int* d_st = d_st_2D + blockIdx.x * len_V;
|
||||
int* d_st_idx = d_st_idx_2D + blockIdx.x * (len_V + 2);
|
||||
|
||||
__shared__ int len_F;
|
||||
__shared__ int len_st;
|
||||
__shared__ int len_st_idx;
|
||||
__shared__ double delta;
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
d_dist[i] = EG_DOUBLE_INF;
|
||||
d_sigma[i] = 0;
|
||||
d_delta[i] = 0;
|
||||
|
||||
d_U[i] = 1;
|
||||
d_lock_flag[i] = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_dist[s] = 0;
|
||||
d_sigma[s] = 1;
|
||||
|
||||
d_U[s] = 0;
|
||||
d_F[0] = s;
|
||||
len_F = 1;
|
||||
d_st[0] = s;
|
||||
len_st = 1;
|
||||
d_st_idx[0] = 0;
|
||||
d_st_idx[1] = 1;
|
||||
len_st_idx = 2;
|
||||
|
||||
delta = 0.0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
int needlock = 1;
|
||||
while (delta < EG_DOUBLE_INF) {
|
||||
for (int j = threadIdx.x; j < len_F * warp_size; j += blockDim.x) {
|
||||
int f = d_F[j / warp_size];
|
||||
int edge_start = d_V[f];
|
||||
int edge_end = d_V[f + 1];
|
||||
double dist = d_dist[f];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
double relax_w = dist + d_W[e + edge_start];
|
||||
needlock = 1;
|
||||
while (needlock) {
|
||||
if (atomicCAS(d_lock_flag + adj, 0, 1) == 0) {
|
||||
if (relax_w < d_dist[adj]) {
|
||||
d_dist[adj] = relax_w;
|
||||
d_sigma[adj] = 0;
|
||||
}
|
||||
if (d_dist[adj] == relax_w) {
|
||||
d_sigma[adj] += d_sigma[f];
|
||||
}
|
||||
atomicExch(d_lock_flag + adj, 0);
|
||||
needlock = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
delta = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] == 1 && dist_i < EG_DOUBLE_INF) {
|
||||
atomicMinDouble(&delta, dist_i + d_min_edge[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
len_F = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] && dist_i < delta && dist_i < EG_DOUBLE_INF) {
|
||||
d_U[i] = 0;
|
||||
int f_idx = atomicAdd(&len_F, 1);
|
||||
d_F[f_idx] = i;
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_F; i += blockDim.x) {
|
||||
int st_idx = atomicAdd(&len_st, 1);
|
||||
d_st[st_idx] = d_F[i];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_st_idx[len_st_idx] = d_st_idx[len_st_idx - 1] + len_F;
|
||||
++len_st_idx;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
__shared__ int depth, st_start, st_end;
|
||||
if (threadIdx.x == 0) {
|
||||
depth = len_st_idx - 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0 && endpoints) {
|
||||
atomicAddDouble(d_BC + s, d_st_idx[depth] - 1);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (depth > 0) {
|
||||
if (threadIdx.x == 0) {
|
||||
st_start = d_st_idx[depth - 1];
|
||||
st_end = d_st_idx[depth];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int j = threadIdx.x; j < (st_end - st_start) * warp_size; j += blockDim.x) {
|
||||
int pred = d_st[st_start + j / warp_size];
|
||||
int edge_start = d_V[pred];
|
||||
int edge_end = d_V[pred + 1];
|
||||
double pred_sigma = d_sigma[pred];
|
||||
double pred_dist = d_dist[pred];
|
||||
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int succ = d_E[e + edge_start];
|
||||
double weight = d_W[e + edge_start];
|
||||
double succ_dist = d_dist[succ];
|
||||
if (succ_dist == pred_dist + weight) {
|
||||
atomicAddDouble(d_delta + pred,
|
||||
pred_sigma / d_sigma[succ] * (1 + d_delta[succ]));
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
|
||||
if (j % warp_size == 0 && s != pred) {
|
||||
atomicAddDouble(d_BC + pred, d_delta[pred] + endpoints);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
--depth;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_rescale(
|
||||
_IN_ int len_V,
|
||||
_IN_ double scale,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
|
||||
if (tid < len_V) {
|
||||
d_BC[tid] *= scale;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static double calc_scale(
|
||||
_IN_ int len_V,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints
|
||||
)
|
||||
{
|
||||
double scale = 1.0;
|
||||
if (normalized) {
|
||||
if (endpoints) {
|
||||
if (len_V < 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / (double(len_V) * (len_V - 1));
|
||||
}
|
||||
} else if (len_V <= 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / ((double(len_V) - 1) * (len_V - 2));
|
||||
}
|
||||
} else {
|
||||
if (!is_directed) {
|
||||
scale = 0.5;
|
||||
} else {
|
||||
scale = 1.0;
|
||||
}
|
||||
}
|
||||
return scale;
|
||||
}
|
||||
|
||||
|
||||
|
||||
int cuda_betweenness_centrality (
|
||||
_IN_ int* V,
|
||||
_IN_ int* E,
|
||||
_IN_ double* W,
|
||||
_IN_ int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* BC
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int block_size = 256;
|
||||
size_t grid_size = len_V / block_size + (len_V % block_size != 0);
|
||||
size_t mem_free = 0, mem_total = 0;
|
||||
|
||||
double scale = calc_scale(len_V, is_directed, normalized, endpoints);
|
||||
|
||||
int *d_V = NULL, *d_E = NULL, *d_sources= NULL, *d_lock_flag_2D = NULL;
|
||||
int *d_U_2D = NULL, *d_F_2D = NULL, *d_st_2D = NULL, *d_st_idx_2D = NULL;
|
||||
double *d_W = NULL, *d_min_edge = NULL, *d_dist_2D = NULL,
|
||||
*d_sigma_2D = NULL, *d_delta_2D = NULL, *d_BC = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemGetInfo(&mem_free, &mem_total));
|
||||
while (true) {
|
||||
size_t mem_needed = sizeof(int) * len_V // d_V
|
||||
+ sizeof(int) * len_E // d_E
|
||||
+ sizeof(int) * len_sources // d_sources
|
||||
+ sizeof(int) * grid_size * len_V // d_lock_flag_2D
|
||||
+ sizeof(int) * grid_size * len_V // d_U_2D
|
||||
+ sizeof(int) * grid_size * len_V // d_F_2D
|
||||
+ sizeof(int) * grid_size * len_V // d_st_2D
|
||||
+ sizeof(int) * grid_size * (len_V + 2) // d_st_idx_2D
|
||||
+ sizeof(double) * len_E // d_W
|
||||
+ sizeof(double) * len_V // d_min_edge
|
||||
+ sizeof(double) * grid_size * len_V // d_dist_2D
|
||||
+ sizeof(double) * grid_size * len_V // d_sigma_2D
|
||||
+ sizeof(double) * grid_size * len_V // d_delta_2D
|
||||
+ sizeof(double) * len_V // d_BC
|
||||
;
|
||||
if (mem_needed < mem_free / 2) {
|
||||
break;
|
||||
} else {
|
||||
grid_size /= 2;
|
||||
}
|
||||
}
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sources, sizeof(int) * len_sources));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_lock_flag_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_U_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_F_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_idx_2D, sizeof(int) * grid_size * (len_V + 2)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, sizeof(double) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_min_edge, sizeof(double) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_dist_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sigma_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_delta_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_BC, sizeof(double) * len_V));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * len_V, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_sources, sources, sizeof(int) * len_sources, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, sizeof(double) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
d_calc_min_edge<<<grid_size, block_size>>>(d_V, d_E, d_W, len_V, len_E, d_min_edge);
|
||||
|
||||
d_dijkstra_bc<<<grid_size, block_size>>>(d_V, d_E, d_W, d_min_edge, d_sources, d_dist_2D, d_sigma_2D,
|
||||
d_delta_2D, d_U_2D, d_F_2D, d_lock_flag_2D, d_st_2D,
|
||||
d_st_idx_2D, len_V, len_E, len_sources, warp_size,
|
||||
endpoints, d_BC);
|
||||
|
||||
if (scale != 1.0) {
|
||||
d_rescale<<<grid_size, block_size>>>(len_V, scale, d_BC);
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(BC, d_BC, sizeof(double) * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_sources);
|
||||
cudaFree(d_lock_flag_2D);
|
||||
cudaFree(d_U_2D);
|
||||
cudaFree(d_F_2D);
|
||||
cudaFree(d_st_2D);
|
||||
cudaFree(d_st_idx_2D);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_min_edge);
|
||||
cudaFree(d_dist_2D);
|
||||
cudaFree(d_sigma_2D);
|
||||
cudaFree(d_delta_2D);
|
||||
cudaFree(d_BC);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,399 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __device__ double atomicAddDouble (
|
||||
_OUT_ double* address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long int* address_as_ull =
|
||||
(unsigned long long int*)address;
|
||||
unsigned long long int old = *address_as_ull, assumed;
|
||||
do {
|
||||
assumed = old;
|
||||
old = atomicCAS(address_as_ull, assumed,
|
||||
__double_as_longlong(val +
|
||||
__longlong_as_double(assumed)));
|
||||
} while (assumed != old);
|
||||
return __longlong_as_double(old);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __device__ double atomicMinDouble (
|
||||
_OUT_ double *address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long ret = __double_as_longlong(*address);
|
||||
while (val < __longlong_as_double(ret))
|
||||
{
|
||||
unsigned long long old = ret;
|
||||
if ((ret = atomicCAS((unsigned long long *)address, old, __double_as_longlong(val))) == old)
|
||||
break;
|
||||
}
|
||||
return __longlong_as_double(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_calc_min_edge (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ double* d_min_edge
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (tid < len_V) {
|
||||
double curr_min = EG_DOUBLE_INF;
|
||||
int edge_start = d_V[tid];
|
||||
int edge_end = d_V[tid + 1];
|
||||
for(int i = edge_start; i < edge_end; ++i) {
|
||||
curr_min = min(curr_min, d_W[i]);
|
||||
}
|
||||
d_min_edge[tid] = curr_min;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_dijkstra_bc (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ double* d_min_edge,
|
||||
_IN_ int* d_sources,
|
||||
_BUFFER_ double* d_dist_2D,
|
||||
_BUFFER_ double* d_sigma_2D,
|
||||
_BUFFER_ double* d_delta_2D,
|
||||
_BUFFER_ int* d_U_2D,
|
||||
_BUFFER_ int* d_F_2D,
|
||||
_BUFFER_ int* d_st_2D,
|
||||
_BUFFER_ int* d_st_idx_2D,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
for (int s_idx = blockIdx.x; s_idx < len_sources; s_idx += gridDim.x) {
|
||||
int s = d_sources[s_idx];
|
||||
|
||||
double* d_dist = d_dist_2D + blockIdx.x * len_V;
|
||||
double* d_sigma = d_sigma_2D + blockIdx.x * len_V;
|
||||
double* d_delta = d_delta_2D + blockIdx.x * len_V;
|
||||
|
||||
int* d_U = d_U_2D + blockIdx.x * len_V;
|
||||
int* d_F = d_F_2D + blockIdx.x * len_V;
|
||||
int* d_st = d_st_2D + blockIdx.x * len_V;
|
||||
int* d_st_idx = d_st_idx_2D + blockIdx.x * (len_V + 2);
|
||||
|
||||
__shared__ int len_F;
|
||||
__shared__ int len_st;
|
||||
__shared__ int len_st_idx;
|
||||
__shared__ double delta;
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
d_dist[i] = EG_DOUBLE_INF;
|
||||
d_sigma[i] = 0;
|
||||
d_delta[i] = 0;
|
||||
|
||||
d_U[i] = 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_dist[s] = 0;
|
||||
d_sigma[s] = 1;
|
||||
|
||||
d_U[s] = 0;
|
||||
d_F[0] = s;
|
||||
len_F = 1;
|
||||
d_st[0] = s;
|
||||
len_st = 1;
|
||||
d_st_idx[0] = 0;
|
||||
d_st_idx[1] = 1;
|
||||
len_st_idx = 2;
|
||||
|
||||
delta = 0.0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (delta < EG_DOUBLE_INF) {
|
||||
for (int j = threadIdx.x; j < len_F * warp_size; j += blockDim.x) {
|
||||
int f = d_F[j / warp_size];
|
||||
int edge_start = d_V[f];
|
||||
int edge_end = d_V[f + 1];
|
||||
double dist = d_dist[f];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
double relax_w = dist + d_W[e + edge_start];
|
||||
atomicMinDouble(d_dist + adj, relax_w);
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
delta = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] == 1 && dist_i < EG_DOUBLE_INF) {
|
||||
atomicMinDouble(&delta, dist_i + d_min_edge[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
len_F = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] && dist_i < delta && dist_i < EG_DOUBLE_INF) {
|
||||
d_U[i] = 0;
|
||||
int f_idx = atomicAdd(&len_F, 1);
|
||||
d_F[f_idx] = i;
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_F; i += blockDim.x) {
|
||||
int st_idx = atomicAdd(&len_st, 1);
|
||||
d_st[st_idx] = d_F[i];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_st_idx[len_st_idx] = d_st_idx[len_st_idx - 1] + len_F;
|
||||
++len_st_idx;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
// calculate single source shortest path END
|
||||
|
||||
// calculate sigma START
|
||||
for (int curr_lvl = 0; curr_lvl + 1 < len_st_idx; ++curr_lvl) {
|
||||
int lvl_start = d_st_idx[curr_lvl];
|
||||
int lvl_end = d_st_idx[curr_lvl + 1];
|
||||
for (int j = threadIdx.x; j < (lvl_end - lvl_start) * warp_size; j += blockDim.x) {
|
||||
int v = d_st[lvl_start + j / warp_size];
|
||||
double dist_v = d_dist[v];
|
||||
int edge_start = d_V[v];
|
||||
int edge_end = d_V[v + 1];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
if (dist_v + d_W[e + edge_start] == d_dist[adj]) {
|
||||
atomicAddDouble(d_sigma + adj, d_sigma[v]);
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
// calculate sigma END
|
||||
|
||||
__shared__ int depth, st_start, st_end;
|
||||
if (threadIdx.x == 0) {
|
||||
depth = len_st_idx - 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0 && endpoints) {
|
||||
atomicAddDouble(d_BC + s, d_st_idx[depth] - 1);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (depth > 0) {
|
||||
if (threadIdx.x == 0) {
|
||||
st_start = d_st_idx[depth - 1];
|
||||
st_end = d_st_idx[depth];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int j = threadIdx.x; j < (st_end - st_start) * warp_size; j += blockDim.x) {
|
||||
int pred = d_st[st_start + j / warp_size];
|
||||
int edge_start = d_V[pred];
|
||||
int edge_end = d_V[pred + 1];
|
||||
double pred_sigma = d_sigma[pred];
|
||||
double pred_dist = d_dist[pred];
|
||||
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int succ = d_E[e + edge_start];
|
||||
double weight = d_W[e + edge_start];
|
||||
double succ_dist = d_dist[succ];
|
||||
if (succ_dist == pred_dist + weight) {
|
||||
atomicAddDouble(d_delta + pred,
|
||||
pred_sigma / d_sigma[succ] * (1 + d_delta[succ]));
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
|
||||
if (j % warp_size == 0 && s != pred) {
|
||||
atomicAddDouble(d_BC + pred, d_delta[pred] + endpoints);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
--depth;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_rescale(
|
||||
_IN_ int len_V,
|
||||
_IN_ double scale,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
int tid = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
|
||||
if (tid < len_V) {
|
||||
d_BC[tid] *= scale;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static double calc_scale(
|
||||
_IN_ int len_V,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints
|
||||
)
|
||||
{
|
||||
double scale = 1.0;
|
||||
if (normalized) {
|
||||
if (endpoints) {
|
||||
if (len_V < 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / (double(len_V) * (len_V - 1));
|
||||
}
|
||||
} else if (len_V <= 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / ((double(len_V) - 1) * (len_V - 2));
|
||||
}
|
||||
} else {
|
||||
if (!is_directed) {
|
||||
scale = 0.5;
|
||||
} else {
|
||||
scale = 1.0;
|
||||
}
|
||||
}
|
||||
return scale;
|
||||
}
|
||||
|
||||
|
||||
|
||||
int cuda_betweenness_centrality (
|
||||
_IN_ int* V,
|
||||
_IN_ int* E,
|
||||
_IN_ double* W,
|
||||
_IN_ int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* BC
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int block_size = 256;
|
||||
size_t grid_size = len_V / block_size + (len_V % block_size != 0);
|
||||
|
||||
double scale = calc_scale(len_V, is_directed, normalized, endpoints);
|
||||
|
||||
int *d_V = NULL, *d_E = NULL, *d_sources= NULL;
|
||||
int *d_U_2D = NULL, *d_F_2D = NULL, *d_st_2D = NULL, *d_st_idx_2D = NULL;
|
||||
double *d_W = NULL, *d_min_edge = NULL, *d_dist_2D = NULL,
|
||||
*d_sigma_2D = NULL, *d_delta_2D = NULL, *d_BC = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * (len_V + 1)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sources, sizeof(int) * len_sources));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_U_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_F_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_2D, sizeof(int) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_idx_2D, sizeof(int) * grid_size * (len_V + 2)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, sizeof(double) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_min_edge, sizeof(double) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_dist_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sigma_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_delta_2D, sizeof(double) * grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_BC, sizeof(double) * len_V));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * (len_V + 1), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_sources, sources, sizeof(int) * len_sources, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, sizeof(double) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
d_calc_min_edge<<<grid_size, block_size>>>(d_V, d_E, d_W, len_V, len_E, d_min_edge);
|
||||
|
||||
d_dijkstra_bc<<<grid_size, block_size>>>(d_V, d_E, d_W, d_min_edge, d_sources, d_dist_2D,
|
||||
d_sigma_2D, d_delta_2D, d_U_2D, d_F_2D, d_st_2D,
|
||||
d_st_idx_2D, len_V, len_E, len_sources, warp_size,
|
||||
endpoints, d_BC);
|
||||
|
||||
if (scale != 1.0) {
|
||||
d_rescale<<<grid_size, block_size>>>(len_V, scale, d_BC);
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(BC, d_BC, sizeof(double) * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_sources);
|
||||
cudaFree(d_U_2D);
|
||||
cudaFree(d_F_2D);
|
||||
cudaFree(d_st_2D);
|
||||
cudaFree(d_st_idx_2D);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_min_edge);
|
||||
cudaFree(d_dist_2D);
|
||||
cudaFree(d_sigma_2D);
|
||||
cudaFree(d_delta_2D);
|
||||
cudaFree(d_BC);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,426 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __device__ double atomicAddDouble (
|
||||
_OUT_ double* address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long int* address_as_ull =
|
||||
(unsigned long long int*)address;
|
||||
unsigned long long int old = *address_as_ull, assumed;
|
||||
do {
|
||||
assumed = old;
|
||||
old = atomicCAS(address_as_ull, assumed,
|
||||
__double_as_longlong(val +
|
||||
__longlong_as_double(assumed)));
|
||||
} while (assumed != old);
|
||||
return __longlong_as_double(old);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __device__ double atomicMinDouble (
|
||||
_OUT_ double *address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long ret = __double_as_longlong(*address);
|
||||
while (val < __longlong_as_double(ret))
|
||||
{
|
||||
unsigned long long old = ret;
|
||||
if ((ret = atomicCAS((unsigned long long *)address, old, __double_as_longlong(val))) == old)
|
||||
break;
|
||||
}
|
||||
return __longlong_as_double(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_calc_min_edge (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ double* d_min_edge
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int tnum = blockDim.x * gridDim.x;
|
||||
|
||||
for (int u = tid; u < len_V; u += tnum) {
|
||||
double curr_min = EG_DOUBLE_INF;
|
||||
int edge_start = d_V[u];
|
||||
int edge_end = d_V[u + 1];
|
||||
for(int v = edge_start; v < edge_end; ++v) {
|
||||
curr_min = min(curr_min, d_W[v]);
|
||||
}
|
||||
d_min_edge[u] = curr_min;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_dijkstra_bc (
|
||||
_IN_ int* d_curr_node,
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ double* d_min_edge,
|
||||
_IN_ int* d_sources,
|
||||
_BUFFER_ double* d_dist_2D,
|
||||
_BUFFER_ double* d_sigma_2D,
|
||||
_BUFFER_ double* d_delta_2D,
|
||||
_BUFFER_ int* d_U_2D,
|
||||
_BUFFER_ int* d_F_2D,
|
||||
_BUFFER_ int* d_st_2D,
|
||||
_BUFFER_ int* d_st_idx_2D,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
//for (int s_idx = blockIdx.x; s_idx < len_sources; s_idx += gridDim.x) {
|
||||
while (1) {
|
||||
__shared__ int curr_node;
|
||||
if (threadIdx.x == 0) {
|
||||
curr_node = atomicAdd(d_curr_node, 1);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (curr_node >= len_sources) {
|
||||
break;
|
||||
}
|
||||
|
||||
int s = d_sources[curr_node];
|
||||
|
||||
double* d_dist = d_dist_2D + blockIdx.x * len_V;
|
||||
double* d_sigma = d_sigma_2D + blockIdx.x * len_V;
|
||||
double* d_delta = d_delta_2D + blockIdx.x * len_V;
|
||||
|
||||
int* d_U = d_U_2D + blockIdx.x * len_V;
|
||||
int* d_F = d_F_2D + blockIdx.x * len_V;
|
||||
int* d_st = d_st_2D + blockIdx.x * len_V;
|
||||
int* d_st_idx = d_st_idx_2D + blockIdx.x * (len_V + 2);
|
||||
|
||||
__shared__ int len_F;
|
||||
__shared__ int len_st;
|
||||
__shared__ int len_st_idx;
|
||||
__shared__ double delta;
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
d_dist[i] = EG_DOUBLE_INF;
|
||||
d_sigma[i] = 0;
|
||||
d_delta[i] = 0;
|
||||
|
||||
d_U[i] = 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_dist[s] = 0;
|
||||
d_sigma[s] = 1;
|
||||
|
||||
d_U[s] = 0;
|
||||
d_F[0] = s;
|
||||
len_F = 1;
|
||||
d_st[0] = s;
|
||||
len_st = 1;
|
||||
d_st_idx[0] = 0;
|
||||
d_st_idx[1] = 1;
|
||||
len_st_idx = 2;
|
||||
|
||||
delta = 0.0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (delta < EG_DOUBLE_INF) {
|
||||
for (int j = threadIdx.x; j < len_F * warp_size; j += blockDim.x) {
|
||||
int f = d_F[j / warp_size];
|
||||
int edge_start = d_V[f];
|
||||
int edge_end = d_V[f + 1];
|
||||
double dist = d_dist[f];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
double relax_w = dist + d_W[e + edge_start];
|
||||
atomicMinDouble(d_dist + adj, relax_w);
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
delta = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] == 1 && dist_i < EG_DOUBLE_INF) {
|
||||
atomicMinDouble(&delta, dist_i + d_min_edge[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
len_F = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] && dist_i < delta && dist_i < EG_DOUBLE_INF) {
|
||||
d_U[i] = 0;
|
||||
int f_idx = atomicAdd(&len_F, 1);
|
||||
d_F[f_idx] = i;
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_F; i += blockDim.x) {
|
||||
int st_idx = atomicAdd(&len_st, 1);
|
||||
d_st[st_idx] = d_F[i];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_st_idx[len_st_idx] = d_st_idx[len_st_idx - 1] + len_F;
|
||||
++len_st_idx;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
// calculate single source shortest path END
|
||||
|
||||
// calculate sigma START
|
||||
for (int curr_lvl = 0; curr_lvl + 1 < len_st_idx; ++curr_lvl) {
|
||||
int lvl_start = d_st_idx[curr_lvl];
|
||||
int lvl_end = d_st_idx[curr_lvl + 1];
|
||||
for (int j = threadIdx.x; j < (lvl_end - lvl_start) * warp_size; j += blockDim.x) {
|
||||
int v = d_st[lvl_start + j / warp_size];
|
||||
double dist_v = d_dist[v];
|
||||
int edge_start = d_V[v];
|
||||
int edge_end = d_V[v + 1];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
if (dist_v + d_W[e + edge_start] == d_dist[adj]) {
|
||||
atomicAddDouble(d_sigma + adj, d_sigma[v]);
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
// calculate sigma END
|
||||
|
||||
__shared__ int depth, st_start, st_end;
|
||||
if (threadIdx.x == 0) {
|
||||
depth = len_st_idx - 1;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0 && endpoints) {
|
||||
atomicAddDouble(d_BC + s, d_st_idx[depth] - 1);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (depth > 0) {
|
||||
if (threadIdx.x == 0) {
|
||||
st_start = d_st_idx[depth - 1];
|
||||
st_end = d_st_idx[depth];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int j = threadIdx.x; j < (st_end - st_start) * warp_size; j += blockDim.x) {
|
||||
int pred = d_st[st_start + j / warp_size];
|
||||
int edge_start = d_V[pred];
|
||||
int edge_end = d_V[pred + 1];
|
||||
double pred_sigma = d_sigma[pred];
|
||||
double pred_dist = d_dist[pred];
|
||||
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int succ = d_E[e + edge_start];
|
||||
double weight = d_W[e + edge_start];
|
||||
double succ_dist = d_dist[succ];
|
||||
if (succ_dist == pred_dist + weight) {
|
||||
atomicAddDouble(d_delta + pred,
|
||||
pred_sigma / d_sigma[succ] * (1 + d_delta[succ]));
|
||||
}
|
||||
}
|
||||
__threadfence_block();
|
||||
|
||||
if (j % warp_size == 0 && s != pred) {
|
||||
atomicAddDouble(d_BC + pred, d_delta[pred] + endpoints);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
--depth;
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_rescale(
|
||||
_IN_ int len_V,
|
||||
_IN_ double scale,
|
||||
_OUT_ double* d_BC
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int tnum = blockDim.x * gridDim.x;
|
||||
|
||||
for (int u = tid; u < len_V; u += tnum) {
|
||||
d_BC[u] *= scale;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static double calc_scale(
|
||||
_IN_ int len_V,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints
|
||||
)
|
||||
{
|
||||
double scale = 1.0;
|
||||
if (normalized) {
|
||||
if (endpoints) {
|
||||
if (len_V < 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / (double(len_V) * (len_V - 1));
|
||||
}
|
||||
} else if (len_V <= 2) {
|
||||
scale = 1.0;
|
||||
} else {
|
||||
scale = 1.0 / ((double(len_V) - 1) * (len_V - 2));
|
||||
}
|
||||
} else {
|
||||
if (!is_directed) {
|
||||
scale = 0.5;
|
||||
} else {
|
||||
scale = 1.0;
|
||||
}
|
||||
}
|
||||
return scale;
|
||||
}
|
||||
|
||||
|
||||
|
||||
int cuda_betweenness_centrality (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* BC
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int min_edge_block_size;
|
||||
int min_edge_grid_size;
|
||||
int dijkstra_block_size;
|
||||
int dijkstra_grid_size;
|
||||
int rescale_block_size;
|
||||
int rescale_grid_size;
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&min_edge_grid_size, &min_edge_block_size, d_calc_min_edge, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&dijkstra_grid_size, &dijkstra_block_size, d_dijkstra_bc, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&rescale_grid_size, &rescale_block_size, d_rescale, 0, 0);
|
||||
|
||||
double scale = calc_scale(len_V, is_directed, normalized, endpoints);
|
||||
|
||||
int *d_curr_node = NULL;
|
||||
int *d_V = NULL, *d_E = NULL, *d_sources= NULL;
|
||||
int *d_U_2D = NULL, *d_F_2D = NULL, *d_st_2D = NULL, *d_st_idx_2D = NULL;
|
||||
double *d_W = NULL, *d_min_edge = NULL, *d_dist_2D = NULL,
|
||||
*d_sigma_2D = NULL, *d_delta_2D = NULL, *d_BC = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_curr_node, sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * (len_V + 1)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sources, sizeof(int) * len_sources));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_U_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_F_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_st_idx_2D, sizeof(int) * dijkstra_grid_size * (len_V + 2)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, sizeof(double) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_min_edge, sizeof(double) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_dist_2D, sizeof(double) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sigma_2D, sizeof(double) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_delta_2D, sizeof(double) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_BC, sizeof(double) * len_V));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemset(d_curr_node, 0, 1));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * (len_V + 1), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_sources, sources, sizeof(int) * len_sources, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, sizeof(double) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
d_calc_min_edge<<<min_edge_grid_size, min_edge_block_size>>>(d_V, d_E, d_W, len_V, len_E, d_min_edge);
|
||||
|
||||
d_dijkstra_bc<<<dijkstra_grid_size, dijkstra_block_size>>>(d_curr_node, d_V, d_E, d_W, d_min_edge,
|
||||
d_sources, d_dist_2D, d_sigma_2D, d_delta_2D, d_U_2D,
|
||||
d_F_2D, d_st_2D, d_st_idx_2D, len_V, len_E, len_sources,
|
||||
warp_size, endpoints, d_BC);
|
||||
|
||||
if (scale != 1.0) {
|
||||
d_rescale<<<rescale_grid_size, rescale_block_size>>>(len_V, scale, d_BC);
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(BC, d_BC, sizeof(double) * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_curr_node);
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_sources);
|
||||
cudaFree(d_U_2D);
|
||||
cudaFree(d_F_2D);
|
||||
cudaFree(d_st_2D);
|
||||
cudaFree(d_st_idx_2D);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_min_edge);
|
||||
cudaFree(d_dist_2D);
|
||||
cudaFree(d_sigma_2D);
|
||||
cudaFree(d_delta_2D);
|
||||
cudaFree(d_BC);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,64 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
/**
|
||||
* description:
|
||||
* use cuda to calculate betweenness_centrality. the graph must be
|
||||
* in CSR format.
|
||||
*
|
||||
* arguments:
|
||||
* V -
|
||||
* the vertices in CSR format
|
||||
*
|
||||
* E -
|
||||
* the edges in CSR format
|
||||
*
|
||||
* W -
|
||||
* the weight of edges in CSR format
|
||||
*
|
||||
* sources -
|
||||
* set of source vertices to consider when calculating shortest paths.
|
||||
*
|
||||
* len_V -
|
||||
* len of V
|
||||
*
|
||||
* len_E -
|
||||
* len of E
|
||||
*
|
||||
* warp_size -
|
||||
* the number of threads assigned to a vertex
|
||||
*
|
||||
* is_directed -
|
||||
* if this graph is directed
|
||||
*
|
||||
* normalized -
|
||||
* if the answer needs to be normalized
|
||||
*
|
||||
* endpoints -
|
||||
* if true include the endpoints in the shortest basic counts.
|
||||
*
|
||||
* BC -
|
||||
* betweenness centrality output
|
||||
*
|
||||
* return:
|
||||
* EG_GPU_STATUS_CODE
|
||||
*/
|
||||
int cuda_betweenness_centrality (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_IN_ int is_directed,
|
||||
_IN_ int normalized,
|
||||
_IN_ int endpoints,
|
||||
_OUT_ double* BC
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,83 @@
|
||||
#include <vector>
|
||||
#include <string>
|
||||
|
||||
#include "centrality/closeness_centrality.cuh"
|
||||
#include "centrality/betweenness_centrality.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::pair;
|
||||
using std::string;
|
||||
using std::vector;
|
||||
|
||||
static int decide_warp_size (
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E
|
||||
)
|
||||
{
|
||||
vector<int> warp_size_cand{1, 2, 4, 8, 16, 32};
|
||||
|
||||
if (len_E / len_V < warp_size_cand.front()) {
|
||||
return warp_size_cand.front();
|
||||
}
|
||||
|
||||
for (int i = 0; i + 1 < warp_size_cand.size(); ++i) {
|
||||
if (warp_size_cand[i] <= len_E / len_V
|
||||
&& len_E / len_V < warp_size_cand[i + 1]) {
|
||||
return warp_size_cand[i + 1];
|
||||
}
|
||||
}
|
||||
return warp_size_cand.back();
|
||||
}
|
||||
|
||||
|
||||
|
||||
int closeness_centrality(
|
||||
_IN_ const std::vector<int>& V,
|
||||
_IN_ const std::vector<int>& E,
|
||||
_IN_ const std::vector<double>& W,
|
||||
_IN_ const std::vector<int>& sources,
|
||||
_OUT_ std::vector<double>& CC
|
||||
) {
|
||||
int len_V = V.size() - 1;
|
||||
int len_E = E.size();
|
||||
|
||||
int warp_size = decide_warp_size(len_V, len_E);
|
||||
|
||||
CC = vector<double>(len_V);
|
||||
|
||||
int r = cuda_closeness_centrality(V.data(), E.data(), W.data(),
|
||||
sources.data(), len_V, len_E, sources.size(),
|
||||
warp_size, CC.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
|
||||
|
||||
int betweenness_centrality(
|
||||
_IN_ const std::vector<int>& V,
|
||||
_IN_ const std::vector<int>& E,
|
||||
_IN_ const std::vector<double>& W,
|
||||
_IN_ const std::vector<int>& sources,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ bool normalized,
|
||||
_IN_ bool endpoints,
|
||||
_OUT_ std::vector<double>& BC
|
||||
) {
|
||||
int len_V = V.size() - 1;
|
||||
int len_E = E.size();
|
||||
|
||||
int warp_size = decide_warp_size(len_V, len_E);
|
||||
|
||||
BC = vector<double>(len_V);
|
||||
|
||||
int r = cuda_betweenness_centrality(V.data(), E.data(), W.data(),
|
||||
sources.data(), len_V, len_E, sources.size(),
|
||||
warp_size, is_directed, normalized, endpoints, BC.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,246 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __device__ double atomicAddDouble (
|
||||
_OUT_ double* address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long int* address_as_ull =
|
||||
(unsigned long long int*)address;
|
||||
unsigned long long int old = *address_as_ull, assumed;
|
||||
do {
|
||||
assumed = old;
|
||||
old = atomicCAS(address_as_ull, assumed,
|
||||
__double_as_longlong(val +
|
||||
__longlong_as_double(assumed)));
|
||||
} while (assumed != old);
|
||||
return __longlong_as_double(old);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __device__ double atomicMinDouble (
|
||||
_OUT_ double *address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long ret = __double_as_longlong(*address);
|
||||
while (val < __longlong_as_double(ret))
|
||||
{
|
||||
unsigned long long old = ret;
|
||||
if ((ret = atomicCAS((unsigned long long *)address, old, __double_as_longlong(val))) == old)
|
||||
break;
|
||||
}
|
||||
return __longlong_as_double(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_calc_min_edge (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ double* d_min_edge
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int tnum = blockDim.x * gridDim.x;
|
||||
|
||||
for (int u = tid; u < len_V; u += tnum) {
|
||||
double curr_min = EG_DOUBLE_INF;
|
||||
int edge_start = d_V[u];
|
||||
int edge_end = d_V[u + 1];
|
||||
for(int v = edge_start; v < edge_end; ++v) {
|
||||
curr_min = min(curr_min, d_W[v]);
|
||||
}
|
||||
d_min_edge[u] = curr_min;
|
||||
}
|
||||
}
|
||||
|
||||
static __global__ void d_dijkstra_cc (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ double* d_min_edge,
|
||||
_IN_ int* d_sources,
|
||||
_BUFFER_ double* d_dist_2D,
|
||||
_BUFFER_ int* d_U_2D,
|
||||
_BUFFER_ int* d_F_2D,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_OUT_ double* d_CC
|
||||
)
|
||||
{
|
||||
for (int s_idx = blockIdx.x; s_idx < len_sources; s_idx += gridDim.x) {
|
||||
int s = d_sources[s_idx];
|
||||
|
||||
int* d_U = d_U_2D + blockIdx.x * len_V;
|
||||
int* d_F = d_F_2D + blockIdx.x * len_V;
|
||||
double* d_dist = d_dist_2D + blockIdx.x * len_V;
|
||||
|
||||
__shared__ int len_F;
|
||||
__shared__ double delta;
|
||||
__shared__ double dist_accum;
|
||||
__shared__ int reachable_cnt;
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
d_U[i] = 1;
|
||||
d_dist[i] = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_dist[s] = 0.0;
|
||||
d_F[0] = s;
|
||||
len_F = 1;
|
||||
delta = 0.0;
|
||||
dist_accum = 0.0;
|
||||
reachable_cnt = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (delta < EG_DOUBLE_INF) {
|
||||
for (int j = threadIdx.x; j < len_F * warp_size; j += blockDim.x) {
|
||||
int f = d_F[j / warp_size];
|
||||
int edge_start = d_V[f];
|
||||
int edge_end = d_V[f + 1];
|
||||
double dist = d_dist[f];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
double relax_w = dist + d_W[e + edge_start];
|
||||
atomicMinDouble(d_dist + adj, relax_w);
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
delta = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] == 1 && dist_i < EG_DOUBLE_INF) {
|
||||
atomicMinDouble(&delta, dist_i + d_min_edge[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
len_F = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] && dist_i <= delta && dist_i < EG_DOUBLE_INF) {
|
||||
d_U[i] = 0;
|
||||
int f_idx = atomicAdd(&len_F, 1);
|
||||
d_F[f_idx] = i;
|
||||
|
||||
atomicAdd(&reachable_cnt, 1);
|
||||
atomicAddDouble(&dist_accum, d_dist[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_CC[s_idx] = dist_accum == 0.0 ? 0.0 :
|
||||
(double)(reachable_cnt - 1) *
|
||||
(double)(reachable_cnt - 1) /
|
||||
((len_V - 1) * dist_accum);
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
// we here use CSR to represent a graph
|
||||
int cuda_closeness_centrality (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_OUT_ double* CC
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int min_edge_block_size;
|
||||
int min_edge_grid_size;
|
||||
int dijkstra_block_size;
|
||||
int dijkstra_grid_size;
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&min_edge_grid_size, &min_edge_block_size, d_calc_min_edge, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&dijkstra_grid_size, &dijkstra_block_size, d_dijkstra_cc, 0, 0);
|
||||
|
||||
int *d_V = NULL, *d_E = NULL, *d_sources= NULL;
|
||||
int *d_U_2D = NULL, *d_F_2D = NULL;
|
||||
double *d_W = NULL, *d_min_edge = NULL, *d_dist_2D = NULL, *d_CC = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * (len_V + 1)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sources, sizeof(int) * len_sources));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_U_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_F_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, sizeof(double) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_min_edge, sizeof(double) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_dist_2D, sizeof(double) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_CC, sizeof(double) * len_V));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * (len_V + 1), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_sources, sources, sizeof(int) * len_sources, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, sizeof(double) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
d_calc_min_edge<<<dijkstra_grid_size, dijkstra_block_size>>>(d_V, d_E, d_W, len_V, len_E, d_min_edge);
|
||||
|
||||
d_dijkstra_cc<<<min_edge_grid_size, min_edge_block_size>>>(d_V, d_E, d_W, d_min_edge, d_sources,
|
||||
d_dist_2D, d_U_2D, d_F_2D, len_V, len_E, len_sources, warp_size, d_CC);
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(CC, d_CC, sizeof(double) * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_sources);
|
||||
cudaFree(d_U_2D);
|
||||
cudaFree(d_F_2D);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_min_edge);
|
||||
cudaFree(d_dist_2D);
|
||||
cudaFree(d_CC);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,53 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
/**
|
||||
* description:
|
||||
* use cuda to calculate closeness_centrality. the graph must be
|
||||
* in CSR format.
|
||||
*
|
||||
* arguments:
|
||||
* V -
|
||||
* the vertices in CSR format
|
||||
*
|
||||
* E -
|
||||
* the edges in CSR format
|
||||
*
|
||||
* W -
|
||||
* the weight of edges in CSR format
|
||||
*
|
||||
* sources -
|
||||
* an array of EG_GPU_NODE_STATUS. the according CC[i] will be
|
||||
* calculated only if sources[i] == EG_GPU_NODE_ACTIVE
|
||||
*
|
||||
* len_V -
|
||||
* len of V
|
||||
*
|
||||
* len_E -
|
||||
* len of E
|
||||
*
|
||||
* warp_size -
|
||||
* the number of threads assigned to a vertex
|
||||
*
|
||||
* CC -
|
||||
* closeness centrality output
|
||||
*
|
||||
* return:
|
||||
* EG_GPU_STATUS_CODE
|
||||
*/
|
||||
int cuda_closeness_centrality (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int warp_size,
|
||||
_OUT_ double* CC
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,25 @@
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "core/k_core.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::vector;
|
||||
|
||||
int k_core(
|
||||
_IN_ const std::vector<int>& V,
|
||||
_IN_ const std::vector<int>& E,
|
||||
_OUT_ std::vector<int>& KC
|
||||
) {
|
||||
int len_V = V.size() - 1;
|
||||
int len_E = E.size();
|
||||
|
||||
KC = vector<int>(len_V, 0);
|
||||
int r = cuda_k_core(V.data(), E.data(), len_V, len_E, KC.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,239 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __global__ void d_calc_deg(
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ int* d_deg
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int tnum = blockDim.x * gridDim.x;
|
||||
for (int u = tid; u < len_V; u += tnum) {
|
||||
d_deg[u] = d_V[u + 1] - d_V[u];
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_k_core_scan(
|
||||
_IN_ int* d_deg,
|
||||
_IN_ int len_V,
|
||||
_IN_ int level,
|
||||
_IN_ int* d_buf_2D,
|
||||
_IN_ int* d_buf_tail_2D
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int threads_num = blockDim.x * gridDim.x;
|
||||
int* d_buf = d_buf_2D + blockIdx.x * len_V;
|
||||
__shared__ int buf_tail;
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
buf_tail = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int base = 0; base < len_V; base += threads_num) {
|
||||
int v = base + tid;
|
||||
|
||||
if (v >= len_V) {
|
||||
continue;
|
||||
}
|
||||
|
||||
if (d_deg[v] == level) {
|
||||
int buf_idx = atomicAdd(&buf_tail, 1);
|
||||
d_buf[buf_idx] = v;
|
||||
}
|
||||
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_buf_tail_2D[blockIdx.x] = buf_tail;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_k_core_loop(
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_OUT_ int* d_deg,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int level,
|
||||
_IN_ int* d_buf_2D,
|
||||
_IN_ int* d_buf_tail_2D,
|
||||
_OUT_ int* d_count
|
||||
)
|
||||
{
|
||||
int warp_size = 32;
|
||||
int tid = threadIdx.x;
|
||||
int* d_buf = d_buf_2D + blockIdx.x * len_V;
|
||||
int warp_id = tid / warp_size;
|
||||
int lane_id = tid % warp_size;
|
||||
int reg_tail;
|
||||
int i;
|
||||
|
||||
__shared__ int buf_tail;
|
||||
__shared__ int base;
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
base = 0;
|
||||
buf_tail = d_buf_tail_2D[blockIdx.x];
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (1) {
|
||||
__syncthreads();
|
||||
|
||||
if (base == buf_tail) {
|
||||
break;
|
||||
}
|
||||
|
||||
i = base + warp_id;
|
||||
reg_tail = buf_tail;
|
||||
__syncthreads();
|
||||
|
||||
if (i >= reg_tail) {
|
||||
continue;
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
base += blockDim.x / warp_size;
|
||||
if (reg_tail < base) {
|
||||
base = reg_tail;
|
||||
}
|
||||
}
|
||||
|
||||
int v = d_buf[i];
|
||||
int edge_start = d_V[v];
|
||||
int edge_end = d_V[v + 1];
|
||||
|
||||
while (1) {
|
||||
__syncwarp();
|
||||
|
||||
if (edge_start >= edge_end) {
|
||||
break;
|
||||
}
|
||||
|
||||
int curr_e = edge_start + lane_id;
|
||||
edge_start += warp_size;
|
||||
|
||||
if (curr_e >= edge_end) {
|
||||
continue;
|
||||
}
|
||||
|
||||
int nbr = d_E[curr_e];
|
||||
if (d_deg[nbr] > level) {
|
||||
int a = atomicSub(d_deg + nbr, 1);
|
||||
|
||||
if (a == level + 1) {
|
||||
int loc = atomicAdd(&buf_tail, 1);
|
||||
d_buf[loc] = nbr;
|
||||
}
|
||||
|
||||
if (a <= level) {
|
||||
atomicAdd(d_deg + nbr, 1);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0 && buf_tail) {
|
||||
atomicAdd(d_count, buf_tail);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
|
||||
int cuda_k_core (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ int* k_core_res
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int calc_deg_block_size;
|
||||
int calc_deg_grid_size;
|
||||
int scan_block_size;
|
||||
int scan_grid_size;
|
||||
int loop_block_size;
|
||||
int loop_grid_size;
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&calc_deg_grid_size, &calc_deg_block_size, d_calc_deg, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&scan_grid_size, &scan_block_size, d_k_core_scan, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&loop_grid_size, &loop_block_size, d_k_core_loop, 0, 0);
|
||||
|
||||
int k_core_grid_size = max(scan_grid_size, loop_grid_size);
|
||||
|
||||
int count = 0, level = 0;
|
||||
|
||||
int *d_V = NULL, *d_E = NULL, *d_deg = NULL, *d_k_core_res = NULL,
|
||||
*d_buf_2D = NULL, *d_buf_tail_2D = NULL, *d_count = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * (len_V + 1)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_deg, sizeof(int) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_k_core_res, sizeof(int) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_buf_2D, sizeof(int) * k_core_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_buf_tail_2D, sizeof(int) * k_core_grid_size));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_count, sizeof(int)));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * (len_V + 1), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemset(d_count, 0, sizeof(int)));
|
||||
|
||||
d_calc_deg<<<calc_deg_grid_size, calc_deg_block_size>>>(d_V, d_E, len_V, len_E, d_deg);
|
||||
while (count < len_V) {
|
||||
EXIT_IF_CUDA_FAILED(cudaMemset(d_buf_tail_2D, 0, sizeof(int) * k_core_grid_size));
|
||||
|
||||
d_k_core_scan<<<k_core_grid_size, scan_block_size>>>(d_deg, len_V, level, d_buf_2D, d_buf_tail_2D);
|
||||
|
||||
d_k_core_loop<<<k_core_grid_size, loop_block_size>>>(d_V, d_E, d_deg, len_V, len_E, level,
|
||||
d_buf_2D, d_buf_tail_2D, d_count);
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(&count, d_count, sizeof(int), cudaMemcpyDeviceToHost));
|
||||
|
||||
++level;
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(k_core_res, d_deg, sizeof(int) * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_deg);
|
||||
cudaFree(d_k_core_res);
|
||||
cudaFree(d_buf_2D);
|
||||
cudaFree(d_buf_tail_2D);
|
||||
cudaFree(d_count);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,39 @@
|
||||
# pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
/**
|
||||
* description:
|
||||
* use cuda to calculate k core. the graph must be
|
||||
* in CSR format.
|
||||
*
|
||||
* arguments:
|
||||
* V -
|
||||
* the vertices in CSR format
|
||||
*
|
||||
* E -
|
||||
* the edges in CSR format
|
||||
*
|
||||
* len_V -
|
||||
* len of V
|
||||
*
|
||||
* len_E -
|
||||
* len of E
|
||||
*
|
||||
* k_core_res -
|
||||
* result of k_core
|
||||
*
|
||||
* return:
|
||||
* EG_GPU_STATUS_CODE
|
||||
*/
|
||||
int cuda_k_core (
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ int* k_core_res
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,63 @@
|
||||
#include <limits>
|
||||
#include <vector>
|
||||
|
||||
#include "path/sssp_dijkstra.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::vector;
|
||||
|
||||
static int decide_warp_size(
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E
|
||||
)
|
||||
{
|
||||
vector<int> warp_size_cand{1, 2, 4, 8, 16, 32};
|
||||
|
||||
if (len_E / len_V < warp_size_cand.front()) {
|
||||
return warp_size_cand.front();
|
||||
}
|
||||
|
||||
for (int i = 0; i + 1 < warp_size_cand.size(); ++i) {
|
||||
if (warp_size_cand[i] <= len_E / len_V
|
||||
&& len_E / len_V < warp_size_cand[i + 1]) {
|
||||
return warp_size_cand[i + 1];
|
||||
}
|
||||
}
|
||||
return warp_size_cand.back();
|
||||
}
|
||||
|
||||
|
||||
|
||||
int sssp_dijkstra(
|
||||
_IN_ const vector<int>& V,
|
||||
_IN_ const vector<int>& E,
|
||||
_IN_ const vector<double>& W,
|
||||
_IN_ const vector<int>& sources,
|
||||
_IN_ int target,
|
||||
_OUT_ vector<double>& res
|
||||
)
|
||||
{
|
||||
int len_V = V.size() - 1;
|
||||
int len_E = E.size();
|
||||
|
||||
int warp_size = decide_warp_size(len_V, len_E);
|
||||
|
||||
res = vector<double>(sources.size() * V.size());
|
||||
|
||||
int r = cuda_sssp_dijkstra(V.data(), E.data(), W.data(),
|
||||
sources.data(), len_V, len_E, sources.size(),
|
||||
target, warp_size, res.data());
|
||||
|
||||
double double_inf = std::numeric_limits<double>::infinity();
|
||||
for (int i = 0; i < res.size(); ++i) {
|
||||
if (res[i] >= EG_DOUBLE_INF) {
|
||||
res[i] = double_inf;
|
||||
}
|
||||
}
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,233 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
static __device__ double atomicMinDouble (
|
||||
_OUT_ double *address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long ret = __double_as_longlong(*address);
|
||||
while (val < __longlong_as_double(ret))
|
||||
{
|
||||
unsigned long long old = ret;
|
||||
if ((ret = atomicCAS((unsigned long long *)address, old, __double_as_longlong(val))) == old)
|
||||
break;
|
||||
}
|
||||
return __longlong_as_double(ret);
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_calc_min_edge (
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_OUT_ double* d_min_edge
|
||||
)
|
||||
{
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int tnum = blockDim.x * gridDim.x;
|
||||
|
||||
for (int u = tid; u < len_V; u += tnum) {
|
||||
double curr_min = EG_DOUBLE_INF;
|
||||
int edge_start = d_V[u];
|
||||
int edge_end = d_V[u + 1];
|
||||
for(int v = edge_start; v < edge_end; ++v) {
|
||||
curr_min = min(curr_min, d_W[v]);
|
||||
}
|
||||
d_min_edge[u] = curr_min;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
static __global__ void d_sssp_dijkstra (
|
||||
_IN_ int* d_curr_node,
|
||||
_IN_ int* d_V,
|
||||
_IN_ int* d_E,
|
||||
_IN_ double* d_W,
|
||||
_IN_ double* d_min_edge,
|
||||
_IN_ int* d_sources,
|
||||
_OUT_ double* d_dist_2D,
|
||||
_BUFFER_ int* d_U_2D,
|
||||
_BUFFER_ int* d_F_2D,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int target,
|
||||
_IN_ int warp_size
|
||||
)
|
||||
{
|
||||
while (1) {
|
||||
__shared__ int curr_node;
|
||||
if (threadIdx.x == 0) {
|
||||
curr_node = atomicAdd(d_curr_node, 1);
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (curr_node >= len_sources) {
|
||||
break;
|
||||
}
|
||||
|
||||
int s = d_sources[curr_node];
|
||||
|
||||
double* d_dist = d_dist_2D + curr_node * len_V;
|
||||
int* d_U = d_U_2D + blockIdx.x * len_V;
|
||||
int* d_F = d_F_2D + blockIdx.x * len_V;
|
||||
|
||||
__shared__ int len_F;
|
||||
__shared__ double delta;
|
||||
__shared__ int target_cnt;
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
d_U[i] = 1;
|
||||
d_dist[i] = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
d_dist[s] = 0.0;
|
||||
d_F[0] = s;
|
||||
len_F = 1;
|
||||
delta = 0.0;
|
||||
target_cnt = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
while (delta < EG_DOUBLE_INF && target_cnt == 0) {
|
||||
for (int j = threadIdx.x; j < len_F * warp_size; j += blockDim.x) {
|
||||
int f = d_F[j / warp_size];
|
||||
int edge_start = d_V[f];
|
||||
int edge_end = d_V[f + 1];
|
||||
double dist = d_dist[f];
|
||||
for (int e = j % warp_size; e < edge_end - edge_start; e += warp_size) {
|
||||
int adj = d_E[e + edge_start];
|
||||
double relax_w = dist + d_W[e + edge_start];
|
||||
atomicMinDouble(d_dist + adj, relax_w);
|
||||
}
|
||||
__threadfence_block();
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
delta = EG_DOUBLE_INF;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] == 1 && dist_i < EG_DOUBLE_INF) {
|
||||
atomicMinDouble(&delta, dist_i + d_min_edge[i]);
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
len_F = 0;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
for (int i = threadIdx.x; i < len_V; i += blockDim.x) {
|
||||
double dist_i = d_dist[i];
|
||||
if (d_U[i] && dist_i <= delta && dist_i < EG_DOUBLE_INF) {
|
||||
d_U[i] = 0;
|
||||
int f_idx = atomicAdd(&len_F, 1);
|
||||
d_F[f_idx] = i;
|
||||
target_cnt += i == target;
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
// we here use CSR to represent a graph
|
||||
int cuda_sssp_dijkstra(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int target,
|
||||
_IN_ int warp_size,
|
||||
_OUT_ double* res
|
||||
)
|
||||
{
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int min_edge_block_size;
|
||||
int min_edge_grid_size;
|
||||
int dijkstra_block_size;
|
||||
int dijkstra_grid_size;
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&min_edge_grid_size, &min_edge_block_size, d_calc_min_edge, 0, 0);
|
||||
cudaOccupancyMaxPotentialBlockSize(&dijkstra_grid_size, &dijkstra_block_size, d_sssp_dijkstra, 0, 0);
|
||||
|
||||
int *d_curr_node = NULL;
|
||||
int *d_V = NULL, *d_E = NULL, *d_sources= NULL;
|
||||
int *d_U_2D = NULL, *d_F_2D = NULL;
|
||||
double *d_W = NULL, *d_min_edge = NULL, *d_dist_2D = NULL;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_curr_node, sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, sizeof(int) * (len_V + 1)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, sizeof(int) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_sources, sizeof(int) * len_sources));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_U_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_F_2D, sizeof(int) * dijkstra_grid_size * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, sizeof(double) * len_E));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_min_edge, sizeof(double) * len_V));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_dist_2D, sizeof(double) * len_sources * len_V));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemset(d_curr_node, 0, sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, sizeof(int) * (len_V + 1), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, sizeof(int) * len_E, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_sources, sources, sizeof(int) * len_sources, cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, sizeof(double) * len_E, cudaMemcpyHostToDevice));
|
||||
|
||||
d_calc_min_edge<<<dijkstra_grid_size, dijkstra_block_size>>>(d_V, d_E, d_W, len_V, len_E, d_min_edge);
|
||||
|
||||
d_sssp_dijkstra<<<min_edge_grid_size, min_edge_block_size>>>(d_curr_node ,d_V, d_E, d_W, d_min_edge, d_sources,
|
||||
d_dist_2D, d_U_2D, d_F_2D, len_V, len_E, len_sources, target, warp_size);
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(res, d_dist_2D, sizeof(double) * len_sources * len_V, cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_curr_node);
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_sources);
|
||||
cudaFree(d_U_2D);
|
||||
cudaFree(d_F_2D);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_min_edge);
|
||||
cudaFree(d_dist_2D);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,20 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
int cuda_sssp_dijkstra(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const double* W,
|
||||
_IN_ const int* sources,
|
||||
_IN_ int len_V,
|
||||
_IN_ int len_E,
|
||||
_IN_ int len_sources,
|
||||
_IN_ int target,
|
||||
_IN_ int warp_size,
|
||||
_OUT_ double* res
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,31 @@
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
|
||||
#include "structural_holes/constraint.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::vector;
|
||||
|
||||
int constraint(
|
||||
_IN_ const vector<int>& V,
|
||||
_IN_ const vector<int>& E,
|
||||
_IN_ const vector<int>& row,
|
||||
_IN_ const vector<int>& col,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ const vector<double>& W,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ vector<int>& node_mask,
|
||||
_OUT_ vector<double>& constraint
|
||||
) {
|
||||
int num_edges = row.size();
|
||||
|
||||
constraint = vector<double>(num_nodes);
|
||||
int r = cuda_constraint(V.data(), E.data(), row.data(), col.data(), W.data(), num_nodes, num_edges, is_directed, node_mask.data(), constraint.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,341 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
#define NODES_PER_BLOCK 1
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
enum norm_t { SUM = 0, MAX = 1 };
|
||||
|
||||
static __device__ double mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv;
|
||||
}
|
||||
|
||||
static __device__ double normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if (norm == SUM) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
} else if (norm == MAX) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double local_constraint(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double direct = normalized_mutual_weight(V,E,W,u,v,SUM);
|
||||
double indirect = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double norm_uw = normalized_mutual_weight(V, E, W, u, neighbor,SUM);
|
||||
double norm_wv = normalized_mutual_weight(V, E, W, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
double local_constraint_of_uv = (direct + indirect) * (direct + indirect);
|
||||
return local_constraint_of_uv;
|
||||
}
|
||||
|
||||
__global__ void calculate_constraints(
|
||||
const int* __restrict__ V,
|
||||
const int* __restrict__ E,
|
||||
const double* __restrict__ W,
|
||||
const int num_nodes,
|
||||
const int* __restrict__ node_mask,
|
||||
double* __restrict__ constraint_results
|
||||
) {
|
||||
int start_node = blockIdx.x * NODES_PER_BLOCK;
|
||||
int end_node = min(start_node + NODES_PER_BLOCK, num_nodes);
|
||||
|
||||
for (int v = start_node; v < end_node; ++v) {
|
||||
if (!node_mask[v]) continue;
|
||||
|
||||
double constraint_of_v = 0.0;
|
||||
bool is_nan = true;
|
||||
|
||||
__shared__ double shared_constraint[256];
|
||||
double local_sum = 0.0;
|
||||
|
||||
for (int i = V[v] + threadIdx.x; i < V[v + 1]; i += blockDim.x) {
|
||||
is_nan = false;
|
||||
int neighbor = E[i];
|
||||
local_sum += local_constraint(V, E, W, v, neighbor);
|
||||
}
|
||||
|
||||
shared_constraint[threadIdx.x] = local_sum;
|
||||
__syncthreads();
|
||||
|
||||
for (int offset = blockDim.x / 2; offset > 0; offset /= 2) {
|
||||
if (threadIdx.x < offset) {
|
||||
shared_constraint[threadIdx.x] += shared_constraint[threadIdx.x + offset];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
constraint_results[v] = (is_nan) ? NAN : shared_constraint[0];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
static __device__ double directed_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0, a_vu = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
for (int i = V[v]; i < V[v+1]; i++) {
|
||||
if (E[i] == u) {
|
||||
a_vu = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv + a_vu;
|
||||
}
|
||||
|
||||
static __device__ double directed_normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_edges,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = directed_mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if(norm==SUM){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_wu;
|
||||
}
|
||||
}
|
||||
}else if(norm==MAX){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_wu);
|
||||
}
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double directed_local_constraint(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_edges,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double direct = directed_normalized_mutual_weight(V,E,row,col,W,num_edges,u,v,SUM);
|
||||
double indirect = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double norm_uw = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, u, neighbor,SUM);
|
||||
double norm_wv = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double norm_uw = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, u, neighbor,SUM);
|
||||
double norm_wv = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
}
|
||||
double local_constraint_of_uv = (direct + indirect) * (direct + indirect);
|
||||
return local_constraint_of_uv;
|
||||
}
|
||||
|
||||
__global__ void directed_calculate_constraints(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_nodes,
|
||||
int num_edges,
|
||||
int* node_mask,
|
||||
double* constraint_results
|
||||
) {
|
||||
int start_node = blockIdx.x * NODES_PER_BLOCK;
|
||||
int end_node = min(start_node + NODES_PER_BLOCK, num_nodes);
|
||||
|
||||
for (int v = start_node; v < end_node; ++v) {
|
||||
if (!node_mask[v]) continue;
|
||||
|
||||
double constraint_of_v = 0.0;
|
||||
bool is_nan = true;
|
||||
|
||||
__shared__ double shared_constraint[256];
|
||||
double local_sum = 0.0;
|
||||
|
||||
for (int i = V[v] + threadIdx.x; i < V[v + 1]; i += blockDim.x) {
|
||||
is_nan = false;
|
||||
int neighbor = E[i];
|
||||
local_sum += directed_local_constraint(V, E, row, col, W, num_edges, v, neighbor);
|
||||
}
|
||||
|
||||
for (int i = threadIdx.x; i < num_edges; i += blockDim.x) {
|
||||
if (col[i] == v) {
|
||||
// is_nan = false;
|
||||
int neighbor = row[i];
|
||||
local_sum += directed_local_constraint(V, E, row, col, W, num_edges, v, neighbor);
|
||||
}
|
||||
}
|
||||
|
||||
shared_constraint[threadIdx.x] = local_sum;
|
||||
__syncthreads();
|
||||
|
||||
for (int offset = blockDim.x / 2; offset > 0; offset /= 2) {
|
||||
if (threadIdx.x < offset) {
|
||||
shared_constraint[threadIdx.x] += shared_constraint[threadIdx.x + offset];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
constraint_results[v] = (is_nan) ? NAN : shared_constraint[0];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
int cuda_constraint(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* constraint_results
|
||||
) {
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
|
||||
int* d_V;
|
||||
int* d_E;
|
||||
int* d_row;
|
||||
int* d_col;
|
||||
double* d_W;
|
||||
int* d_node_mask;
|
||||
double* d_constraint_results;
|
||||
int block_size = 256;
|
||||
int grid_size = (num_nodes + NODES_PER_BLOCK - 1) / NODES_PER_BLOCK;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, (num_nodes+1) * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_row, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_col, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, num_edges * sizeof(double)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_node_mask, num_nodes * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_constraint_results, num_nodes * sizeof(double)));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, (num_nodes+1) * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_row, row, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_col, col, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_node_mask, node_mask, num_nodes * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, num_edges * sizeof(double), cudaMemcpyHostToDevice));
|
||||
|
||||
if(is_directed){
|
||||
directed_calculate_constraints<<<grid_size, block_size>>>(d_V, d_E, d_row, d_col, d_W, num_nodes, num_edges, d_node_mask, d_constraint_results);
|
||||
}else{
|
||||
calculate_constraints<<<grid_size, block_size>>>(d_V, d_E, d_W, num_nodes, d_node_mask, d_constraint_results);
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(constraint_results, d_constraint_results, num_nodes * sizeof(double), cudaMemcpyDeviceToHost));
|
||||
exit:
|
||||
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_row);
|
||||
cudaFree(d_col);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_node_mask);
|
||||
cudaFree(d_constraint_results);
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,20 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
int cuda_constraint(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* constraint_results
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,31 @@
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
|
||||
#include "structural_holes/effective_size.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::vector;
|
||||
|
||||
int effective_size(
|
||||
_IN_ const vector<int>& V,
|
||||
_IN_ const vector<int>& E,
|
||||
_IN_ const vector<int>& row,
|
||||
_IN_ const vector<int>& col,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ const vector<double>& W,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ vector<int>& node_mask,
|
||||
_OUT_ vector<double>& effective_size
|
||||
) {
|
||||
int num_edges = row.size();
|
||||
|
||||
effective_size = vector<double>(num_nodes);
|
||||
int r = cuda_effective_size(V.data(), E.data(), row.data(), col.data(), W.data(), num_nodes, num_edges, is_directed, node_mask.data(), effective_size.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,346 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
#define NODES_PER_BLOCK 1
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
enum norm_t { SUM = 0, MAX = 1 };
|
||||
|
||||
static __device__ double mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv;
|
||||
}
|
||||
|
||||
static __device__ double normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if (norm == SUM) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
} else if (norm == MAX) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double directed_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0, a_vu = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
for (int i = V[v]; i < V[v+1]; i++) {
|
||||
if (E[i] == u) {
|
||||
a_vu = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv + a_vu;
|
||||
}
|
||||
|
||||
static __device__ double directed_normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_edges,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = directed_mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if(norm==SUM){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_wu;
|
||||
}
|
||||
}
|
||||
}else if(norm==MAX){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_wu);
|
||||
}
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double redundancy(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
const int num_nodes,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double r = 0.0;
|
||||
for (int i = V[v]; i < V[v + 1]; i++) {
|
||||
int w = E[i];
|
||||
r += normalized_mutual_weight(V, E, W, u, w, SUM) * normalized_mutual_weight(V, E, W, v, w, MAX);
|
||||
}
|
||||
return 1-r;
|
||||
}
|
||||
|
||||
|
||||
__inline__ __device__ double warp_reduce_sum(double val)
|
||||
{
|
||||
for (int offset = warpSize / 2; offset > 0; offset /= 2)
|
||||
{
|
||||
val += __shfl_down_sync(0xffffffff, val, offset);
|
||||
}
|
||||
return val;
|
||||
}
|
||||
|
||||
__inline__ __device__ double block_reduce_sum(double val)
|
||||
{
|
||||
val = warp_reduce_sum(val);
|
||||
|
||||
__shared__ double shared[32];
|
||||
int warp_id = threadIdx.x / warpSize;
|
||||
if (threadIdx.x % warpSize == 0)
|
||||
{
|
||||
shared[warp_id] = val;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
if (warp_id == 0)
|
||||
{
|
||||
val = (threadIdx.x < (blockDim.x / warpSize)) ? shared[threadIdx.x] : 0.0;
|
||||
val = warp_reduce_sum(val);
|
||||
}
|
||||
return val;
|
||||
}
|
||||
|
||||
__global__ void calculate_effective_size(
|
||||
const int* __restrict__ V,
|
||||
const int* __restrict__ E,
|
||||
const double* __restrict__ W,
|
||||
const int num_nodes,
|
||||
const int* __restrict__ node_mask,
|
||||
double* __restrict__ effective_size_results
|
||||
) {
|
||||
int u = blockIdx.x;
|
||||
if (u >= num_nodes || !node_mask[u]) return;
|
||||
|
||||
int neighbor_start = V[u];
|
||||
int neighbor_end = V[u + 1];
|
||||
int degree = neighbor_end - neighbor_start;
|
||||
|
||||
int threads_per_block = blockDim.x;
|
||||
|
||||
double redundancy_sum = 0.0;
|
||||
for (int idx = threadIdx.x; idx < degree; idx += threads_per_block) {
|
||||
int i = neighbor_start + idx;
|
||||
int v = E[i];
|
||||
if (v != u) {
|
||||
double r = 0.0;
|
||||
for (int j = V[v]; j < V[v + 1]; j++) {
|
||||
int w = E[j];
|
||||
r += normalized_mutual_weight(V, E, W, u, w, SUM) *
|
||||
normalized_mutual_weight(V, E, W, v, w, MAX);
|
||||
}
|
||||
redundancy_sum += 1 - r;
|
||||
}
|
||||
}
|
||||
|
||||
redundancy_sum = block_reduce_sum(redundancy_sum);
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
effective_size_results[u] = redundancy_sum;
|
||||
}
|
||||
}
|
||||
|
||||
static __device__ double directed_redundancy(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
const int num_nodes,
|
||||
const int num_edges,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double r = 0.0;
|
||||
for (int i = V[v]; i < V[v + 1]; i++) {
|
||||
int w = E[i];
|
||||
r += directed_normalized_mutual_weight(V, E, row,col,W,num_edges, u, w,SUM) * directed_normalized_mutual_weight(V, E, row,col,W, num_edges, v,w,MAX);
|
||||
}
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == v) {
|
||||
int w = row[i];
|
||||
r += directed_normalized_mutual_weight(V, E, row,col,W,num_edges, u, w,SUM) * directed_normalized_mutual_weight(V, E, row,col,W, num_edges, v,w,MAX);
|
||||
}
|
||||
}
|
||||
return 1-r;
|
||||
}
|
||||
|
||||
__global__ void directed_calculate_effective_size(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
const int num_nodes,
|
||||
const int num_edges,
|
||||
const int* node_mask,
|
||||
double* effective_size_results
|
||||
) {
|
||||
int u = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (u >= num_nodes || !node_mask[u]) return;
|
||||
double redundancy_sum = 0.0;
|
||||
bool is_nan = true;
|
||||
|
||||
for (int i = V[u]; i < V[u + 1]; i++) {
|
||||
int v = E[i];
|
||||
if (v == u) continue;
|
||||
is_nan = false;
|
||||
redundancy_sum += directed_redundancy(V,E,row,col,W,num_nodes,num_edges,u,v);
|
||||
}
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int v = row[i];
|
||||
redundancy_sum += directed_redundancy(V,E,row,col,W,num_nodes,num_edges,u,v);
|
||||
}
|
||||
}
|
||||
effective_size_results[u] = is_nan ? NAN : redundancy_sum;
|
||||
}
|
||||
|
||||
|
||||
int cuda_effective_size(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* effective_size_results
|
||||
) {
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
int min_grid_size = 0;
|
||||
int block_size = 0;
|
||||
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&min_grid_size, &block_size, calculate_effective_size, 0, 0);
|
||||
|
||||
int grid_size = (num_nodes + block_size * NODES_PER_BLOCK - 1) / (block_size * NODES_PER_BLOCK);
|
||||
|
||||
int* d_V;
|
||||
int* d_E;
|
||||
int* d_row;
|
||||
int* d_col;
|
||||
double* d_W;
|
||||
int* d_node_mask;
|
||||
double* d_effective_size_results;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, (num_nodes+1) * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_row, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_col, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, num_edges * sizeof(double)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_node_mask, num_nodes * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_effective_size_results, num_nodes * sizeof(double)));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, (num_nodes+1) * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_row, row, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_col, col, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_node_mask, node_mask, num_nodes * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, num_edges * sizeof(double), cudaMemcpyHostToDevice));
|
||||
|
||||
if(is_directed){
|
||||
directed_calculate_effective_size<<<grid_size, block_size>>>(d_V, d_E, d_row, d_col, d_W, num_nodes, num_edges, d_node_mask, d_effective_size_results);
|
||||
}else{
|
||||
int block_size = 256;
|
||||
int grid_size = (num_nodes + NODES_PER_BLOCK - 1) / NODES_PER_BLOCK;
|
||||
calculate_effective_size<<<grid_size, block_size>>>(d_V, d_E, d_W, num_nodes, d_node_mask, d_effective_size_results);
|
||||
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(effective_size_results, d_effective_size_results, num_nodes * sizeof(double), cudaMemcpyDeviceToHost));
|
||||
|
||||
exit:
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_row);
|
||||
cudaFree(d_col);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_node_mask);
|
||||
cudaFree(d_effective_size_results);
|
||||
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,20 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
int cuda_effective_size(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* effective_size_results
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,31 @@
|
||||
#include <string>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
|
||||
#include "structural_holes/hierarchy.cuh"
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
using std::vector;
|
||||
|
||||
int hierarchy(
|
||||
_IN_ const vector<int>& V,
|
||||
_IN_ const vector<int>& E,
|
||||
_IN_ const vector<int>& row,
|
||||
_IN_ const vector<int>& col,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ const vector<double>& W,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ vector<int>& node_mask,
|
||||
_OUT_ vector<double>& hierarchy
|
||||
) {
|
||||
int num_edges = row.size();
|
||||
|
||||
hierarchy = vector<double>(num_nodes);
|
||||
int r = cuda_hierarchy(V.data(), E.data(), row.data(), col.data(), W.data(), num_nodes, num_edges, is_directed, node_mask.data(), hierarchy.data());
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,369 @@
|
||||
#include <cuda.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include "common.h"
|
||||
#define NODES_PER_BLOCK 1
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
enum norm_t { SUM = 0, MAX = 1 };
|
||||
|
||||
static __device__ double mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv;
|
||||
}
|
||||
|
||||
static __device__ double normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if (norm == SUM) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
} else if (norm == MAX) {
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double directed_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double a_uv = 0.0, a_vu = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
if (E[i] == v) {
|
||||
a_uv = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
for (int i = V[v]; i < V[v+1]; i++) {
|
||||
if (E[i] == u) {
|
||||
a_vu = W[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return a_uv + a_vu;
|
||||
}
|
||||
|
||||
static __device__ double directed_normalized_mutual_weight(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_edges,
|
||||
int u,
|
||||
int v,
|
||||
norm_t norm
|
||||
) {
|
||||
double weight_uv = directed_mutual_weight(V, E, W, u, v);
|
||||
|
||||
double scale = 0.0;
|
||||
if(norm==SUM){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_uw;
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale += weight_wu;
|
||||
}
|
||||
}
|
||||
}else if(norm==MAX){
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double weight_uw = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_uw);
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double weight_wu = directed_mutual_weight(V, E, W, u, neighbor);
|
||||
scale = fmax(scale,weight_wu);
|
||||
}
|
||||
}
|
||||
}
|
||||
return (scale==0.0) ? 0.0 : (weight_uv / scale);
|
||||
}
|
||||
|
||||
static __device__ double atomicAdd (
|
||||
_OUT_ double* address,
|
||||
_IN_ double val
|
||||
)
|
||||
{
|
||||
unsigned long long int* address_as_ull =
|
||||
(unsigned long long int*)address;
|
||||
unsigned long long int old = *address_as_ull, assumed;
|
||||
do {
|
||||
assumed = old;
|
||||
old = atomicCAS(address_as_ull, assumed,
|
||||
__double_as_longlong(val +
|
||||
__longlong_as_double(assumed)));
|
||||
} while (assumed != old);
|
||||
return __longlong_as_double(old);
|
||||
}
|
||||
|
||||
static __device__ double local_constraint(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double direct = normalized_mutual_weight(V,E,W,u,v,SUM);
|
||||
double indirect = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double norm_uw = normalized_mutual_weight(V, E, W, u, neighbor,SUM);
|
||||
double norm_wv = normalized_mutual_weight(V, E, W, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
double local_constraint_of_uv = (direct + indirect) * (direct + indirect);
|
||||
return local_constraint_of_uv;
|
||||
}
|
||||
|
||||
__global__ void calculate_hierarchy(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const double* W,
|
||||
int num_nodes,
|
||||
const int* node_mask,
|
||||
double* hierarchy_results
|
||||
) {
|
||||
int start_node = blockIdx.x * NODES_PER_BLOCK;
|
||||
int end_node = min(start_node + NODES_PER_BLOCK, num_nodes);
|
||||
|
||||
extern __shared__ double shared_mem[];
|
||||
double* shared_c = shared_mem;
|
||||
double* shared_C = &shared_mem[blockDim.x];
|
||||
|
||||
for (int v = start_node; v < end_node; ++v) {
|
||||
if (!node_mask[v]) continue;
|
||||
|
||||
int n = V[v + 1] - V[v];
|
||||
if (n <= 1) {
|
||||
hierarchy_results[v] = 0.0;
|
||||
continue;
|
||||
}
|
||||
if (threadIdx.x == 0) shared_C[0] = 0.0;
|
||||
__syncthreads();
|
||||
|
||||
double local_C = 0.0;
|
||||
|
||||
for (int i = V[v] + threadIdx.x; i < V[v + 1]; i += blockDim.x) {
|
||||
int w = E[i];
|
||||
double constraint = local_constraint(V, E, W, v, w);
|
||||
shared_c[threadIdx.x] = constraint;
|
||||
local_C += constraint;
|
||||
}
|
||||
|
||||
atomicAdd(&shared_C[0], local_C);
|
||||
__syncthreads();
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
double C = shared_C[0];
|
||||
double hierarchy_sum = 0.0;
|
||||
for (int i = 0; i < n; i++) {
|
||||
double normalized_c = shared_c[i] / C;
|
||||
hierarchy_sum += normalized_c * n * logf(normalized_c * n) / (n * logf(n));
|
||||
}
|
||||
hierarchy_results[v] = hierarchy_sum;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
}
|
||||
}
|
||||
|
||||
static __device__ double directed_local_constraint(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
int num_edges,
|
||||
int u,
|
||||
int v
|
||||
) {
|
||||
double direct = directed_normalized_mutual_weight(V,E,row,col,W,num_edges,u,v,SUM);
|
||||
double indirect = 0.0;
|
||||
for (int i = V[u]; i < V[u+1]; i++) {
|
||||
int neighbor = E[i];
|
||||
double norm_uw = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, u, neighbor,SUM);
|
||||
double norm_wv = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == u) {
|
||||
int neighbor = row[i];
|
||||
double norm_uw = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, u, neighbor,SUM);
|
||||
double norm_wv = directed_normalized_mutual_weight(V, E, row, col, W, num_edges, neighbor, v,SUM);
|
||||
indirect += norm_uw * norm_wv;
|
||||
}
|
||||
}
|
||||
double local_constraint_of_uv = (direct + indirect) * (direct + indirect);
|
||||
return local_constraint_of_uv;
|
||||
}
|
||||
|
||||
__global__ void directed_calculate_hierarchy(
|
||||
const int* V,
|
||||
const int* E,
|
||||
const int* row,
|
||||
const int* col,
|
||||
const double* W,
|
||||
const int num_nodes,
|
||||
const int num_edges,
|
||||
const int* node_mask,
|
||||
double* hierarchy_results
|
||||
) {
|
||||
int v = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (v >= num_nodes || !node_mask[v]) return;
|
||||
int in_neighbor = V[v + 1] - V[v];
|
||||
int out_neighbor = 0;
|
||||
double C = 0.0;
|
||||
double hierarchy_sum = 0.0;
|
||||
int neighbor = 0;
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == v) {
|
||||
out_neighbor++;
|
||||
}
|
||||
}
|
||||
double *c = new double[in_neighbor+out_neighbor];
|
||||
for (int i = V[v]; i < V[v + 1]; i++) {
|
||||
int w = E[i];
|
||||
c[neighbor] = directed_local_constraint(V, E, row, col, W, num_edges, v, w);
|
||||
C += c[neighbor];
|
||||
neighbor++;
|
||||
}
|
||||
for (int i = 0; i < num_edges; i++) {
|
||||
if (col[i] == v) {
|
||||
int w = row[i];
|
||||
c[neighbor] = directed_local_constraint(V, E, row, col, W, num_edges, v, w);
|
||||
C += c[neighbor];
|
||||
neighbor++;
|
||||
}
|
||||
}
|
||||
__syncthreads();
|
||||
if (neighbor > 1) {
|
||||
for (int i = 0; i < neighbor; i++) {
|
||||
hierarchy_sum += (c[i] / C) * neighbor * logf((c[i] / C) * neighbor) / (neighbor * logf(neighbor));
|
||||
}
|
||||
hierarchy_results[v] = hierarchy_sum;
|
||||
}else{
|
||||
hierarchy_results[v] = 0;
|
||||
}
|
||||
delete[] c;
|
||||
}
|
||||
|
||||
int cuda_hierarchy(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* hierarchy_results
|
||||
){
|
||||
int cuda_ret = cudaSuccess;
|
||||
int EG_ret = EG_GPU_SUCC;
|
||||
int min_grid_size = 0;
|
||||
int block_size = 0;
|
||||
|
||||
cudaOccupancyMaxPotentialBlockSize(&min_grid_size, &block_size, calculate_hierarchy, 0, 0);
|
||||
int grid_size = (num_nodes + block_size - 1) / block_size;
|
||||
|
||||
int* d_V;
|
||||
int* d_E;
|
||||
int* d_row;
|
||||
int* d_col;
|
||||
double* d_W;
|
||||
int* d_node_mask;
|
||||
double* d_hierarchy_results;
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_V, (num_nodes+1) * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_E, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_row, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_col, num_edges * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_W, num_edges * sizeof(double)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_node_mask, num_nodes * sizeof(int)));
|
||||
EXIT_IF_CUDA_FAILED(cudaMalloc((void**)&d_hierarchy_results, num_nodes * sizeof(double)));
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_V, V, (num_nodes+1) * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_E, E, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_row, row, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_col, col, num_edges * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_node_mask, node_mask, num_nodes * sizeof(int), cudaMemcpyHostToDevice));
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(d_W, W, num_edges * sizeof(double), cudaMemcpyHostToDevice));
|
||||
|
||||
if(is_directed){
|
||||
directed_calculate_hierarchy<<<grid_size, block_size>>>(d_V, d_E, d_row, d_col, d_W, num_nodes, num_edges, d_node_mask, d_hierarchy_results);
|
||||
}else{
|
||||
int block_size = 256;
|
||||
int grid_size = (num_nodes + NODES_PER_BLOCK - 1) / NODES_PER_BLOCK;
|
||||
int shared_memory_size = 2 * sizeof(double) * block_size;
|
||||
calculate_hierarchy<<<grid_size, block_size, shared_memory_size>>>(d_V, d_E, d_W, num_nodes, d_node_mask, d_hierarchy_results);
|
||||
}
|
||||
|
||||
EXIT_IF_CUDA_FAILED(cudaMemcpy(hierarchy_results, d_hierarchy_results, num_nodes * sizeof(double), cudaMemcpyDeviceToHost));
|
||||
exit:
|
||||
|
||||
cudaFree(d_V);
|
||||
cudaFree(d_E);
|
||||
cudaFree(d_row);
|
||||
cudaFree(d_col);
|
||||
cudaFree(d_W);
|
||||
cudaFree(d_node_mask);
|
||||
cudaFree(d_hierarchy_results);
|
||||
if (cuda_ret != cudaSuccess) {
|
||||
switch (cuda_ret) {
|
||||
case cudaErrorMemoryAllocation:
|
||||
EG_ret = EG_GPU_FAILED_TO_ALLOCATE_DEVICE_MEM;
|
||||
break;
|
||||
default:
|
||||
EG_ret = EG_GPU_DEVICE_ERR;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
return EG_ret;
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,19 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
int cuda_hierarchy(
|
||||
_IN_ const int* V,
|
||||
_IN_ const int* E,
|
||||
_IN_ const int* row,
|
||||
_IN_ const int* col,
|
||||
_IN_ const double* W,
|
||||
_IN_ int num_nodes,
|
||||
_IN_ int num_edges,
|
||||
_IN_ bool is_directed,
|
||||
_IN_ int* node_mask,
|
||||
_OUT_ double* hierarchy_results
|
||||
);
|
||||
} // namespace gpu_easygraph
|
||||
@@ -0,0 +1,100 @@
|
||||
#include <vector>
|
||||
|
||||
#include "./common/err.h"
|
||||
|
||||
namespace gpu_easygraph {
|
||||
|
||||
int closeness_centrality(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<double>& W,
|
||||
const std::vector<int>& sources,
|
||||
std::vector<double>& CC
|
||||
);
|
||||
|
||||
|
||||
|
||||
int betweenness_centrality(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<double>& W,
|
||||
const std::vector<int>& sources,
|
||||
bool is_directed,
|
||||
bool normalized,
|
||||
bool endpoints,
|
||||
std::vector<double>& BC
|
||||
);
|
||||
|
||||
|
||||
|
||||
int k_core(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
std::vector<int>& KC
|
||||
);
|
||||
|
||||
|
||||
|
||||
int sssp_dijkstra(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<double>& W,
|
||||
const std::vector<int>& sources,
|
||||
int target,
|
||||
std::vector<double>& res
|
||||
);
|
||||
|
||||
|
||||
|
||||
int pagerank(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
double alpha,
|
||||
int max_iter_num,
|
||||
double threshold,
|
||||
std::vector<double>& PR
|
||||
);
|
||||
|
||||
|
||||
|
||||
int constraint(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<int>& row,
|
||||
const std::vector<int>& col,
|
||||
int num_nodes,
|
||||
const std::vector<double>& W,
|
||||
bool is_directed,
|
||||
std::vector<int>& node_mask,
|
||||
std::vector<double>& constraint
|
||||
);
|
||||
|
||||
|
||||
|
||||
int hierarchy(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<int>& row,
|
||||
const std::vector<int>& col,
|
||||
int num_nodes,
|
||||
const std::vector<double>& W,
|
||||
bool is_directed,
|
||||
std::vector<int>& node_mask,
|
||||
std::vector<double>& hierarchy
|
||||
);
|
||||
|
||||
|
||||
|
||||
int effective_size(
|
||||
const std::vector<int>& V,
|
||||
const std::vector<int>& E,
|
||||
const std::vector<int>& row,
|
||||
const std::vector<int>& col,
|
||||
int num_nodes,
|
||||
const std::vector<double>& W,
|
||||
bool is_directed,
|
||||
std::vector<int>& node_mask,
|
||||
std::vector<double>& effective_size
|
||||
);
|
||||
|
||||
} // namespace gpu_easygraph
|
||||
Reference in New Issue
Block a user