Skip to content
Snippets Groups Projects
Commit d907c4ec authored by Recolic Keghart's avatar Recolic Keghart
Browse files

insert debug things

parent a605807b
No related branches found
No related tags found
No related merge requests found
......@@ -5,6 +5,7 @@
#include <thrust/remove.h>
#include <thrust/sort.h>
#include <cub/cub.cuh>
#include "rdebug.hpp"
#define cErr(errcode) { gpuAssert((errcode), __FILE__, __LINE__); }
__inline__ __host__ __device__
......@@ -736,10 +737,12 @@ void significant_insert(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &upd
__host__
void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_values) {
SIZE_TYPE ous = update_keys.size();
LOG_TIME("enter_update_gpma")
// step1: sort update keys with values
thrust::sort_by_key(update_keys.begin(), update_keys.end(), update_values.begin());
cErr(cudaDeviceSynchronize());
LOG_TIME("1-2")
// step2: get leaf node of each update (execute del and mod)
DEV_VEC_SIZE update_nodes(update_keys.size());
......@@ -747,6 +750,7 @@ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_val
locate_leaf_batch(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), gpma.keys.size(), gpma.segment_length, gpma.tree_height,
RAW_PTR(update_keys), RAW_PTR(update_values), update_keys.size(), RAW_PTR(update_nodes));
cErr(cudaDeviceSynchronize());
LOG_TIME("2-3")
// step3: extract insertions
DEV_VEC_SIZE unique_update_nodes(update_keys.size());
......@@ -757,6 +761,7 @@ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_val
compact_insertions(update_nodes, update_keys, update_values, update_size);
compress_insertions_by_node(update_nodes, update_size, unique_update_nodes, update_offset, unique_node_size);
cErr(cudaDeviceSynchronize());
LOG_TIME("3-4")
// step4: rebuild for significant update
int threshold = 5 * 1000 * 1000;
......@@ -764,6 +769,7 @@ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_val
significant_insert(gpma, update_keys, update_values, update_size);
return;
}
LOG_TIME("4-5")
// step5: rebalance each tree level
for (SIZE_TYPE level = 0; level <= gpma.tree_height && update_size; level++) {
......@@ -785,6 +791,7 @@ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_val
compress_insertions_by_node(update_nodes, update_size, unique_update_nodes, update_offset,
unique_node_size);
}
LOG_TIME("5-6")
// step6: rebalance the root node if necessary
if (update_size > 0) {
......@@ -800,6 +807,7 @@ void update_gpma(GPMA &gpma, DEV_VEC_KEY &update_keys, DEV_VEC_VALUE &update_val
RAW_PTR(update_keys), RAW_PTR(update_values), update_size, RAW_PTR(unique_update_nodes),
RAW_PTR(update_offset), unique_node_size, lower_bound, upper_bound, RAW_PTR(gpma.row_offset));
}
LOG_TIME("6-7 LEAVE update_gpma")
cErr(cudaDeviceSynchronize());
}
......
......@@ -279,6 +279,8 @@ void gpma_bfs(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE *row_offsets,
SIZE_TYPE level = 1;
const SIZE_TYPE THREADS_NUM = 256;
while (true) {
printf("DEBUG: gpma_bfs::ITERATION-\n");
// gather
SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, host_num[0]);
host_num[0] = 0;
......
......@@ -61,19 +61,23 @@ int main(int argc, char **argv) {
int num_slide = 100;
int step = half / num_slide;
LOG_TIME("before init_csr_gpma")
GPMA gpma;
init_csr_gpma(gpma, node_size);
cudaDeviceSynchronize();
LOG_TIME("before update_gpma 1")
update_gpma(gpma, base_keys, base_values);
thrust::device_vector<SIZE_TYPE> bfs_result(node_size);
cudaDeviceSynchronize();
LOG_TIME("before first bfs")
gpma_bfs(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(gpma.row_offset), node_size,
edge_size, bfs_start_node, RAW_PTR(bfs_result));
int reach_nodes = node_size - thrust::count(bfs_result.begin(), bfs_result.end(), 0);
printf("start from node %d, number of reachable nodes: %d\n", bfs_start_node, reach_nodes);
LOG_TIME("before main loop")
for (int i = 0; i < num_slide; i++) {
thrust::host_vector<KEY_TYPE> hk(step * 2);
for (int j = 0; j < step; j++) {
......@@ -95,11 +99,13 @@ int main(int argc, char **argv) {
cudaDeviceSynchronize();
}
printf("Graph is updated.\n");
LOG_TIME("before second bfs")
gpma_bfs(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(gpma.row_offset), node_size,
edge_size, bfs_start_node, RAW_PTR(bfs_result));
reach_nodes = node_size - thrust::count(bfs_result.begin(), bfs_result.end(), 0);
printf("start from node %d, number of reachable nodes: %d\n", bfs_start_node, reach_nodes);
LOG_TIME("after second loop")
return 0;
}
#ifndef RLIB_IMPL_GPMA_DEBUG_HPP_
#define RLIB_IMPL_GPMA_DEBUG_HPP_ 1
#include <time.h>
#include <cstdio>
#include <cstdint>
inline int64_t get_time_in_us() {
struct timespec t;
clock_gettime(CLOCK_MONOTONIC, &t);
return ((int64_t)(t.tv_sec) * (int64_t)1000000000 + (int64_t)(t.tv_nsec)) / 1000;
}
#define LOG_TIME(msg) printf("T+%lld - " msg "\n", get_time_in_us());
#endif
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment