Moved documentation from external repo to be part of the code repo.

.readthedocs.yaml

@@ -0,0 +1,7 @@
+version: 2
+
+sphinx:
+  configuration: docs/sphinx/source/conf.py
+
+conda:
+  environment: docs/sphinx/environment.yml

README.md

@@ -22,6 +22,7 @@ Jungwon Kim, Seyong Lee, Beau Johnston, and Jeffrey S. Vetter. 2021. IRIS: A Por
   doi={}
 }
 ```
+
 # More Information
 
 If you would like more information on IRIS, please either submit an issue at https://github.com/ORNL/iris/issues or email vetter@computer.org. 

docs/sphinx/Makefile

@@ -0,0 +1,20 @@
+# Minimal makefile for Sphinx documentation
+#
+
+# You can set these variables from the command line, and also
+# from the environment for the first two.
+SPHINXOPTS    ?=
+SPHINXBUILD   ?= sphinx-build
+SOURCEDIR     = source
+BUILDDIR      = build
+
+# Put it first so that "make" without argument is like "make help".
+help:
+	@$(SPHINXBUILD) -M help "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)
+
+.PHONY: help Makefile
+
+# Catch-all target: route all unknown targets to Sphinx using the new
+# "make mode" option.  $(O) is meant as a shortcut for $(SPHINXOPTS).
+%: Makefile
+	@$(SPHINXBUILD) -M $@ "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O)

docs/sphinx/environment.yml

@@ -0,0 +1,14 @@
+name: IRIS-docs
+
+channels:
+  - conda-forge
+  - defaults
+
+dependencies:
+  - python
+  - breathe
+  - pip
+  - pip:
+    - sphinxcontrib-contentui
+    - sphinx_rtd_theme
+

docs/sphinx/make.bat

@@ -0,0 +1,35 @@
+@ECHO OFF
+
+pushd %~dp0
+
+REM Command file for Sphinx documentation
+
+if "%SPHINXBUILD%" == "" (
+	set SPHINXBUILD=sphinx-build
+)
+set SOURCEDIR=source
+set BUILDDIR=build
+
+if "%1" == "" goto help
+
+%SPHINXBUILD% >NUL 2>NUL
+if errorlevel 9009 (
+	echo.
+	echo.The 'sphinx-build' command was not found. Make sure you have Sphinx
+	echo.installed, then set the SPHINXBUILD environment variable to point
+	echo.to the full path of the 'sphinx-build' executable. Alternatively you
+	echo.may add the Sphinx directory to PATH.
+	echo.
+	echo.If you don't have Sphinx installed, grab it from
+	echo.http://sphinx-doc.org/
+	exit /b 1
+)
+
+%SPHINXBUILD% -M %1 %SOURCEDIR% %BUILDDIR% %SPHINXOPTS% %O%
+goto end
+
+:help
+%SPHINXBUILD% -M help %SOURCEDIR% %BUILDDIR% %SPHINXOPTS% %O%
+
+:end
+popd

docs/sphinx/source/_code/helloworld/helloworld.c

@@ -0,0 +1,34 @@
+#include <iris/iris.h>
+#include <stdio.h>
+
+char a[12] = "hello world";
+char b[12];
+size_t size = 12;
+
+int main(int argc, char** argv) {
+  iris_init(&argc, &argv, 1);
+
+  iris_mem mem_a;
+  iris_mem mem_b;
+  iris_mem_create(size, &mem_a);
+  iris_mem_create(size, &mem_b);
+
+  iris_task task;
+  iris_task_create(&task);
+  iris_task_h2d(task, mem_a, 0, size, a);
+  void* params[2] = { mem_b, mem_a };
+  int params_info[2] = { iris_w, iris_r };
+  iris_task_kernel(task, "uppercase", 1, NULL, &size, NULL, 2, params, params_info);
+  iris_task_d2h(task, mem_b, 0, size, b);
+  iris_task_submit(task, iris_roundrobin, NULL, 1);
+
+  printf("%s\n", b);
+
+  iris_task_release(task);
+  iris_mem_release(mem_a);
+  iris_mem_release(mem_b);
+
+  iris_finalize();
+
+  return 0;
+}

docs/sphinx/source/_code/helloworld/helloworld.cpp

@@ -0,0 +1,28 @@
+#include <iris/iris.hpp>
+#include <stdio.h>
+
+char a[12] = "hello world";
+char b[12];
+size_t size = 12;
+
+int main(int argc, char** argv) {
+  iris::Platform platform;
+  platform.init(&argc, &argv, true);
+
+  iris::Mem mem_a(size);
+  iris::Mem mem_b(size);
+
+  iris::Task task;
+  task.h2d(&mem_a, 0, size, a);
+  void* params[2] = { &mem_b, &mem_a };
+  int params_info[2] = { iris_w, iris_r };
+  task.kernel("uppercase", 1, NULL, &size, NULL, 2, params, params_info);
+  task.d2h(&mem_b, 0, size, b);
+  task.submit(iris_roundrobin, NULL, true);
+
+  printf("%s\n", b);
+
+  platform.finalize();
+
+  return 0;
+}

docs/sphinx/source/_code/helloworld/kernel.cl

@@ -0,0 +1,5 @@
+__kernel void uppercase(__global char* b, __global char* a) {
+  int i = get_global_id(0);
+  if (a[i] >= 'a' && a[i] <= 'z') b[i] = a[i] + 'A' - 'a';
+  else b[i] = a[i];
+}

docs/sphinx/source/_code/helloworld/kernel.cu

@@ -0,0 +1,5 @@
+extern "C" __global__ void uppercase(char* b, char* a) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (a[i] >= 'a' && a[i] <= 'z') b[i] = a[i] + 'A' - 'a';
+  else b[i] = a[i];
+}

docs/sphinx/source/_code/helloworld/kernel.hexagon.c

@@ -0,0 +1,10 @@
+#include <iris/iris_hexagon_imp.h>
+
+AEEResult irishexagon_uppercase(char* b, int blen, char* a, int alen, IRIS_HEXAGON_KERNEL_ARGS) {
+  int32 i = 0;
+  IRIS_HEXAGON_KERNEL_BEGIN(i)
+  if (a[i] >= 'a' && a[i] <= 'z') b[i] = a[i] + 'A' - 'a';
+  else b[i] = a[i];
+  IRIS_HEXAGON_KERNEL_END
+  return AEE_SUCCESS;
+}

docs/sphinx/source/_code/helloworld/kernel.hip.cpp

@@ -0,0 +1,7 @@
+#include <hip/hip_runtime.h>
+
+extern "C" __global__ void uppercase(char* b, char* a) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (a[i] >= 'a' && a[i] <= 'z') b[i] = a[i] + 'A' - 'a';
+  else b[i] = a[i];
+}

docs/sphinx/source/_code/helloworld/kernel.omp.h

@@ -0,0 +1,10 @@
+#include <iris/iris_openmp.h>
+
+static void uppercase(char* b, char* a, IRIS_OPENMP_KERNEL_ARGS) {
+  int i = 0;
+#pragma omp parallel for shared(b, a) private(i)
+  IRIS_OPENMP_KERNEL_BEGIN(i)
+  if (a[i] >= 'a' && a[i] <= 'z') b[i] = a[i] + 'A' - 'a';
+  else b[i] = a[i];
+  IRIS_OPENMP_KERNEL_END
+}

docs/sphinx/source/_code/saxpy/kernel.cl

@@ -0,0 +1,4 @@
+__kernel void saxpy(__global float* S, float A, __global float* X, __global float* Y) {
+  int i = get_global_id(0);
+  S[i] = A * X[i] + Y[i];
+}

docs/sphinx/source/_code/saxpy/kernel.cu

@@ -0,0 +1,4 @@
+extern "C" __global__ void saxpy(float* S, float A, float* X, float* Y) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  S[i] = A * X[i] + Y[i];
+}

docs/sphinx/source/_code/saxpy/kernel.hexagon.c

@@ -0,0 +1,9 @@
+#include <iris/iris_hexagon_imp.h>
+
+AEEResult irishexagon_saxpy(float* S, int Slen, float A, float* X, int Xlen, float* Y, int Ylen, IRIS_HEXAGON_KERNEL_ARGS) {
+  int32 i = 0;
+  IRIS_HEXAGON_KERNEL_BEGIN(i)
+  S[i] = A * X[i] + Y[i];
+  IRIS_HEXAGON_KERNEL_END
+  return AEE_SUCCESS;
+}

docs/sphinx/source/_code/saxpy/kernel.hip.cpp

@@ -0,0 +1,6 @@
+#include <hip/hip_runtime.h>
+
+extern "C" __global__ void saxpy(float* S, float A, float* X, float* Y) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  S[i] = A * X[i] + Y[i];
+}

docs/sphinx/source/_code/saxpy/kernel.omp.h

@@ -0,0 +1,9 @@
+#include <iris/iris_openmp.h>
+
+static void saxpy(float* S, float A, float* X, float* Y, IRIS_OPENMP_KERNEL_ARGS) {
+  int i = 0;
+#pragma omp parallel for shared(S, A, X, Y) private(i)
+  IRIS_OPENMP_KERNEL_BEGIN(i)
+  S[i] = A * X[i] + Y[i];
+  IRIS_OPENMP_KERNEL_END
+}

docs/sphinx/source/_code/saxpy/kernel.openmp.c

@@ -0,0 +1,69 @@
+#include <brisbane/brisbane_openmp.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef struct {
+  __global float * Z;
+  float A;
+  __global float * X;
+  __global float * Y;
+} brisbane_openmp_saxpy_args;
+brisbane_openmp_saxpy_args saxpy_args;
+
+static int brisbane_openmp_saxpy_setarg(int idx, size_t size, void* value) {
+  switch (idx) {
+    case 1: memcpy(&saxpy_args.A, value, size); break;
+    default: return BRISBANE_ERR;
+  }
+  return BRISBANE_OK;
+}
+
+static int brisbane_openmp_saxpy_setmem(int idx, void* mem) {
+  switch (idx) {
+    case 0: saxpy_args.Z = (__global float *__restrict) mem; break;
+    case 2: saxpy_args.X = (__global float *__restrict) mem; break;
+    case 3: saxpy_args.Y = (__global float *__restrict) mem; break;
+    default: return BRISBANE_ERR;
+  }
+  return BRISBANE_OK;
+}
+
+#include "kernel.openmp.h"
+
+int brisbane_openmp_kernel(const char* name) {
+  brisbane_openmp_lock();
+  if (strcmp(name, "saxpy") == 0) {
+    brisbane_openmp_kernel_idx = 0;
+    return BRISBANE_OK;
+  }
+  return BRISBANE_ERR;
+}
+
+int brisbane_openmp_setarg(int idx, size_t size, void* value) {
+  switch (brisbane_openmp_kernel_idx) {
+    case 0: return brisbane_openmp_saxpy_setarg(idx, size, value);
+  }
+  return BRISBANE_ERR;
+}
+
+int brisbane_openmp_setmem(int idx, void* mem) {
+  switch (brisbane_openmp_kernel_idx) {
+    case 0: return brisbane_openmp_saxpy_setmem(idx, mem);
+  }
+  return BRISBANE_ERR;
+}
+
+int brisbane_openmp_launch(int dim, size_t off, size_t ndr) {
+  switch (brisbane_openmp_kernel_idx) {
+    case 0: saxpy(saxpy_args.Z, saxpy_args.A, saxpy_args.X, saxpy_args.Y, off, ndr); break;
+  }
+  brisbane_openmp_unlock();
+  return BRISBANE_OK;
+}
+
+#ifdef __cplusplus
+} /* end of extern "C" */
+#endif
+

docs/sphinx/source/_code/saxpy/kernel.openmp.h

@@ -0,0 +1,10 @@
+#include <iris/iris_openmp.h>
+
+static void saxpy(float* Z, float A, float* X, float* Y, IRIS_OPENMP_KERNEL_ARGS) {
+  size_t i;
+#pragma omp parallel for shared(Z, A, X, Y) private(i)
+  IRIS_OPENMP_KERNEL_BEGIN(i)
+  Z[i] = A * X[i] + Y[i];
+  IRIS_OPENMP_KERNEL_END
+}
+

docs/sphinx/source/_code/saxpy/saxpy.c

@@ -0,0 +1,84 @@
+#include <iris/iris.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <malloc.h>
+
+int main(int argc, char** argv) {
+  iris_init(&argc, &argv, 1);
+
+  size_t SIZE;
+  int TARGET;
+  int VERBOSE;
+  float *X, *Y, *Z;
+  float A = 10;
+  int ERROR = 0;
+
+  SIZE = argc > 1 ? atol(argv[1]) : 8;
+  TARGET = argc > 2 ? atol(argv[2]) : 0;
+  VERBOSE = argc > 3 ? atol(argv[3]) : 1;
+
+  printf("[%s:%d] SIZE[%zu] TARGET[%d] VERBOSE[%d]\n", __FILE__, __LINE__, SIZE, TARGET, VERBOSE);
+
+  X = (float*) malloc(SIZE * sizeof(float));
+  Y = (float*) malloc(SIZE * sizeof(float));
+  Z = (float*) malloc(SIZE * sizeof(float));
+
+  if (VERBOSE) {
+
+  for (int i = 0; i < SIZE; i++) {
+    X[i] = i;
+    Y[i] = i;
+  }
+
+  printf("X [");
+  for (int i = 0; i < SIZE; i++) printf(" %2.0f.", X[i]);
+  printf("]\n");
+  printf("Y [");
+  for (int i = 0; i < SIZE; i++) printf(" %2.0f.", Y[i]);
+  printf("]\n");
+
+  }
+
+  iris_mem mem_X;
+  iris_mem mem_Y;
+  iris_mem mem_Z;
+  iris_mem_create(SIZE * sizeof(float), &mem_X);
+  iris_mem_create(SIZE * sizeof(float), &mem_Y);
+  iris_mem_create(SIZE * sizeof(float), &mem_Z);
+
+  iris_task task0;
+  iris_task_create(&task0);
+  iris_task_h2d_full(task0, mem_X, X);
+  iris_task_h2d_full(task0, mem_Y, Y);
+  void* saxpy_params[4] = { mem_Z, &A, mem_X, mem_Y };
+  int saxpy_params_info[4] = { iris_w, sizeof(A), iris_r, iris_r };
+  iris_task_kernel(task0, "saxpy", 1, NULL, &SIZE, NULL, 4, saxpy_params, saxpy_params_info);
+  iris_task_d2h_full(task0, mem_Z, Z);
+  iris_task_submit(task0, TARGET, NULL, 1);
+
+  if (VERBOSE) {
+
+  for (int i = 0; i < SIZE; i++) {
+    if (Z[i] != A * X[i] + Y[i]) ERROR++;
+  }
+
+  printf("S = %f * X + Y [", A);
+  for (int i = 0; i < SIZE; i++) printf(" %3.0f.", Z[i]);
+  printf("]\n");
+
+  }
+
+  iris_mem_release(mem_X);
+  iris_mem_release(mem_Y);
+  iris_mem_release(mem_Z);
+
+  free(X);
+  free(Y);
+  free(Z);
+
+  iris_task_release(task0);
+
+  iris_finalize();
+
+  return 0;
+}