diff --git a/libgomp/config.h.in b/libgomp/config.h.in
index 14c7e2a..67f5420 100644
--- a/libgomp/config.h.in
+++ b/libgomp/config.h.in
@@ -30,6 +30,9 @@
 /* Define to 1 if you have the <inttypes.h> header file. */
 #undef HAVE_INTTYPES_H
 
+/* Define to 1 if you have the `dl' library (-ldl). */
+#undef HAVE_LIBDL
+
 /* Define to 1 if you have the <memory.h> header file. */
 #undef HAVE_MEMORY_H
 
@@ -107,6 +110,9 @@
 /* Define to the version of this package. */
 #undef PACKAGE_VERSION
 
+/* Define if all infrastructure, needed for plugins, is supported. */
+#undef PLUGIN_SUPPORT
+
 /* The size of `char', as computed by sizeof. */
 #undef SIZEOF_CHAR
 
diff --git a/libgomp/configure b/libgomp/configure
index 238b1af..f4f71a4 100755
--- a/libgomp/configure
+++ b/libgomp/configure
@@ -15046,6 +15046,69 @@ fi
 rm -f core conftest.err conftest.$ac_objext \
     conftest$ac_exeext conftest.$ac_ext
 
+plugin_support=yes
+{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for dlsym in -ldl" >&5
+$as_echo_n "checking for dlsym in -ldl... " >&6; }
+if test "${ac_cv_lib_dl_dlsym+set}" = set; then :
+  $as_echo_n "(cached) " >&6
+else
+  ac_check_lib_save_LIBS=$LIBS
+LIBS="-ldl  $LIBS"
+cat confdefs.h - <<_ACEOF >conftest.$ac_ext
+/* end confdefs.h.  */
+
+/* Override any GCC internal prototype to avoid an error.
+   Use char because int might match the return type of a GCC
+   builtin and then its argument prototype would still apply.  */
+#ifdef __cplusplus
+extern "C"
+#endif
+char dlsym ();
+int
+main ()
+{
+return dlsym ();
+  ;
+  return 0;
+}
+_ACEOF
+if ac_fn_c_try_link "$LINENO"; then :
+  ac_cv_lib_dl_dlsym=yes
+else
+  ac_cv_lib_dl_dlsym=no
+fi
+rm -f core conftest.err conftest.$ac_objext \
+    conftest$ac_exeext conftest.$ac_ext
+LIBS=$ac_check_lib_save_LIBS
+fi
+{ $as_echo "$as_me:${as_lineno-$LINENO}: result: $ac_cv_lib_dl_dlsym" >&5
+$as_echo "$ac_cv_lib_dl_dlsym" >&6; }
+if test "x$ac_cv_lib_dl_dlsym" = x""yes; then :
+  cat >>confdefs.h <<_ACEOF
+#define HAVE_LIBDL 1
+_ACEOF
+
+  LIBS="-ldl $LIBS"
+
+else
+  plugin_support=no
+fi
+
+ac_fn_c_check_header_mongrel "$LINENO" "dirent.h" "ac_cv_header_dirent_h" "$ac_includes_default"
+if test "x$ac_cv_header_dirent_h" = x""yes; then :
+
+else
+  plugin_support=no
+fi
+
+
+
+if test x$plugin_support = xyes; then
+
+$as_echo "#define PLUGIN_SUPPORT 1" >>confdefs.h
+
+fi
+
 # Check for functions needed.
 for ac_func in getloadavg clock_gettime strtoull
 do :
diff --git a/libgomp/configure.ac b/libgomp/configure.ac
index d87ed29..85ecbcf 100644
--- a/libgomp/configure.ac
+++ b/libgomp/configure.ac
@@ -193,6 +193,15 @@ AC_LINK_IFELSE(
    [],
    [AC_MSG_ERROR([Pthreads are required to build libgomp])])])
 
+plugin_support=yes
+AC_CHECK_LIB(dl, dlsym, , [plugin_support=no])
+AC_CHECK_HEADER(dirent.h, , [plugin_support=no])
+
+if test x$plugin_support = xyes; then
+  AC_DEFINE(PLUGIN_SUPPORT, 1,
+    [Define if all infrastructure, needed for plugins, is supported.])
+fi
+
 # Check for functions needed.
 AC_CHECK_FUNCS(getloadavg clock_gettime strtoull)
 
