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

adjust scripts for gpma_bfs<CPU>, but discovered update_gpma<CPU> bug

parent e6fbb568
No related branches found
No related tags found
1 merge request!2Bfs cpu
......@@ -28,7 +28,7 @@ with open('soc-pokec-relationships.txt', 'r') as f:
node_size = max(node_size, b)
edge_size = len(edges)
random.shuffle(edges)
with open('pokec.txt', 'w') as f:
with open('$NODE_MAX.pokec.txt', 'w') as f:
f.write('{} {}\n'.format(node_size, edge_size))
for a, b in edges:
f.write('{} {}\n'.format(a - 1, b - 1))
......
......@@ -80,11 +80,25 @@ class GPMA {
init_gpma_members(row_num);
init_gpma_insertions();
}
GPMA() = default;
void print_status(std::string prefix = "DBG") const {
DEBUG_PRINTFLN(prefix + ": GPMA_DUMP: keys={}, values={}, row_offset={}, seg_length,tree_height,row_num={},{},{}", keys.size(), values.size(), row_offset.size(), segment_length, tree_height, row_num);
DEBUG_PRINTFLN(prefix + ": GPMA_DUMP: keys={}, values={}, row_offset={}", rlib::printable_iter(keys), rlib::printable_iter(values_for_print(values)), rlib::printable_iter(row_offset));
}
auto mirror() const {
std::conditional_t<DEV == GPU, GPMA<CPU>, GPMA<GPU>> result;
result.keys = keys;
result.values = values;
result.row_offset = row_offset;
result.row_num = row_num;
result.segment_length = segment_length;
result.tree_height = tree_height;
result.lower_element = lower_element;
result.upper_element = upper_element;
return result;
}
};
/* returns the index of highest nonzero bit.
......
......@@ -281,7 +281,12 @@ __host__ void gpma_bfs(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE *row_offset
SIZE_TYPE BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, host_num[0]);
host_num[0] = 0;
anyMemcpy<CPU, DEV>(edge_queue_offset, host_num, sizeof(SIZE_TYPE));
gpma_bfs_gather_kernel<THREADS_NUM><<<BLOCKS_NUM, THREADS_NUM>>>(node_queue, node_queue_offset, edge_queue, edge_queue_offset, keys, values, row_offsets);
if(DEV == GPU) {
gpma_bfs_gather_kernel<THREADS_NUM><<<BLOCKS_NUM, THREADS_NUM>>>(node_queue, node_queue_offset, edge_queue, edge_queue_offset, keys, values, row_offsets);
}
else {
}
// contract
level++;
......@@ -289,7 +294,12 @@ __host__ void gpma_bfs(KEY_TYPE *keys, VALUE_TYPE *values, SIZE_TYPE *row_offset
anyMemcpy<DEV, CPU>(host_num, edge_queue_offset, sizeof(SIZE_TYPE));
BLOCKS_NUM = CALC_BLOCKS_NUM(THREADS_NUM, host_num[0]);
gpma_bfs_contract_kernel<THREADS_NUM><<<BLOCKS_NUM, THREADS_NUM>>>(edge_queue, edge_queue_offset, node_queue, node_queue_offset, level, results, bitmap);
if(DEV == GPU) {
gpma_bfs_contract_kernel<THREADS_NUM><<<BLOCKS_NUM, THREADS_NUM>>>(edge_queue, edge_queue_offset, node_queue, node_queue_offset, level, results, bitmap);
}
else {
}
anyMemcpy<DEV, CPU>(host_num, node_queue_offset, sizeof(SIZE_TYPE));
if (0 == host_num[0])
......
......@@ -53,15 +53,15 @@ int main(int argc, char **argv) {
h_base_keys[i] = ((KEY_TYPE)host_x[i] << 32) + host_y[i];
}
NATIVE_VEC_KEY<GPU> base_keys = h_base_keys;
NATIVE_VEC_VALUE<GPU> base_values(half, 1);
NATIVE_VEC_KEY<CPU> base_keys = h_base_keys;
NATIVE_VEC_VALUE<CPU> base_values(half, 1);
cudaDeviceSynchronize();
int num_slide = 100;
int step = half / num_slide;
LOG_TIME("before init_csr_gpma")
GPMA<GPU> gpma(node_size);
GPMA<CPU> gpma(node_size);
cudaDeviceSynchronize();
LOG_TIME("before update_gpma 1")
......@@ -70,7 +70,10 @@ int main(int argc, char **argv) {
cudaDeviceSynchronize();
LOG_TIME("before first bfs")
gpma_bfs<GPU>(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(gpma.row_offset), node_size, edge_size, bfs_start_node, RAW_PTR(bfs_result));
{
auto gpma_mirror = gpma.mirror();
gpma_bfs<GPU>(RAW_PTR(gpma_mirror.keys), RAW_PTR(gpma_mirror.values), RAW_PTR(gpma_mirror.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);
......@@ -86,10 +89,10 @@ int main(int argc, char **argv) {
hk[j + step] = ((KEY_TYPE)host_x[idx] << 32) + host_y[idx];
}
NATIVE_VEC_VALUE<GPU> update_values(step * 2);
NATIVE_VEC_VALUE<CPU> update_values(step * 2);
thrust::fill(update_values.begin(), update_values.begin() + step, 1);
thrust::fill(update_values.begin() + step, update_values.end(), VALUE_NONE);
NATIVE_VEC_KEY<GPU> update_keys = hk;
NATIVE_VEC_KEY<CPU> update_keys = hk;
cudaDeviceSynchronize();
update_gpma(gpma, update_keys, update_values);
......@@ -98,7 +101,10 @@ int main(int argc, char **argv) {
printf("Graph is updated.\n");
LOG_TIME("before second bfs")
gpma_bfs<GPU>(RAW_PTR(gpma.keys), RAW_PTR(gpma.values), RAW_PTR(gpma.row_offset), node_size, edge_size, bfs_start_node, RAW_PTR(bfs_result));
{
auto gpma_mirror = gpma.mirror();
gpma_bfs<GPU>(RAW_PTR(gpma_mirror.keys), RAW_PTR(gpma_mirror.values), RAW_PTR(gpma_mirror.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 bfs")
......
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