Merge branch 'fhennig/runtime-integration' into 'v2.0-dev'

Move custom GPU atomics into separate namespace. Add a script to print the include path.

See merge request !477

src/pystencils/backend/platforms/cuda.py

@@ -109,8 +109,14 @@ class CudaPlatform(GenericGpu):
             cond = is_valid_thread
 
         # use atomic operation
+        match actual_reduction_op:
+            case ReductionOp.Min | ReductionOp.Max | ReductionOp.Mul:
+                impl_namespace = "pystencils::runtime::gpu::"
+            case _:
+                impl_namespace = ""
+
         func = CFunction(
-            f"atomic{actual_reduction_op.name}", [ptrtype, stype], PsCustomType("void")
+            f"{impl_namespace}atomic{actual_reduction_op.name}", [ptrtype, stype], PsCustomType("void")
         )
         func_args = (ptr_expr, symbol_expr)
 

src/pystencils/include/__main__.py

+import argparse
+
+
+def main():
+    parser = argparse.ArgumentParser(
+        "Print the include path for the pystencils runtime headers"
+    )
+    parser.add_argument(
+        "-I", dest="dash_i", action="store_true", help="Emit as `-I` compiler argument"
+    )
+    parser.add_argument(
+        "-s", dest="strip", action="store_true", help="Emit without trailing newline"
+    )
+    args = parser.parse_args()
+
+    from . import get_pystencils_include_path
+
+    include_path = get_pystencils_include_path()
+
+    if args.dash_i:
+        print(f"-I{include_path}", end="")
+    else:
+        end = "" if args.strip else "\n"
+        print(get_pystencils_include_path(), end=end)
+
+
+main()

src/pystencils/include/pystencils_runtime/bits/gpu_atomics.h

 #pragma once
 
+namespace pystencils::runtime::gpu {
+
 // No direct implementation for all atomic operations available
 // -> add support by custom implementations using a CAS mechanism
 
@@ -79,4 +81,6 @@ __device__ __forceinline__ float atomicMax(float *address, float val)
             break;
     }
     return __int_as_float(ret);
-}
\ No newline at end of file
+}
+
+} // pystencils::runtime::gpu

src/pystencils/jit/gpu_cupy.py

@@ -248,7 +248,7 @@ class CupyJit(JitBase):
         return CupyKernelWrapper(kernel, raw_kernel)
 
     def _compiler_options(self) -> tuple[str, ...]:
-        options = ["-w", "-std=c++11"]
+        options = ["-w", "-std=c++17"]
         options.append("-I" + get_pystencils_include_path())
         return tuple(options)