diff --git a/libgomp/target.c b/libgomp/target.c
index 8b445bc..e1213b1 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -30,6 +30,15 @@
 #include <stdlib.h>
 #include <string.h>
 
+#ifdef PLUGIN_SUPPORT
+# include <dlfcn.h>
+# include <dirent.h>
+#endif
+
+static void gomp_target_init (void);
+
+static pthread_once_t gomp_is_initialized = PTHREAD_ONCE_INIT;
+
 /* Forward declaration for a node in the tree.  */
 typedef struct splay_tree_node_s *splay_tree_node;
 typedef struct splay_tree_s *splay_tree;
@@ -50,6 +59,10 @@ struct target_mem_desc {
   struct target_mem_desc *prev;
   /* Number of items in following list.  */
   size_t list_count;
+
+  /* Corresponding target device descriptor.  */
+  struct gomp_device_descr *device_descr;
+
   /* List of splay keys to remove (or decrease refcount)
      at the end of region.  */
   splay_tree_key list[];
@@ -70,6 +83,12 @@ struct splay_tree_key_s {
   bool copy_from;
 };
 
+/* Array of descriptors of all available devices.  */
+static struct gomp_device_descr *devices;
+
+/* Total number of available devices.  */
+static int num_devices;
+
 /* The comparison function.  */
 
 static int
@@ -87,33 +106,55 @@ splay_compare (splay_tree_key x, splay_tree_key y)
 
 #include "splay-tree.h"
 
+/* This structure describes accelerator device.
+   It contains name of the corresponding libgomp plugin, function handlers for
+   interaction with the device, ID-number of the device, and information about
+   mapped memory.  */
+struct gomp_device_descr
+{
+  /* This is the ID number of device.  It could be specified in DEVICE-clause of
+     TARGET construct.  */
+  int id;
+
+  /* Plugin file handler.  */
+  void *plugin_handle;
+
+  /* Function handlers.  */
+  bool (*device_available_func) (void);
+
+  /* Splay tree containing information about mapped memory regions.  */
+  struct splay_tree_s dev_splay_tree;
+
+  /* Mutex for operating with the splay tree and other shared structures.  */
+  gomp_mutex_t dev_env_lock;
+};
+
 attribute_hidden int
 gomp_get_num_devices (void)
 {
-  /* FIXME: Scan supported accelerators when called the first time.  */
-  return 0;
+  (void) pthread_once (&gomp_is_initialized, gomp_target_init);
+  return num_devices;
 }
 
-static int
-resolve_device (int device)
+static struct gomp_device_descr *
+resolve_device (int device_id)
 {
-  if (device == -1)
+  if (device_id == -1)
     {
       struct gomp_task_icv *icv = gomp_icv (false);
-      device = icv->default_device_var;
+      device_id = icv->default_device_var;
     }
+  if (device_id >= gomp_get_num_devices ()
+      && device_id != 257)
+    return NULL;
+
   /* FIXME: Temporary hack for testing non-shared address spaces on host.  */
-  if (device == 257)
-    return 257;
-  if (device >= gomp_get_num_devices ())
-    return -1;
-  return -1;
+  if (device_id == 257)
+    return &devices[0];
+
+  return &devices[device_id];
 }
 
-/* These variables would be per-accelerator (which doesn't have shared address
-   space.  */
-static struct splay_tree_s dev_splay_tree;
-static gomp_mutex_t dev_env_lock;
 
 /* Handle the case where splay_tree_lookup found oldn for newn.
    Helper function of gomp_map_vars.  */
@@ -137,8 +178,9 @@ gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn,
 }
 
 static struct target_mem_desc *
-gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
-	       unsigned char *kinds, bool is_target)
+gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
+	       void **hostaddrs, size_t *sizes, unsigned char *kinds,
+	       bool is_target)
 {
   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
   struct splay_tree_key_s cur_node;
@@ -146,6 +188,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
   tgt->list_count = mapnum;
   tgt->refcount = 1;
+  tgt->device_descr = devicep;
 
   if (mapnum == 0)
     return tgt;
@@ -159,7 +202,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
       tgt_size = mapnum * sizeof (void *);
     }
 
