Hello World

Setting the following environment variables is highly recommended to make life easier.

$ export IRIS=<install_path> # install_path would be $HOME/.local
$ export CPATH=$CPATH:$IRIS/include
$ export LIBRARY_PATH=$LIBRARY_PATH:$IRIS/lib:$IRIS/lib64
$ export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$IRIS/lib:$IRIS/lib64
$ export PYTHONPATH=$PYTHONPATH:$IRIS/include

The “Hello World” program is the first step towards learning IRIS. This program displays the message “HELLO WORLD” on the screen.

$ cd iris/apps/helloworld
$ make
$ ./helloworld
HELLO WORLD
$

Host Code

C

#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;
}

C++

#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;
}

Kernels

CUDA

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];
}

HIP

#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];
}

OpenCL

__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];
}

OpenMP

#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
}

Hexagon

#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;
}

SAXPY

SAXPY stands for “Single-precision A * X Plus Y”. It is a combination of scalar multiplication and vector addition.

$ cd iris/apps/saxpy
$ make
$ ./saxpy-c
X [  0.  1.  2.  3.  4.  5.  6.  7.]
Y [  0.  1.  2.  3.  4.  5.  6.  7.]
S = 10.000000 * X + Y [   0.  11.  22.  33.  44.  55.  66.  77.]
$