Rewriting large portion for simpler approach

5 files changed, 210 insertions, 26 deletions

diff --git a/examples/heterogeneous_computing/.nix/derivation.nix b/examples/heterogeneous_computing/.nix/derivation.nix
new file mode 100644
index 0000000..48dfc7a
--- /dev/null
+++ b/examples/heterogeneous_computing/.nix/derivation.nix
@@ -0,0 +1,39 @@
+{ lib
+, stdenv
+, scons
+, clang-tools
+, forstio
+, pname
+, version
+, kel-lbm
+}:
+
+stdenv.mkDerivation {
+	pname = pname + "-examples-" + "heterogeneous_computing";
+	inherit version;
+	src = ./..;
+
+	nativeBuildInputs = [
+		scons
+		clang-tools
+	];
+
+	buildInputs = [
+		forstio.core
+		forstio.async
+		forstio.io
+		forstio.io_codec
+		forstio.codec
+		forstio.codec-unit
+		forstio.codec-json
+		forstio.remote
+		forstio.remote-filesystem
+		forstio.remote-io
+		forstio.remote-sycl
+		kel-lbm.core
+	];
+
+	preferLocalBuild = true;
+
+	outputs = [ "out" "dev" ];
+}
diff --git a/examples/heterogeneous_computing/SConscript b/examples/heterogeneous_computing/SConscript
new file mode 100644
index 0000000..1e52d88
--- /dev/null
+++ b/examples/heterogeneous_computing/SConscript
@@ -0,0 +1,34 @@
+#!/bin/false
+
+import os
+import os.path
+import glob
+
+
+Import('env')
+
+dir_path = Dir('.').abspath
+
+# Environment for base library
+examples_env = env.Clone();
+examples_env['CXX'] = 'syclcc-clang';
+examples_env['CXXFLAGS'] += ['-O3'];
+
+examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp"))
+examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp"))
+
+env.sources += examples_env.sources;
+env.headers += examples_env.headers;
+
+# Cavity2D
+examples_objects = [];
+examples_env.add_source_files(examples_objects, ['heterogeneous_computing.cpp'], shared=False);
+examples_env.heterogeneous_computing = examples_env.Program('#bin/heterogeneous_computing', [examples_objects]);
+
+# Set Alias
+env.examples = [
+	examples_env.heterogeneous_computing
+];
+env.Alias('examples', env.examples);
+env.targets += ['examples'];
+env.Install('$prefix/bin/', env.examples);
diff --git a/examples/heterogeneous_computing/SConstruct b/examples/heterogeneous_computing/SConstruct
new file mode 100644
index 0000000..fe206e1
--- /dev/null
+++ b/examples/heterogeneous_computing/SConstruct
@@ -0,0 +1,82 @@
+#!/usr/bin/env python3
+
+import sys
+import os
+import os.path
+import glob
+import re
+
+
+if sys.version_info < (3,):
+    def isbasestring(s):
+        return isinstance(s,basestring)
+else:
+    def isbasestring(s):
+        return isinstance(s, (str,bytes))
+
+def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""):
+
+    if isbasestring(filetype):
+        dir_path = self.Dir('.').abspath
+        filetype = sorted(glob.glob(dir_path+"/"+filetype))
+
+    for path in filetype:
+        target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path )
+        if shared:
+            target_name+='.os'
+            sources.append( self.SharedObject( target=target_name, source=path ) )
+        else:
+            target_name+='.o'
+            sources.append( self.StaticObject( target=target_name, source=path ) )
+    pass
+
+def isAbsolutePath(key, dirname, env):
+	assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,)
+
+env_vars = Variables(
+	args=ARGUMENTS
+)
+
+env_vars.Add('prefix',
+	help='Installation target location of build results and headers',
+	default='/usr/local/',
+	validator=isAbsolutePath
+)
+
+env_vars.Add('build_examples',
+	help='If examples should be built',
+	default="true"
+)
+
+env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[],
+    CPPDEFINES=['SAW_UNIX'],
+    CXXFLAGS=[
+        '-std=c++20',
+        '-g',
+        '-Wall',
+        '-Wextra',
+        '-isystem', 'AdaptiveCpp'
+    ],
+    LIBS=[
+        'forstio-core',
+        'acpp-rt',
+        'omp'
+    ]
+);
+env.__class__.add_source_files = add_kel_source_files
+env.Tool('compilation_db');
+env.cdb = env.CompilationDatabase('compile_commands.json');
+
+env.objects = [];
+env.sources = [];
+env.headers = [];
+env.targets = [];
+
+Export('env')
+SConscript('SConscript')
+
+env.Alias('cdb', env.cdb);
+env.Alias('all', [env.targets]);
+env.Default('all');
+
+env.Alias('install', '$prefix')
diff --git a/examples/heterogeneous_computing/heterogeneous_computing.cpp b/examples/heterogeneous_computing/heterogeneous_computing.cpp
new file mode 100644
index 0000000..990652a
--- /dev/null
+++ b/examples/heterogeneous_computing/heterogeneous_computing.cpp
@@ -0,0 +1,33 @@
+#include <forstio/error.hpp>
+#include <iostream>
+
+namespace kel {
+namespace sch {
+using namespace saw::schema;
+using KelConfig = Struct<
+	Member<String,"resolution">
+>;
+}
+
+saw::error_or<void> real_main(int argc, char** argv){
+	return saw::make_void();
+}
+}
+
+int main(int argc, char** argv){
+	auto eov = kel::kel_main(argc, argv);
+	if(eov.is_error()){
+		auto& err = eov.get_error();
+		auto err_msg = err.get_message();
+		std::cerr<<"[Error]: "<<err.get_category();
+
+		if(not err_msg.empty()){
+			std::cerr<<" - "<<err_msg;
+		}
+		std::cerr<<std::endl;
+
+		return err.get_id();
+	}
+
+	return 0;
+}
diff --git a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
index 56bcdaf..6dda469 100644
--- a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
+++ b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp
@@ -70,6 +70,10 @@ saw::error_or<void> set_geometry(
 	acpp::sycl::queue& sycl_q
 ){
 	using namespace kel::lbm;
+	
+	saw::data<sch::T> rho{1.0};
+	saw::data<sch::FixedArray<sch::T,sch::D2Q9::D>> vel{{0.0,0.0}};
+	auto eq = equilibrium<sch::T,sch::D2Q9>(rho, vel);
 
 	sycl_q.submit([&](acpp::sycl::handler& h){
 		h.parallel_for(acpp::sycl::range<2>{meta.at({0}).get(), meta.at({1}).get()},[=](acpp::sycl::id<2> idx){
@@ -79,6 +83,8 @@ saw::error_or<void> set_geometry(
 			size_t acc_id = j * meta.at({0u}).get() + i;
 			auto& c = cells[acc_id];
 			auto& info = c.template get<"info">()({0});
+			auto& dfs = c.template get<"dfs">();
+			auto& dfs_old = c.template get<"dfs_old">();
 
 			if(i >= 2u and j >= 2u and (i+2u) < meta.at({0u}).get() and (j+2u) < meta.at({1u}).get()){
 				// Fluid
@@ -95,23 +101,16 @@ saw::error_or<void> set_geometry(
 			}else {
 				info.set({0u});
 			}
+			for(saw::data<sch::UInt64> k{0u}; k < saw::data<sch::UInt64>{Desc::Q}; ++k){
+				dfs(k) = eq.at(k);
+				dfs_old(k) = eq.at(k);
+			}
 		});
 	});
 
-	return saw::make_void();
-}
-
-void set_initial_conditions(
-	saw::data<sch::CellStruct>* cells,
-	const saw::data<sch::FixedArray<sch::UInt64,sch::D2Q9::D>>& meta,
-	acpp::sycl::queue& sycl_q
-){
-	using namespace kel::lbm;
-
-	saw::data<sch::T> rho{1.0};
-	saw::data<sch::FixedArray<sch::T,sch::D2Q9::D>> vel{{0.0,0.0}};
-	auto eq = equilibrium<sch::T,sch::D2Q9>(rho, vel);
+	sycl_q.wait();
 
+	return saw::make_void();
 }
 
 void lbm_step(
@@ -178,26 +177,23 @@ saw::error_or<void> kel_main(int argc, char** argv){
 		}
 	}
 
-	sycl_q.wait();
-	sycl_q.memcpy(&host_cells[0u], cells, x_d * y_d * sizeof(saw::data<lbm::sch::CellStruct>) );
-	sycl_q.wait();
-	acpp::sycl::free(cells, sycl_q);
-	sycl_q.wait();
-
 	std::string vtk_f_name{"tmp/poiseulle_2d_gpu_"};
 	vtk_f_name += std::to_string(0u) + ".vtk";
 	// write_vtk_file(vtk_f_name,host_cells);
 
-	for(uint64_t i = 0u; i < x_d; ++i){
-		for(uint64_t j = 0u; j < y_d; ++j){
+	for(uint64_t i = 0u; i < 1024u*1204u; ++i){
+		if(i%128u == 0u){
 
-			size_t acc_id = j * x_d + i;
-
-			std::cout<<static_cast<uint64_t>(host_cells.at(acc_id).template get<"info">()({0u}).get())<<" ";
 		}
-		std::cout<<"\n";
+		lbm_step(cells,meta,i,sycl_q);
+
 	}
-	std::cout<<std::endl;
+
+	sycl_q.wait();
+	sycl_q.memcpy(&host_cells[0u], cells, x_d * y_d * sizeof(saw::data<lbm::sch::CellStruct>) );
+	sycl_q.wait();
+	acpp::sycl::free(cells, sycl_q);
+	sycl_q.wait();
 
 	return saw::make_void();
 }