-  gomp_mutex_lock (&dev_env_lock);
+  gomp_mutex_lock (&devicep->dev_env_lock);
   for (i = 0; i < mapnum; i++)
     {
       cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -167,7 +210,8 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
 	cur_node.host_end = cur_node.host_start + sizes[i];
       else
 	cur_node.host_end = cur_node.host_start + sizeof (void *);
-      splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node);
+      splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree,
+					    &cur_node);
       if (n)
 	{
 	  tgt->list[i] = n;
@@ -215,7 +259,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
 	    else
 	      k->host_end = k->host_start + sizeof (void *);
 	    splay_tree_key n
-	      = splay_tree_lookup (&dev_splay_tree, k);
+	      = splay_tree_lookup (&devicep->dev_splay_tree, k);
 	    if (n)
 	      {
 		tgt->list[i] = n;
@@ -235,7 +279,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
 		tgt->refcount++;
 		array->left = NULL;
 		array->right = NULL;
-		splay_tree_insert (&dev_splay_tree, array);
+		splay_tree_insert (&devicep->dev_splay_tree, array);
 		switch (kinds[i] & 7)
 		  {
 		  case 0: /* ALLOC */
@@ -257,16 +301,19 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
 		    /* Add bias to the pointer value.  */
 		    cur_node.host_start += sizes[i];
 		    cur_node.host_end = cur_node.host_start + 1;
-		    n = splay_tree_lookup (&dev_splay_tree, &cur_node);
+		    n = splay_tree_lookup (&devicep->dev_splay_tree,
+					   &cur_node);
 		    if (n == NULL)
 		      {
 			/* Could be possibly zero size array section.  */
 			cur_node.host_end--;
-			n = splay_tree_lookup (&dev_splay_tree, &cur_node);
+			n = splay_tree_lookup (&devicep->dev_splay_tree,
+					       &cur_node);
 			if (n == NULL)
 			  {
 			    cur_node.host_start--;
-			    n = splay_tree_lookup (&dev_splay_tree, &cur_node);
+			    n = splay_tree_lookup (&devicep->dev_splay_tree,
+						   &cur_node);
 			    cur_node.host_start++;
 			  }
 		      }
@@ -303,7 +350,7 @@ gomp_map_vars (size_t mapnum, void **hostaddrs, size_t *sizes,
 	}
     }
 
-  gomp_mutex_unlock (&dev_env_lock);
+  gomp_mutex_unlock (&devicep->dev_env_lock);
   return tgt;
 }
 
@@ -322,6 +369,8 @@ gomp_unmap_tgt (struct target_mem_desc *tgt)
 static void
 gomp_unmap_vars (struct target_mem_desc *tgt)
 {
+  struct gomp_device_descr *devicep = tgt->device_descr;
+
   if (tgt->list_count == 0)
     {
       free (tgt);
@@ -329,7 +378,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt)
     }
 
   size_t i;
-  gomp_mutex_lock (&dev_env_lock);
+  gomp_mutex_lock (&devicep->dev_env_lock);
   for (i = 0; i < tgt->list_count; i++)
     if (tgt->list[i]->refcount > 1)
       tgt->list[i]->refcount--;
@@ -341,7 +390,7 @@ gomp_unmap_vars (struct target_mem_desc *tgt)
 	  memcpy ((void *) k->host_start,
 		  (void *) (k->tgt->tgt_start + k->tgt_offset),
 		  k->host_end - k->host_start);
-	splay_tree_remove (&dev_splay_tree, k);
+	splay_tree_remove (&devicep->dev_splay_tree, k);
 	if (k->tgt->refcount > 1)
 	  k->tgt->refcount--;
 	else
@@ -352,26 +401,30 @@ gomp_unmap_vars (struct target_mem_desc *tgt)
     tgt->refcount--;
   else
     gomp_unmap_tgt (tgt);
-  gomp_mutex_unlock (&dev_env_lock);
+  gomp_mutex_unlock (&devicep->dev_env_lock);
 }
 
 static void
