{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## __interop__ Construct"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "The __interop__ construct allows OpenMP to interoperate with foreign runtime environments. In the example below, asynchronous cuda memory copies and a  _cublasDaxpy_  routine are executed  in a cuda stream. Also, an asynchronous target task execution (having a __nowait__ clause)  and two explicit tasks are executed through OpenMP directives.  Scheduling dependences (synchronization) are imposed on the foreign stream and the OpenMP tasks through __depend__ clauses."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "First, an interop object,  _obj_ , is initialized for synchronization by including the __targetsync__  _interop-type_  in the interop __init__ clause  (__init(__ __targetsync,obj__ __)__).   The object provides access to the foreign runtime. The __depend__ clause provides a dependence behavior for foreign tasks associated with a valid object."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Next, the __omp_get_interop_int__ routine is used to extract the foreign  runtime id (__omp_ipr_fr_id__), and a test in the next statement ensures  that the cuda runtime (__omp_ifr_cuda__) is available."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Within the block for executing the  _cublasDaxpy_  routine, a stream is acquired  with the __omp_get_interop_ptr__ routine, which returns a cuda stream ( _s_ ). The stream is included in the cublas handle, and used directly in the asynchronous memory routines.  The following __interop__ construct, with the __destroy__ clause,  ensures that the foreign tasks have completed."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "//%compiler: clang\n",
    "//%cflags: -fopenmp\n",
    "\n",
    "/*\n",
    "* name:       interop.1\n",
    "* type:       C\n",
    "* version:    omp_5.1\n",
    "*/\n",
    "#include <omp.h>\n",
    "#include <stdio.h>\n",
    "#include <stdlib.h>\n",
    "#include <cublas_v2.h>\n",
    "#include <cuda_runtime_api.h>\n",
    "\n",
    "#define N 16384\n",
    "\n",
    "void myVectorSet(int n, double s, double *x)\n",
    "{\n",
    "   for(int i=0; i<n; ++i) x[i] = s*(i+1);\n",
    "}\n",
    "void myDaxpy(int n, double s, double *x, double *y)\n",
    "{\n",
    "   for(int i=0; i<n; ++i) y[i] = s*x[i]+y[i];\n",
    "}\n",
    "void myDscal(int n, double s, double *x)\n",
    "{\n",
    "   for(int i=0; i<n; ++i) x[i] = s*x[i];\n",
    "}\n",
    "\n",
    "\n",
    "int main(){\n",
    "  const  double scalar=2.0;\n",
    "  double *x, *y, *d_x, *d_y;\n",
    "  int    dev;\n",
    "\n",
    "  omp_interop_t obj=omp_interop_none;\n",
    "  intptr_t  type;\n",
    "\n",
    "  // Async Memcpy requires pinned memory\n",
    "  cudaMallocHost( (void**)&x,   N*sizeof(double) );\n",
    "  cudaMallocHost( (void**)&y,   N*sizeof(double) );\n",
    "  cudaMalloc(     (void**)&d_x, N*sizeof(double) );\n",
    "  cudaMalloc(     (void**)&d_y, N*sizeof(double) );\n",
    "\n",
    "  dev = omp_get_default_device();\n",
    "  omp_target_associate_ptr(&x[0], d_x, sizeof(double)*N, 0, dev);\n",
    "  omp_target_associate_ptr(&y[0], d_y, sizeof(double)*N, 0, dev);\n",
    "\n",
    "  #pragma omp target nowait depend(out: x[0:N]) \\\n",
    "                     map(from: x[0:N]) device(dev)\n",
    "  myVectorSet(N,  1.0, x);\n",
    "\n",
    "  #pragma omp task depend(out: y[0:N])\n",
    "  myVectorSet(N, -1.0, y);\n",
    "\n",
    "  // get obj for syncing\n",
    "  #pragma omp interop init(targetsync: obj) device(dev) \\\n",
    "                 depend(in: x[0:N]) depend(inout: y[0:N])\n",
    "\n",
    "                                        //foreign rt id and string name\n",
    "  int        id = (int  )omp_get_interop_int(obj, omp_ipr_fr_id,   NULL);\n",
    "  char* rt_name = (char*)omp_get_interop_str(obj, omp_ipr_fr_name, NULL);\n",
    "\n",
    "  if(obj != omp_interop_none && id == omp_ifr_cuda) {\n",
    "\n",
    "    printf(\" OpenMP working with %s runtime to execute cublas daxpy.\\n\",\n",
    "           rt_name);\n",
    "    cublasHandle_t handle;\n",
    "    int rc;\n",
    "    cublasCreate(&handle);\n",
    "\n",
    "    cudaStream_t s=\n",
    "        (cudaStream_t)omp_get_interop_ptr(obj, omp_ipr_targetsync, &rc);\n",
    "    if(rc != omp_irc_success) {\n",
    "       fprintf(stderr,\"ERROR: Failed to get %s stream, rt error= %d.\\n\",\n",
    "               rt_name, rc);\n",
    "       if(rc == omp_irc_no_value)\n",
    "          fprintf(stderr,\n",
    "                  \"Parameters valid, no meaningful value available.\");\n",
    "       exit(1);\n",
    "    }\n",
    "\n",
    "    cublasSetStream( handle,s );\n",
    "    cudaMemcpyAsync( d_x, x, N*sizeof(double),\n",
    "                     cudaMemcpyHostToDevice, s );\n",
    "    cudaMemcpyAsync( d_y, y, N*sizeof(double),\n",
    "                     cudaMemcpyHostToDevice, s );\n",
    "    cublasDaxpy(     handle, N, &scalar, &d_x[0], 1, &d_y[0], 1 ) ;\n",
    "    cudaMemcpyAsync( y, d_y, N*sizeof(double),\n",
    "                     cudaMemcpyDeviceToHost, s );\n",
    "\n",
    "  } else {     // Execute as OpenMP offload.\n",
    "\n",
    "    printf(\" Notice: Offloading myDaxpy to perform daxpy calculation.\\n\");\n",
    "\n",
    "    #pragma omp target depend(inout: y[0:N]) depend(in: x[0:N]) nowait \\\n",
    "                       map(to: x[0:N]) map(tofrom: y[0:N]) device(dev)\n",
    "    myDaxpy(N, scalar, x, y);\n",
    "\n",
    "  }\n",
    "\n",
    "   // This also ensures foreign tasks complete.\n",
    "  #pragma omp interop destroy(obj) nowait depend(out: y[0:N])\n",
    "\n",
    "  #pragma omp target depend(inout: x[0:N])\n",
    "  myDscal(N, scalar, x);\n",
    "\n",
    "  #pragma omp taskwait\n",
    "  printf(\"(-1:-16384) %f:%f\\n\", y[0], y[N-1]);\n",
    "  printf(\"(-2:-32768) %f:%f\\n\", x[0], x[N-1]);\n",
    "\n",
    "}"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Native",
   "language": "native",
   "name": "native"
  },
  "language_info": {
   "file_extension": ".c",
   "mimetype": "text/plain",
   "name": "c"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 4
}