Skip to content

Commence Phase 2 #48

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 32 commits into
base: WorkingBranch
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
32 commits
Select commit Hold shift + click to select a range
5fdd9c9
add new starplat programs. update gitignore
lazyprop Jun 10, 2023
2f5d2b3
update readme. add todo sectoin
lazyprop Jun 10, 2023
a29d50f
add license and results
lazyprop Jun 11, 2023
bff6967
update todo
lazyprop Jun 11, 2023
fa63b89
add parser task to todo
lazyprop Jun 11, 2023
cd32226
setup benchmark, fix formatting and add readme
lazyprop Jun 11, 2023
fa58bb4
change adjacency list to use std::vector instead of std::map
lazyprop Jun 13, 2023
e24efde
mmap input file for faster parsing
lazyprop Jun 13, 2023
18b1530
update readme and cleanup
lazyprop Jun 13, 2023
bf362c1
Merge branch 'opencl-dev' into parser-dev
lazyprop Jun 14, 2023
a9748d2
Merge pull request #1 from lazyprop/parser-dev
lazyprop Jun 14, 2023
7e85fd2
updated bench
lazyprop Jun 15, 2023
47da432
set edge weight while parsing
lazyprop Jun 15, 2023
32a117a
Merge branch 'parser-dev' into opencl-dev
lazyprop Jun 15, 2023
138bc64
clang-format
lazyprop Jun 17, 2023
15a4c69
don't double copy the csr
lazyprop Jun 17, 2023
7165ab7
create list of csr array names to reduce redundancy
lazyprop Jun 17, 2023
fa7461c
minor typo fix
lazyprop Jun 17, 2023
d6160c7
update gitignore
lazyprop Jun 24, 2023
2badda5
create output folder in directory of source for better debugging
lazyprop Jun 24, 2023
51c8911
run clang-formt on dsl_cpp_generator amd backend
lazyprop Jun 24, 2023
8b70631
added debugutils. implemented << for Type.
lazyprop Jun 25, 2023
f582cb9
parses dsu declaration
lazyprop Aug 2, 2023
6e71ebe
updated makefile. now printing proc_call
lazyprop Aug 2, 2023
743bac5
better error messages for proccall codegen
lazyprop Aug 2, 2023
55bec30
make flex parser verboes
lazyprop Aug 2, 2023
dbe1f2d
now generating main kernel
lazyprop Aug 2, 2023
f80e2b9
generate calls for dsu methods
lazyprop Aug 2, 2023
eaa6239
commence phase 2. add support for union-find. add boruvka's algorithm
lazyprop Oct 23, 2023
39a69d8
fix typo
lazyprop Dec 8, 2023
6c1328b
re-add borvuka
lazyprop Jan 14, 2024
be9b3b1
rename boruvka
lazyprop Jan 14, 2024
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions .editorconfig
Original file line number Diff line number Diff line change
@@ -0,0 +1,6 @@
root = true

[*.{c,cpp,h,hpp}]
tab_width = 2
indent_size = 2
indent_style = space
8 changes: 7 additions & 1 deletion .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -21,4 +21,10 @@ dump/
*.starplat
runtests.sh
cudadump/
testcases/testAMD/*
testcases/testAMD/*
test-graphs/
graphcode/graphlib-bench/inp.txt
test-graphs.zip
graphcode/phase2/sssp-out
graphcode/phase2/boruvka-out

34 changes: 29 additions & 5 deletions README.md
Original file line number Diff line number Diff line change
@@ -1,11 +1,32 @@
# StarPlat: A Versatile DSL for Graph Analytics

# GraphDSL
[![CUDA Branch](https://github.com/nibeditabh/GraphDSL/actions/workflows/cuda.yml/badge.svg?branch=cuda)](https://github.com/nibeditabh/GraphDSL/actions/workflows/cuda.yml)


## TODO:

# StarPlat: A Versatile DSL for Graph Analytics
# GraphDSL MPI Backend
### High Priority

- [ ] Improve `parseGraph()` in `graphcode/graph.hpp`.
- [ ] The generated OpenCL code segfaults while releasing OpenCL/device
resources.
- [X] Write script for benchmarking generated OpenCL code.
[Low Priority: Extend it to support all backends]
- [ ] Make every file generated OpenCL program use `cl_assert.h` for better
debugging.
- [ ] Improve README. Update outdated instructions. Add instructions for
building OpenCL programs.


### Low Priority

- [ ] All generated files should be in a folder and not litter
`graphcode/generatedXYZ`.
- [ ] Every occurence of "AMD" in the OpenCL backend should be replaced by
"OpenCL".


## MPI Backend

### How to compile Generated MPI Codes on Aqua Cluster
```
Expand Down Expand Up @@ -142,8 +163,11 @@ Graph DSL for basic graph algorithms
- Phase1 (static) **SSSP, BC, PR, and TC**

## Acknowledgments
```
This project is funded by India's National Supercomputing Mission, whose timely funding and regular reviews are greatly appreciated.

[![License](http://img.shields.io/:license-mit-blue.svg?style=flat-square)](http://badges.mit-license.org)
## License

The StarPlat codebase is made available for academic and non-commercial use.
Commercial use of the code needs licensing. The interested may contact
[email protected].

51 changes: 51 additions & 0 deletions graphcode/generated_amd/sssp_dslV2.cl
Original file line number Diff line number Diff line change
@@ -0,0 +1,51 @@
#include "libOpenCL.h"
__kernel void initdist_kernel( __global int *d_dist , int value, int v){
unsigned int id = get_global_id(0);
if(id<v){
d_dist[id] = value;
}
}
__kernel void initmodified_kernel( __global int *d_modified , int value, int v){
unsigned int id = get_global_id(0);
if(id<v){
d_modified[id] = value;
}
}

__kernel void initIndexmodified_kernel(__global int *d_modified, int src , int val){
d_modified[src] = val ;
}

__kernel void initIndexdist_kernel(__global int *d_dist, int src , int val){
d_dist[src] = val ;
}
__kernel void initd_modified_next_kernel(__global int *d_modified_next, int val , int v){
unsigned int id = get_global_id(0);
if(id<v)
{
d_modified_next[id] = val;
}
}
__kernel void Compute_SSSP_kernel(int V, int E, __global int* d_meta, __global int* d_data, __global int* d_src,
__global int* d_weight,__global int* d_rev_meta ,__global int *finished,__global int * d_modified_next,__global int * d_modified,__global int * d_dist){ // BEGIN KER FUN via ADDKERNEL
unsigned int v = get_global_id(0);
if(v >= V) return;
if (d_modified[v] == true){ // if filter begin

//ForAll started here
for (int edge = d_meta[v]; edge < d_meta[v+1]; edge++) { // FOR NBR ITR
int nbr = d_data[edge];
int e = edge;
int dist_new = d_dist[v] + d_weight[e];
int modified_new = true;
if(d_dist[v]!= INT_MAX && d_dist[nbr] > dist_new)
{
atomic_min(&d_dist[nbr],dist_new);
d_modified_next[nbr] = modified_new;
*finished = false ;
}

} // end FOR NBR ITR. TMP FIX!

} // if filter end
} // end KER FUNC
Loading