-gomp_update (size_t mapnum, void **hostaddrs, size_t *sizes,
-	     unsigned char *kinds)
+gomp_update (struct gomp_device_descr *devicep, size_t mapnum,
+	     void **hostaddrs, size_t *sizes, unsigned char *kinds)
 {
   size_t i;
   struct splay_tree_key_s cur_node;
 
+  if (!devicep)
+    return;
+
   if (mapnum == 0)
     return;
 
-  gomp_mutex_lock (&dev_env_lock);
+  gomp_mutex_lock (&devicep->dev_env_lock);
   for (i = 0; i < mapnum; i++)
     if (sizes[i])
       {
 	cur_node.host_start = (uintptr_t) hostaddrs[i];
 	cur_node.host_end = cur_node.host_start + sizes[i];
-	splay_tree_key n = splay_tree_lookup (&dev_splay_tree, &cur_node);
+	splay_tree_key n = splay_tree_lookup (&devicep->dev_splay_tree,
+					      &cur_node);
 	if (n)
 	  {
 	    if (n->host_start > cur_node.host_start
@@ -400,7 +453,7 @@ gomp_update (size_t mapnum, void **hostaddrs, size_t *sizes,
 		      (void *) cur_node.host_start,
 		      (void *) cur_node.host_end);
       }
-  gomp_mutex_unlock (&dev_env_lock);
+  gomp_mutex_unlock (&devicep->dev_env_lock);
 }
 
 /* Called when encountering a target directive.  If DEVICE
@@ -418,28 +471,26 @@ GOMP_target (int device, void (*fn) (void *), const void *openmp_target,
 	     size_t mapnum, void **hostaddrs, size_t *sizes,
 	     unsigned char *kinds)
 {
-  device = resolve_device (device);
-  if (device == -1)
+  struct gomp_device_descr *devicep = resolve_device (device);
+  if (devicep == NULL)
     {
       /* Host fallback.  */
       fn (hostaddrs);
       return;
     }
-  if (device == 257)
-    {
-      struct target_mem_desc *tgt
-	= gomp_map_vars (mapnum, hostaddrs, sizes, kinds, true);
-      fn ((void *) tgt->tgt_start);
-      gomp_unmap_vars (tgt);
-    }
+
+  struct target_mem_desc *tgt
+    = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, true);
+  fn ((void *) tgt->tgt_start);
+  gomp_unmap_vars (tgt);
 }
 
 void
 GOMP_target_data (int device, const void *openmp_target, size_t mapnum,
 		  void **hostaddrs, size_t *sizes, unsigned char *kinds)
 {
-  device = resolve_device (device);
-  if (device == -1)
+  struct gomp_device_descr *devicep = resolve_device (device);
+  if (devicep == NULL)
     {
       /* Host fallback.  */
       struct gomp_task_icv *icv = gomp_icv (false);
@@ -450,21 +501,18 @@ GOMP_target_data (int device, const void *openmp_target, size_t mapnum,
 	     new #pragma omp target data, otherwise GOMP_target_end_data
 	     would get out of sync.  */
 	  struct target_mem_desc *tgt
-	    = gomp_map_vars (0, NULL, NULL, NULL, false);
+	    = gomp_map_vars (NULL, 0, NULL, NULL, NULL, false);
 	  tgt->prev = icv->target_data;
 	  icv->target_data = tgt;
 	}
       return;
     }
 
-  if (device == 257)
-    {
-      struct target_mem_desc *tgt
-	= gomp_map_vars (mapnum, hostaddrs, sizes, kinds, false);
-      struct gomp_task_icv *icv = gomp_icv (true);
-      tgt->prev = icv->target_data;
-      icv->target_data = tgt;
-    }
+  struct target_mem_desc *tgt
+    = gomp_map_vars (devicep, mapnum, hostaddrs, sizes, kinds, false);
+  struct gomp_task_icv *icv = gomp_icv (true);
+  tgt->prev = icv->target_data;
+  icv->target_data = tgt;
 }
 
 void
@@ -483,15 +531,156 @@ void
 GOMP_target_update (int device, const void *openmp_target, size_t mapnum,
 		    void **hostaddrs, size_t *sizes, unsigned char *kinds)
 {
-  device = resolve_device (device);
-  if (device == -1)
+  struct gomp_device_descr *devicep = resolve_device (device);
+  if (devicep == NULL)
     return;
 
-  if (device == 257)
-    gomp_update (mapnum, hostaddrs, sizes, kinds);
+  gomp_update (devicep, mapnum, hostaddrs, sizes, kinds);
 }
 
 void
 GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
 {
 }
