6 files changed, 187 insertions, 3 deletions
diff --git a/lid_driven_cavity/cpp/shell.nix b/lid_driven_cavity/cpp/shell.nix
new file mode 100644
index 0000000..d2d157d
--- /dev/null
+++ b/lid_driven_cavity/cpp/shell.nix
@@ -0,0 +1,40 @@
+{ pkgs ? import <nixpkgs> { }, ... }:
+
+pkgs.stdenvNoCC.mkDerivation rec {
+  name = "boltzgen-env";
+  env = pkgs.buildEnv { name = name; paths = buildInputs; };
+
+  buildInputs = let
+    boltzgen = pkgs.python3.pkgs.buildPythonPackage rec {
+      pname = "boltzgen";
+      version = "0.1";
+
+      src = builtins.fetchGit {
+        url = "https://code.kummerlaender.eu/boltzgen/";
+        rev = "4a2885ad3ae0396486d288df94339d0c45e6db8b";
+      };
+
+      propagatedBuildInputs = with pkgs.python37Packages; [
+        sympy
+        numpy
+        Mako
+      ];
+
+      doCheck = false;
+    };
+
+    local-python = pkgs.python3.withPackages (python-packages: with python-packages; [
+      boltzgen
+    ]);
+
+  in with pkgs; [
+    local-python
+    gcc9
+    cmake
+  ];
+
+  shellHook = ''
+    export NIX_SHELL_NAME="${name}"
+    export PYTHONPATH="$PWD:$PYTHONPATH"
+  '';
+}
diff --git a/lid_driven_cavity/cuda/shell.nix b/lid_driven_cavity/cuda/shell.nix
new file mode 100644
index 0000000..86f707f
--- /dev/null
+++ b/lid_driven_cavity/cuda/shell.nix
@@ -0,0 +1,42 @@
+{ pkgs ? import <nixpkgs> { }, ... }:
+
+pkgs.stdenvNoCC.mkDerivation rec {
+  name = "boltzgen-env";
+  env = pkgs.buildEnv { name = name; paths = buildInputs; };
+
+  buildInputs = let
+    boltzgen = pkgs.python3.pkgs.buildPythonPackage rec {
+      pname = "boltzgen";
+      version = "0.1";
+
+      src = builtins.fetchGit {
+        url = "https://code.kummerlaender.eu/boltzgen/";
+        rev = "4a2885ad3ae0396486d288df94339d0c45e6db8b";
+      };
+
+      propagatedBuildInputs = with pkgs.python37Packages; [
+        sympy
+        numpy
+        Mako
+      ];
+
+      doCheck = false;
+    };
+
+    local-python = pkgs.python3.withPackages (python-packages: with python-packages; [
+      boltzgen
+    ]);
+
+  in with pkgs; [
+    local-python
+    cmake
+    cudatoolkit
+    linuxPackages.nvidia_x11 
+  ];
+
+  shellHook = ''
+    export NIX_SHELL_NAME="${name}"
+    export CUDA_PATH="${pkgs.cudatoolkit}"
+    export PYTHONPATH="$PWD:$PYTHONPATH"
+  '';
+}
diff --git a/lid_driven_cavity/opencl/AA.py b/lid_driven_cavity/opencl/AA.py
index f317cd0..74979c5 100644
--- a/lid_driven_cavity/opencl/AA.py
+++ b/lid_driven_cavity/opencl/AA.py
@@ -83,7 +83,7 @@ class Lattice:
         moments = numpy.ndarray(shape=(self.memory.volume*(self.descriptor.d+1),1), dtype=self.float_type[0])
 
         self.program.collect_moments_all(
-            self.queue, self.geometry.size(), self.layout, self.memory.cl_pop, self.memory.cl_moments)
+            self.queue, self.geometry.inner_size(), self.layout, self.memory.cl_pop, self.memory.cl_moments)
 
         cl.enqueue_copy(self.queue, moments, self.memory.cl_moments).wait();
 