+
+#ifdef PLUGIN_SUPPORT
+
+/* This function checks if the given string FNAME matches
+   "libgomp-plugin-*.so.1".  */
+static bool
+gomp_check_plugin_file_name (const char *fname)
+{
+  const char *prefix = "libgomp-plugin-";
+  const char *suffix = ".so.1";
+  if (!fname)
+    return false;
+  if (strncmp (fname, prefix, strlen (prefix)) != 0)
+    return false;
+  if (strncmp (fname + strlen (fname) - strlen (suffix), suffix,
+	       strlen (suffix)) != 0)
+    return false;
+  return true;
+}
+
+/* This function tries to load plugin for DEVICE.  Name of plugin is passed
+   in PLUGIN_NAME.
+   Plugin handle and handles of the found functions are stored in the
+   corresponding fields of DEVICE.
+   The function returns TRUE on success and FALSE otherwise.  */
+static bool
+gomp_load_plugin_for_device (struct gomp_device_descr *device,
+			     const char *plugin_name)
+{
+  if (!device || !plugin_name)
+    return false;
+
+  device->plugin_handle = dlopen (plugin_name, RTLD_LAZY);
+  if (!device->plugin_handle)
+    return false;
+
+  /* Clear any existing error.  */
+  dlerror ();
+
+  /* Check if all required functions are available in the plugin and store
+     their handlers.
+     TODO: check for other routines as well.  */
+  device->device_available_func = dlsym (device->plugin_handle,
+					 "device_available");
+  if (dlerror () != NULL)
+    {
+      dlclose (device->plugin_handle);
+      return false;
+    }
+
+  return true;
+}
+
+/* This functions scans folder, specified in environment variable
+   LIBGOMP_PLUGIN_PATH, and loads all suitable libgomp plugins from this folder.
+   For a plugin to be suitable, its name should be "libgomp-plugin-*.so.1" and
+   it should implement a certain set of functions.
+   Result of this function is properly initialized variable NUM_DEVICES and
+   array DEVICES, containing all plugins and their callback handles.  */
+static void
+gomp_find_available_plugins (void)
+{
+  char *plugin_path = NULL;
+  DIR *dir = NULL;
+  struct dirent *ent;
+  char plugin_name[PATH_MAX];
+
+  num_devices = 0;
+  devices = NULL;
+
+  plugin_path = getenv ("LIBGOMP_PLUGIN_PATH");
+  if (!plugin_path)
+    return;
+
+  dir = opendir (plugin_path);
+  if (!dir)
+    return;
+
+  while ((ent = readdir (dir)) != NULL)
+    {
+      struct gomp_device_descr current_device;
+      if (!gomp_check_plugin_file_name (ent->d_name))
+	continue;
+      if (strlen (plugin_path) + 1 + strlen (ent->d_name) >= PATH_MAX)
+	continue;
+      strcpy (plugin_name, plugin_path);
+      strcat (plugin_name, "/");
+      strcat (plugin_name, ent->d_name);
+      if (!gomp_load_plugin_for_device (&current_device, plugin_name))
+	continue;
+      devices = realloc (devices, (num_devices + 1)
+				  * sizeof (struct gomp_device_descr));
+      if (devices == NULL)
+	{
+	  num_devices = 0;
+	  closedir (dir);
+	  return;
+	}
+
+      devices[num_devices] = current_device;
+      devices[num_devices].id = num_devices + 1;
+      devices[num_devices].dev_splay_tree.root = NULL;
+      gomp_mutex_init (&devices[num_devices].dev_env_lock);
+      num_devices++;
+    }
+  closedir (dir);
+
+  /* FIXME: Temporary hack for testing non-shared address spaces on host.
+     We create device 257 just to check memory mapping.  */
+  if (num_devices == 0)
+    {
+      num_devices = 1;
+      devices = malloc (sizeof (struct gomp_device_descr));
+      if (devices == NULL)
+	{
+	  num_devices = 0;
+	  return;
+	}
+      devices[0].plugin_handle = NULL;
+      devices[0].device_available_func = NULL;
+      devices[0].dev_splay_tree.root = NULL;
+      gomp_mutex_init (&devices[0].dev_env_lock);
+    }
+  devices[0].id = 257;
+}
+
+/* This function initializes runtime needed for offloading.
+   It loads plugins, sets up a connection with devices, etc.  */
+static void
+gomp_target_init (void)
+{
+  gomp_find_available_plugins ();
+}
+
+#else /* PLUGIN_SUPPORT */
+/* If dlfcn.h is unavailable we always fallback to host execution.
+   GOMP_target* routines are just stubs for this case.  */
+static void
+gomp_target_init (void)
+{
+}
+#endif /* PLUGIN_SUPPORT */