@@ -99,7 +99,7 @@ __kernel void equilibrilize_all(__global ${float_type}* f_next)
 __kernel void collect_moments_all(__global ${float_type}* f,
                                   __global ${float_type}* moments)
 {
-    const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
+    const unsigned int gid = ${index.gid('get_global_id(0)+1', 'get_global_id(1)+1')};
     collect_moments_tock(f, gid, moments);
 }
 """
diff --git a/lid_driven_cavity/opencl/ldc_2d.py b/lid_driven_cavity/opencl/ldc_2d.py
index 44802ef..7380aa9 100644
--- a/lid_driven_cavity/opencl/ldc_2d.py
+++ b/lid_driven_cavity/opencl/ldc_2d.py
@@ -10,7 +10,7 @@ from common import CellList, generate_moment_plots
 nUpdates = 100000
 nStat    = 10000
 
-geometry = Geometry(512, 512)
+geometry = Geometry(256, 256)
 
 print("Generating kernel using boltzgen...\n")
 
diff --git a/lid_driven_cavity/opencl/shell.nix b/lid_driven_cavity/opencl/shell.nix
new file mode 100644
index 0000000..de9d59f
--- /dev/null
+++ b/lid_driven_cavity/opencl/shell.nix
@@ -0,0 +1,43 @@
+{ pkgs ? import <nixpkgs> { }, ... }:
+
+pkgs.stdenvNoCC.mkDerivation rec {
+  name = "boltzgen-env";
+  env = pkgs.buildEnv { name = name; paths = buildInputs; };
+
+  buildInputs = let
+    boltzgen = pkgs.python3.pkgs.buildPythonPackage rec {
+      pname = "boltzgen";
+      version = "0.1";
+
+      src = builtins.fetchGit {
+        url = "https://code.kummerlaender.eu/boltzgen/";
+        rev = "4a2885ad3ae0396486d288df94339d0c45e6db8b";
+      };
+
+      propagatedBuildInputs = with pkgs.python37Packages; [
+        sympy
+        numpy
+        Mako
+      ];
+
+      doCheck = false;
+    };
+
+    local-python = pkgs.python3.withPackages (python-packages: with python-packages; [
+      boltzgen
+      numpy
+      pyopencl setuptools
+      matplotlib
+    ]);
+
+  in with pkgs; [
+    local-python
+    opencl-info
+  ];
+
+  shellHook = ''
+    export NIX_SHELL_NAME="${name}"
+    export PYOPENCL_COMPILER_OUTPUT=1
+    export PYTHONPATH="$PWD:$PYTHONPATH"
+  '';
+}
diff --git a/lid_driven_cavity/opencl_gl_interop/shell.nix b/lid_driven_cavity/opencl_gl_interop/shell.nix
new file mode 100644
index 0000000..42feaaf
--- /dev/null
+++ b/lid_driven_cavity/opencl_gl_interop/shell.nix
@@ -0,0 +1,59 @@
+{ pkgs ? import <nixpkgs> { }, ... }:
+
+pkgs.stdenvNoCC.mkDerivation rec {
+  name = "boltzgen-env";
+  env = pkgs.buildEnv { name = name; paths = buildInputs; };
+
+  buildInputs = let
+    custom-python = let
+      packageOverrides = self: super: {
+        pyopencl = super.pyopencl.overridePythonAttrs(old: rec {
+          buildInputs = with pkgs; [
+            opencl-headers ocl-icd python37Packages.pybind11
+            libGLU_combined
+          ];
+        # Enable OpenGL integration and fix build
+          preBuild = ''
+            python configure.py --cl-enable-gl
+            export HOME=/tmp/pyopencl
+          '';
+        });
+      };
+    in pkgs.python3.override { inherit packageOverrides; };
+
+    boltzgen = pkgs.python3.pkgs.buildPythonPackage rec {
+      pname = "boltzgen";
+      version = "0.1";
+
+      src = builtins.fetchGit {
+        url = "https://code.kummerlaender.eu/boltzgen/";
+        rev = "4a2885ad3ae0396486d288df94339d0c45e6db8b";
+      };
+
+      propagatedBuildInputs = with pkgs.python37Packages; [
+        sympy
+        numpy
+        Mako
+      ];
+
+      doCheck = false;
+    };
+
+    local-python = custom-python.withPackages (python-packages: with python-packages; [
+      boltzgen
+      numpy
+      pyopencl setuptools
+      pyopengl pyrr
+    ]);
+
+  in with pkgs; [
+    local-python
+    opencl-info
+  ];
+
+  shellHook = ''
+    export NIX_SHELL_NAME="${name}"
+    export PYOPENCL_COMPILER_OUTPUT=1
+    export PYTHONPATH="$PWD:$PYTHONPATH"
+  '';
+}