From 699ba6ca46b74eb17e34ccda510901d140d54b63 Mon Sep 17 00:00:00 2001
From: Thomas Roehl <thomas.roehl@fau.de>
Date: Sat, 9 Nov 2019 23:32:03 +0100
Subject: [PATCH] Update Nvidia GPU backend to working state

---
 Makefile                             |    2 +
 config.mk                            |   14 +-
 make/config_checks.mk                |    4 +-
 src/applications/likwid-perfctr.lua  |  214 +++--
 src/applications/likwid-topology.lua |   45 +-
 src/applications/likwid.lua          |    1 +
 src/cpustring.c                      |    4 +
 src/includes/likwid-gpumarker.h      |    2 +-
 src/includes/likwid.h                |    9 +-
 src/includes/nvmon_cupti.h           | 1000 +++++++++++++++++++++
 src/includes/nvmon_nvml.h            |   50 ++
 src/includes/nvmon_perfworks.h       |  274 ++++++
 src/includes/nvmon_types.h           |   24 +
 src/luawid.c                         |   96 +-
 src/nvmon.c                          | 1208 ++++----------------------
 src/topology_gpu.c                   |   25 +-
 test/Makefile                        |    9 +-
 test/cupti/matmul.cu                 |  370 ++++++++
 test/cupti/matmul.h                  |   17 +
 test/cupti/matmul_marker.c           |   27 +
 test/cupti/realtime.h                |   74 ++
 21 files changed, 2297 insertions(+), 1172 deletions(-)
 create mode 100644 src/includes/nvmon_cupti.h
 create mode 100644 src/includes/nvmon_nvml.h
 create mode 100644 src/includes/nvmon_perfworks.h
 create mode 100644 test/cupti/matmul.cu
 create mode 100644 test/cupti/matmul.h
 create mode 100644 test/cupti/matmul_marker.c
 create mode 100644 test/cupti/realtime.h

diff --git a/Makefile b/Makefile
index f06293d79..6662b44f8 100644
--- a/Makefile
+++ b/Makefile
@@ -91,6 +91,8 @@ OBJ := $(filter-out $(BUILD_DIR)/loadDataARM.o,$(OBJ))
 endif
 ifneq ($(NVIDIA_INTERFACE), true)
 OBJ := $(filter-out $(BUILD_DIR)/nvmon.o,$(OBJ))
+OBJ := $(filter-out $(BUILD_DIR)/topology_gpu.o,$(OBJ))
+OBJ := $(filter-out $(BUILD_DIR)/libnvctr.o,$(OBJ))
 endif
 PERFMONHEADERS  = $(patsubst $(SRC_DIR)/includes/%.txt, $(BUILD_DIR)/%.h,$(wildcard $(SRC_DIR)/includes/*.txt))
 OBJ_LUA    =  $(wildcard ./ext/lua/$(COMPILER)/*.o)
diff --git a/config.mk b/config.mk
index b25dfd847..82d2969a9 100644
--- a/config.mk
+++ b/config.mk
@@ -10,7 +10,7 @@
 COMPILER = GCC#NO SPACE
 
 # Path were to install likwid
-PREFIX ?= /usr/local#NO SPACE
+PREFIX ?= /apps/likwid/4.3.4#NO SPACE
 
 # Set the default mode for MSR access.
 # This can usually be overriden on the commandline.
@@ -26,6 +26,7 @@ FORTRAN_INTERFACE = false#NO SPACE
 INSTRUMENT_BENCH = true#NO SPACE
 
 # Build LIKWID with NVIDIA interface (CUDA, CUPTI)
+# For configuring include paths, go to CUDA section
 NVIDIA_INTERFACE = true#NO SPACE
 
 #################################################################
@@ -85,7 +86,7 @@ INSTALL_CHOWN = -g root -o root#NO SPACE
 SHARED_LIBRARY = true#NO SPACE
 
 # Build LIKWID with debug flags
-DEBUG = false#NO SPACE
+DEBUG = true#NO SPACE
 
 # Basic configuration for some internal arrays.
 # Maximal number of hardware threads
@@ -140,12 +141,11 @@ LIKWIDFILTERPATH = $(abspath $(INSTALLED_PREFIX)/share/likwid/filter)
 # $HOME/.likwid/groups
 LIKWIDGROUPPATH = $(abspath $(INSTALLED_PREFIX)/share/likwid/perfgroups)
 
-# CUDA info
+# CUDA / CUPTI build data
+# LIKWID requires CUDA and CUPTI to be present only for compilation with
+# NVIDIA_INTERFACE=true. At runtime, the CUDA and the CUPTI library have
+# to be in the LD_LIBRARY_PATH to dynamically load the libraries.
 # Include directory for CUDA headers
 CUDAINCLUDE = $(CUDA_HOME)/include
-# Library directory for CUDA libs
-CUDALIBDIR = $(CUDA_HOME)/lib64
 # Include directory for CUPTI headers
 CUPTIINCLUDE = $(CUDA_HOME)/extras/CUPTI/include
-# Library directory for CUPTI libs
-CUPTILIBDIR = $(CUDA_HOME)/extras/CUPTI/lib64
diff --git a/make/config_checks.mk b/make/config_checks.mk
index d3be955af..d05d77b69 100644
--- a/make/config_checks.mk
+++ b/make/config_checks.mk
@@ -77,7 +77,7 @@ endif
 endif
 
 ifeq ($(strip $(NVIDIA_INTERFACE)), true)
-LIBS+= -lcuda -ldl
+#LIBS+= -lcuda -ldl
 INCLUDES += -I$(CUDAINCLUDE) -I$(CUPTIINCLUDE)
-CPPFLAGS += -L$(CUDALIBDIR) -L$(CUPTILIBDIR)
+#CPPFLAGS += -L$(CUDALIBDIR) -L$(CUPTILIBDIR)
 endif
diff --git a/src/applications/likwid-perfctr.lua b/src/applications/likwid-perfctr.lua
index 6ca26908a..f5de7f097 100644
--- a/src/applications/likwid-perfctr.lua
+++ b/src/applications/likwid-perfctr.lua
@@ -49,13 +49,19 @@ local function examples()
     io.stdout:write("likwid-perfctr -e\n")
     io.stdout:write("List all events and suitable counters for events with 'L2' in them:\n")
     io.stdout:write("likwid-perfctr -E L2\n")
-    io.stdout:write("Run command on CPU 2 and measure performance group TEST:\n")
-    io.stdout:write("likwid-perfctr -C 2 -g TEST ./a.out\n")
+    io.stdout:write("Run command on CPU 2 and measure performance group CLOCK:\n")
+    io.stdout:write("likwid-perfctr -C 2 -g CLOCK ./a.out\n")
+    if likwid.gpuSupported() then
+        io.stdout:write("Run command and measure on GPU 1 the performance group FLOPS_DP (Only with NVMarkerAPI):\n")
+        io.stdout:write("likwid-perfctr -G 1 -W FLOPS_DP -m ./a.out\n")
+        io.stdout:write("It is possible to combine CPU and GPU measurements (with MarkerAPI and NVMarkerAPI):\n)
+        io.stdout:write("likwid-perfctr -C 2 -g CLOCK -G 1 -W FLOPS_DP -m ./a.out\n")
+    end
 end
 
 local function usage()
     version()
-    io.stdout:write("A tool to read out performance counter registers on x86 processors\n\n")
+    io.stdout:write("A tool to read out performance counter registers on x86, ARM and POWER processors\n\n")
     io.stdout:write("Options:\n")
     io.stdout:write("-h, --help\t\t Help message\n")
     io.stdout:write("-v, --version\t\t Version information\n")
@@ -63,7 +69,13 @@ local function usage()
     io.stdout:write("-c <list>\t\t Processor ids to measure (required), e.g. 1,2-4,8\n")
     io.stdout:write("-C <list>\t\t Processor ids to pin threads and measure, e.g. 1,2-4,8\n")
     io.stdout:write("\t\t\t For information about the <list> syntax, see likwid-pin\n")
-    io.stdout:write("-g, --group <string>\t Performance group or custom event set string\n")
+    if likwid.gpuSupported() then
+        io.stdout:write("-G, --gpus <list>\t List of GPUs to monitor\n")
+    end
+    io.stdout:write("-g, --group <string>\t Performance group or custom event set string for CPU monitoring\n")
+    if likwid.gpuSupported() then
+        io.stdout:write("-W, --gpugroup <string>\t Performance group or custom event set string for GPU monitoring\n")
+    end
     io.stdout:write("-H\t\t\t Get group help (together with -g switch)\n")
     io.stdout:write("-s, --skip <hex>\t Bitmask with threads to skip\n")
     io.stdout:write("-M <0|1>\t\t Set how MSR registers are accessed, 0=direct, 1=accessDaemon\n")
@@ -73,8 +85,6 @@ local function usage()
     io.stdout:write("-i, --info\t\t Print CPU info\n")
     io.stdout:write("-T <time>\t\t Switch eventsets with given frequency\n")
     io.stdout:write("-f, --force\t\t Force overwrite of registers if they are in use\n")
-    io.stdout:write("-G, --gpus <list>\t List of GPUs to monitor\n")
-    io.stdout:write("-w, --gpugroup <string>\t Performance group or custom event set string for GPU monitoring\n")
     io.stdout:write("Modes:")
     io.stdout:write("-S <time>\t\t Stethoscope mode with duration in s, ms or us, e.g 20ms\n")
     io.stdout:write("-t <time>\t\t Timeline mode with frequency in s, ms or us, e.g. 300ms\n")
@@ -137,7 +147,8 @@ outfile = nil
 outfile_orig = nil
 forceOverwrite = 0
 gotC = false
-markerFile = string.format("/tmp/likwid_%d.txt",likwid.getpid())
+markerFolder = "/tmp"
+markerFile = string.format("%s/likwid_%d.txt", markerFolder, likwid.getpid())
 cpuClock = 1
 execpid = false
 if config["daemonMode"] == -1 then
@@ -149,10 +160,11 @@ nan2value = '-'
 
 
 ---------------------------
+gpusSupported = likwid.gpuSupported()
 num_gpus = 0
 gpulist = {}
 gpu_event_string_list = {}
-gpuMarkerFile = string.format("/tmp/likwid_gpu_%d.txt",likwid.getpid())
+gpuMarkerFile = string.format("%s/likwid_gpu_%d.txt", markerFolder, likwid.getpid())
 gotG = false
 gpugroups = {}
 ---------------------------
@@ -331,7 +343,7 @@ for opt,arg in likwid.getopt(arg, {"a", "c:", "C:", "e", "E:", "g:", "h", "H", "
     elseif (opt == "stats") then
         print_stats = true
 ---------------------------
-    elseif (opt == "G") then
+    elseif gpusSupported and (opt == "G") then
         if arg ~= nil then
             num_gpus, gpulist = likwid.gpustr_to_gpulist(arg)
         else
@@ -342,7 +354,7 @@ for opt,arg in likwid.getopt(arg, {"a", "c:", "C:", "e", "E:", "g:", "h", "H", "
             os.exit(1)
         end
         gotG = true
-    elseif opt == "W" or opt == "gpugroup" then
+    elseif gpusSupported and (opt == "W" or opt == "gpugroup") then
         if arg ~= nil then
             table.insert(gpu_event_string_list, arg)
         else
@@ -376,7 +388,10 @@ io.stdout:setvbuf("no")
 cpuinfo = likwid.getCpuInfo()
 cputopo = likwid.getCpuTopology()
 ---------------------------
-gputopo = likwid.getGpuTopology()
+gputopo = nil
+if gpusSupported then
+    gputopo = likwid.getGpuTopology()
+end
 ---------------------------
 
 if num_cpus == 0 and
@@ -409,9 +424,10 @@ elseif num_cpus == 0 and
 end
 
 ---------------------------
-if num_gpus == 0 and
+if gpusSupported and
+   num_gpus == 0 and
    not gotG and
-   gpuinfo and
+   gputopo and
    not print_events and
    print_event == nil and
    not print_groups and
@@ -441,7 +457,7 @@ if num_cpus > 0 then
 end
 
 ---------------------------
-if gpuinfo and num_gpus > 0 then
+if gpuSupported and gputopo and num_gpus > 0 then
     for i,gpu1 in pairs(gpulist) do
         for j, gpu2 in pairs(gpulist) do
             if i ~= j and gpu1 == gpu2 then
@@ -481,9 +497,15 @@ if print_events == true then
         print_stdout(outstr)
     end
 ---------------------------
-    if gputopo then
+    if gpuSupported and gputopo then
+        local cudahome = os.getenv("CUDA_HOME")
+        if cudahome and cudahome:len() > 0 then
+            ldpath = os.getenv("LD_LIBRARY_PATH")
+            local cuptilib = string.format("%s/extras/CUPTI/lib64", cudahome)
+            likwid.setenv("LD_LIBRARY_PATH", cuptilib..":"..ldpath)
+        end
         tab = likwid.getGpuEventsAndCounters()
-        for d=0,tab["numDevices"]-1,1 do
+        for d=0,tab["numDevices"],1 do
             if tab["devices"][d] then
                 print_stdout("\n\n")
                 print_stdout(string.format("The GPUs %d provides %d events.", d, #tab["devices"][d]))
@@ -528,21 +550,29 @@ if print_event ~= nil then
         end
     end
 ---------------------------
-    if gputopo then
-        tab = likwid.getGpuEventsAndCounters()
-        for d=0,tab["numDevices"]-1,1 do
-            for _,e in pairs(tab["devices"][d]) do
-                if e["Name"]:match(case_insensitive_pattern(print_event)) then
-                    local f = false
-                    for _,x in pairs(events) do
-                        if e["Name"] == x["Name"] then
-                            f = true
-                            break
+    if gpuSupported and gputopo then
+        local cudahome = os.getenv("CUDA_HOME")
+        if cudahome and cudahome:len() > 0 then
+            ldpath = os.getenv("LD_LIBRARY_PATH")
+            local cuptilib = string.format("%s/extras/CUPTI/lib64", cudahome)
+            likwid.setenv("LD_LIBRARY_PATH", cuptilib..":"..ldpath)
+        end
+        if cudahome then
+            tab = likwid.getGpuEventsAndCounters()
+            for d=0,tab["numDevices"]-1,1 do
+                for _,e in pairs(tab["devices"][d]) do
+                    if e["Name"]:match(case_insensitive_pattern(print_event)) then
+                        local f = false
+                        for _,x in pairs(events) do
+                            if e["Name"] == x["Name"] then
+                                f = true
+                                break
+                            end
+                        end
+                        if not f then
+                            table.insert(events, e)
+                            counters["GPU"] = {["Name"] = "GPU", ["TypeName"] = "Nvidia GPU counters"}
                         end
-                    end
-                    if not f then
-                        table.insert(events, e)
-                        counters["GPU"] = {["Name"] = "GPU", ["TypeName"] = "Nvidia GPU counters"}
                     end
                 end
             end
@@ -590,8 +620,7 @@ if print_groups == true then
     else
         print_stdout(string.format("No groups defined for %s",cpuinfo["name"]))
     end
-    print(gputopo)
-    if gputopo then
+    if gpuSupported and gputopo then
         avail_groups = likwid.getGpuGroups()
         if avail_groups then
             local max_len = 0
@@ -671,6 +700,13 @@ if print_info or verbose > 0 then
         print_stdout(string.format("PERFMON width of counters:\t%u",cpuinfo["perf_width_ctr"]))
         print_stdout(string.format("PERFMON number of fixed counters:\t%u",cpuinfo["perf_num_fixed_ctr"]))
     end
+    if gpuSupported and gputopo then
+        print_stdout(likwid.hline)
+        for i=1, gputopo["numDevices"] do
+            gpu = gputopo["devices"][i]
+            print_stdout(string.format("NVMON GPU %d Compute capability:\t%.%d", gpu["id"], gpu["ccapMajor"], gpu["ccapMinor"]))
+        end
+    end
     print_stdout(likwid.hline)
     if print_info then
         likwid.printSupportedCPUs()
@@ -712,12 +748,12 @@ end
 if use_marker then
     if likwid.access(markerFile, "rw") ~= -1 then
         print_stderr(string.format("ERROR: MarkerAPI file %s not accessible. Maybe a remaining file of another user.", markerFile))
-        print_stderr("Please purge all MarkerAPI files from /tmp.")
+        print_stderr(string.format("Please purge all MarkerAPI files from %s.", markerFolder))
         os.exit(1)
     end
-    if #gpulist and likwid.access(gpuMarkerFile, "rw") ~= -1 then
+    if gpuSupported and #gpulist and likwid.access(gpuMarkerFile, "rw") ~= -1 then
         print_stderr(string.format("ERROR: GPUMarkerAPI file %s not accessible. Maybe a remaining file of another user.", gpuMarkerFile))
-        print_stderr("Please purge all GPUMarkerAPI files from /tmp.")
+        print_stderr(string.format("Please purge all GPUMarkerAPI files from %s.", markerFolder))
         os.exit(1)
     end
     if not pin_cpus and #cpulist > 0 and #event_string_list > 0 then
@@ -790,7 +826,7 @@ if use_marker == true then
     likwid.setenv("LIKWID_THREADS", table.concat(cpulist,","))
     likwid.setenv("LIKWID_FORCE", "-1")
     likwid.setenv("KMP_INIT_AT_FORK", "FALSE")
-    if #gpulist > 0 and #gpu_event_string_list > 0 then
+    if gpusSupported and #gpulist > 0 and #gpu_event_string_list > 0 then
         likwid.setenv("LIKWID_GPUS", table.concat(gpulist,","))
         str = table.concat(gpu_event_string_list, "|")
         likwid.setenv("LIKWID_GEVENTS", str)
@@ -826,14 +862,14 @@ if #event_string_list > 0 then
         os.exit(1)
     end
 end
-
-if #gpu_event_string_list > 0 then
+---------------------------
+if gpuSupported and #gpu_event_string_list > 0 then
     if likwid.gpuInit(num_gpus, gpulist) < 0 then
         likwid.putGpuTopology()
         os.exit(1)
     end
 end
-
+---------------------------
 
 if verbose > 0 then
     print_stdout(string.format("Executing: %s",table.concat(execList," ")))
@@ -845,7 +881,18 @@ if ldpath == nil then
 elseif not ldpath:match(libpath) then
     likwid.setenv("LD_LIBRARY_PATH", libpath..":"..ldpath)
 end
-
+---------------------------
+if gpusSupported then
+    local cudahome = os.getenv("CUDA_HOME")
+    if cudahome then
+        ldpath = os.getenv("LD_LIBRARY_PATH")
+        local cuptilib = string.format("%s/extras/CUPTI/lib64", cudahome)
+        if not ldpath:match(cuptilib) then
+            likwid.setenv("LD_LIBRARY_PATH", cuptilib..":"..ldpath)
+        end
+    end
+end
+---------------------------
 
 
 local pid = nil
@@ -895,16 +942,18 @@ for i, event_string in pairs(event_string_list) do
         table.insert(group_ids, gid)
     end
 end
-for i, event_string in pairs(gpu_event_string_list) do
-    if event_string:len() > 0 then
-        local gid = likwid.gpuAddEventSet(event_string)
-        if gid < 0 then
-            likwid.putGpuTopology()
-            likwid.putConfiguration()
-            likwid.gpuFinalize()
-            os.exit(1)
+if gpusSupported then
+    for i, event_string in pairs(gpu_event_string_list) do
+        if event_string:len() > 0 then
+            local gid = likwid.gpuAddEventSet(event_string)
+            if gid < 0 then
+                likwid.putGpuTopology()
+                likwid.putConfiguration()
+                likwid.gpuFinalize()
+                os.exit(1)
+            end
+            table.insert(gpugroups, gid)
         end
-        table.insert(gpugroups, gid)
     end
 end
 if #group_ids == 0 and not (#gpu_event_string_list > 0 and use_marker) then
@@ -927,6 +976,13 @@ if #event_string_list > 0 then
     end
 end
 
+if gpusSupported and #gpu_event_string_list > 0 then
+    activeNvGroup = gpugroups[1]
+    if outfile == nil then
+        print_stdout(likwid.hline)
+    end
+end
+
 
 
 
@@ -1092,36 +1148,40 @@ end
 
 
 if use_marker == true then
-    if likwid.access(markerFile, "e") >= 0 then
-        results, metrics = likwid.getMarkerResults(markerFile, cpulist, nan2value)
-        if not results then
-            print_stderr("Failure reading Marker API result file.")
-        elseif #results == 0 then
-            print_stderr("No regions could be found in Marker API result file.")
-        else
-            for r=1, #results do
-                likwid.printOutput(results[r], metrics[r], cpulist, r, print_stats)
+    if #event_string_list > 0 then
+        if likwid.access(markerFile, "e") >= 0 then
+            results, metrics = likwid.getMarkerResults(markerFile, cpulist, nan2value)
+            if not results then
+                print_stderr("Failure reading Marker API result file.")
+            elseif #results == 0 then
+                print_stderr("No regions could be found in Marker API result file.")
+            else
+                for r=1, #results do
+                    likwid.printOutput(results[r], metrics[r], cpulist, r, print_stats)
+                end
             end
+            os.remove(markerFile)
+        else
+            print_stderr("Marker API result file does not exist. This may happen if the application has not called LIKWID_MARKER_CLOSE.")
         end
-        os.remove(markerFile)
-    else
-        print_stderr("Marker API result file does not exist. This may happen if the application has not called LIKWID_MARKER_CLOSE.")
     end
-    if likwid.access(gpuMarkerFile, "e") >= 0 then
-        results, metrics = likwid.getGpuMarkerResults(gpuMarkerFile, markergpulist, nan2value)
-        if not results then
-            print_stderr("Failure reading GPU Marker API result file.")
-        elseif #results == 0 then
-            print_stderr("No regions could be found in GPU Marker API result file.")
-        else
-            for r=1, #results do
-                likwid.printGpuOutput(results[r], metrics[r], gpulist, r, print_stats)
+    if gpusSupported and #gpu_event_string_list > 0 then
+        if likwid.access(gpuMarkerFile, "e") >= 0 then
+            results, metrics = likwid.getGpuMarkerResults(gpuMarkerFile, markergpulist, nan2value)
+            if not results then
+                print_stderr("Failure reading GPU Marker API result file.")
+            elseif #results == 0 then
+                print_stderr("No regions could be found in GPU Marker API result file.")
+            else
+                for r=1, #results do
+                    likwid.printGpuOutput(results[r], metrics[r], gpulist, r, print_stats)
+                end
             end
+            likwid.destroyGpuMarkerFile()
+            os.remove(gpuMarkerFile)
+        else
+            print_stderr("GPU Marker API result file does not exist. This may happen if the application has not called LIKWID_GPUMARKER_CLOSE.")
         end
-        likwid.destroyGpuMarkerFile()
-        os.remove(gpuMarkerFile)
-    else
-        print_stderr("GPU Marker API result file does not exist. This may happen if the application has not called LIKWID_GPUMARKER_CLOSE.")
     end
 elseif use_timeline == false then
     results = likwid.getResults(nan2value)
@@ -1167,9 +1227,9 @@ if outfile then
     end
 end
 
-if #gpu_event_string_list > 0 then
-    likwid.gpuFinalize()
-end
+--if gpusSupported and #gpu_event_string_list > 0 then
+--    likwid.gpuFinalize()
+--end
 likwid.finalize()
 likwid.putTopology()
 likwid.putNumaInfo()
diff --git a/src/applications/likwid-topology.lua b/src/applications/likwid-topology.lua
index ad615cf18..469e9f881 100644
--- a/src/applications/likwid-topology.lua
+++ b/src/applications/likwid-topology.lua
@@ -259,31 +259,32 @@ else
     end
 end
 
-
-gputopo = likwid.getGpuTopology()
-if gputopo then
-    table.insert(output_csv, likwid.sline)
-    table.insert(output_csv, "GPU Topology")
-    table.insert(output_csv, likwid.sline)
-    table.insert(output_csv, string.format("GPU count:\t\t%d", gputopo["numDevices"]))
-    table.insert(output_csv, likwid.hline)
-
-    for i=1, gputopo["numDevices"] do
-        gpu = gputopo["devices"][i]
-        table.insert(output_csv, string.format("STRUCT,GPU Topology %d,9", gpu["id"]))
-        table.insert(output_csv, string.format("ID:\t\t\t%d", gpu["id"]))
-        table.insert(output_csv, string.format("Name:\t\t\t%s", gpu["name"]))
-        table.insert(output_csv, string.format("Compute capability:\t%d.%d", gpu["ccapMajor"], gpu["ccapMinor"]))
-        table.insert(output_csv, string.format("L2 size:\t\t%.2f MB", gpu["l2Size"]/(1024*1024)))
-        table.insert(output_csv, string.format("Memory:\t\t\t%.2f GB", gpu["memory"]/(1024*1024*1024)))
-        table.insert(output_csv, string.format("SIMD width:\t\t%d", gpu["simdWidth"]))
-        table.insert(output_csv, string.format("Clock rate:\t\t%d kHz", gpu["clockRatekHz"]))
-        table.insert(output_csv, string.format("Memory clock rate:\t%d kHz", gpu["memClockRatekHz"]))
-        table.insert(output_csv, string.format("Attached to NUMA node:\t%d", gpu["numaNode"]))
+if likwid.gpuSupported() then
+    gputopo = likwid.getGpuTopology()
+    if gputopo then
+        table.insert(output_csv, likwid.sline)
+        table.insert(output_csv, "GPU Topology")
+        table.insert(output_csv, likwid.sline)
+        table.insert(output_csv, string.format("GPU count:\t\t%d", gputopo["numDevices"]))
         table.insert(output_csv, likwid.hline)
+
+        for i=1, gputopo["numDevices"] do
+            gpu = gputopo["devices"][i]
+            table.insert(output_csv, string.format("STRUCT,GPU Topology %d,9", gpu["id"]))
+            table.insert(output_csv, string.format("ID:\t\t\t%d", gpu["id"]))
+            table.insert(output_csv, string.format("Name:\t\t\t%s", gpu["name"]))
+            table.insert(output_csv, string.format("Compute capability:\t%d.%d", gpu["ccapMajor"], gpu["ccapMinor"]))
+            table.insert(output_csv, string.format("L2 size:\t\t%.2f MB", gpu["l2Size"]/(1024*1024)))
+            table.insert(output_csv, string.format("Memory:\t\t\t%.2f GB", gpu["memory"]/(1024*1024*1024)))
+            table.insert(output_csv, string.format("SIMD width:\t\t%d", gpu["simdWidth"]))
+            table.insert(output_csv, string.format("Clock rate:\t\t%d kHz", gpu["clockRatekHz"]))
+            table.insert(output_csv, string.format("Memory clock rate:\t%d kHz", gpu["memClockRatekHz"]))
+            table.insert(output_csv, string.format("Attached to NUMA node:\t%d", gpu["numaNode"]))
+            table.insert(output_csv, likwid.hline)
+        end
     end
+    likwid.putGpuTopology()
 end
-likwid.putGpuTopology()
 
 if print_csv then
     longest_line = 0
diff --git a/src/applications/likwid.lua b/src/applications/likwid.lua
index bacf65e06..b27e8701e 100644
--- a/src/applications/likwid.lua
+++ b/src/applications/likwid.lua
@@ -183,6 +183,7 @@ likwid.gpuGetNameOfMetric = likwid_gpuGetNameOfMetric
 likwid.gpuGetNameOfEvent = likwid_gpuGetNameOfEvent
 likwid.gpuGetNameOfCounter = likwid_gpuGetNameOfCounter
 
+likwid.gpuSupported = likwid_gpuSupported
 likwid.readGpuMarkerFile = likwid_readGpuMarkerFile
 likwid.destroyGpuMarkerFile = likwid_destroyGpuMarkerFile
 likwid.gpuMarkerNumRegions = likwid_gpuMarkerNumRegions
diff --git a/src/cpustring.c b/src/cpustring.c
index b7fdad3eb..fadac6a2f 100644
--- a/src/cpustring.c
+++ b/src/cpustring.c
@@ -949,6 +949,8 @@ sockstr_to_socklist(const char* sockstr, int* sockets, int length)
     return ret;
 }
 
+#ifdef LIKWID_WITH_NVMON
+
 static int valid_gpu(GpuTopology_t topo, int id)
 {
     for (int i = 0; i < topo->numDevices; i++)
@@ -1000,3 +1002,5 @@ gpustr_to_gpulist(const char* gpustr, int* gpulist, int length)
     }
     return insert;
 }
+
+#endif /* LIKWID_WITH_NVMON */
diff --git a/src/includes/likwid-gpumarker.h b/src/includes/likwid-gpumarker.h
index b5d458714..3c514f82c 100644
--- a/src/includes/likwid-gpumarker.h
+++ b/src/includes/likwid-gpumarker.h
@@ -31,7 +31,6 @@
 #ifndef LIKWID_GPUMARKER_H
 #define LIKWID_GPUMARKER_H
 
-#include <likwid.h>
 
 /** \addtogroup MarkerAPI Marker API module
 *  @{
@@ -75,6 +74,7 @@ Shortcut for likwid_gpuMarkerClose() if compiled with -DLIKWID_NVMON. Otherwise
 /** @}*/
 
 #ifdef LIKWID_NVMON
+#include <likwid.h>
 #define LIKWID_GPUMARKER_INIT likwid_gpuMarkerInit()
 #define LIKWID_GPUMARKER_THREADINIT likwid_gpuMarkerThreadInit()
 #define LIKWID_GPUMARKER_SWITCH likwid_gpuMarkerNextGroup()
diff --git a/src/includes/likwid.h b/src/includes/likwid.h
index 1fb641c0e..fbd053fa5 100644
--- a/src/includes/likwid.h
+++ b/src/includes/likwid.h
@@ -682,6 +682,7 @@ defined in the selection string.
 */
 extern int sockstr_to_socklist(const char* sockstr, int* sockets, int length)  __attribute__ ((visibility ("default") ));
 
+#ifdef LIKWID_WITH_NVMON
 /*! \brief Read GPU selection string and resolve to available GPUs numbers
 
 Reads the GPU selection string and fills the given list with the GPU numbers defined in the selection string.
@@ -692,6 +693,8 @@ Reads the GPU selection string and fills the given list with the GPU numbers def
 */
 extern int gpustr_to_gpulist(const char* gpustr, int* gpulist, int length)  __attribute__ ((visibility ("default") ));
 
+#endif /* LIKWID_WITH_NVMON */
+
 /** @}*/
 
 /*
@@ -1603,11 +1606,13 @@ extern uint64_t freq_getUncoreFreqCur(const int socket_id) __attribute__ ((visib
  *  @{
  */
 
+#ifdef LIKWID_WITH_NVMON
+
 typedef struct {
     int devid;
     int numaNode;
     char* name;
-    size_t mem;
+    unsigned long long mem;
     int ccapMajor;
     int ccapMinor;
     int maxThreadsPerBlock;
@@ -1784,6 +1789,8 @@ int nvmon_getNumberOfMetrics(int groupId) __attribute__ ((visibility ("default")
 
 /** @}*/
 
+#endif /* LIKWID_WITH_NVMON */
+
 #ifdef __cplusplus
 }
 #endif
diff --git a/src/includes/nvmon_cupti.h b/src/includes/nvmon_cupti.h
new file mode 100644
index 000000000..428721332
--- /dev/null
+++ b/src/includes/nvmon_cupti.h
@@ -0,0 +1,1000 @@
+/*
+ * =======================================================================================
+ *
+ *      Filename:  nvmon_cupti.h
+ *
+ *      Description:  Header File of nvmon module (CUPTI backend).
+ *
+ *      Version:   <VERSION>
+ *      Released:  <DATE>
+ *
+ *      Author:   Thomas Gruber (tg), thomas.gruber@googlemail.com
+ *      Project:  likwid
+ *
+ *      Copyright (C) 2019 RRZE, University Erlangen-Nuremberg
+ *
+ *      This program is free software: you can redistribute it and/or modify it under
+ *      the terms of the GNU General Public License as published by the Free Software
+ *      Foundation, either version 3 of the License, or (at your option) any later
+ *      version.
+ *
+ *      This program is distributed in the hope that it will be useful, but WITHOUT ANY
+ *      WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A
+ *      PARTICULAR PURPOSE.  See the GNU General Public License for more details.
+ *
+ *      You should have received a copy of the GNU General Public License along with
+ *      this program.  If not, see <http://www.gnu.org/licenses/>.
+ *
+ * =======================================================================================
+ */
+
+#ifndef LIKWID_NVMON_CUPTI_H
+#define LIKWID_NVMON_CUPTI_H
+
+
+#include <cuda.h>
+#include <cupti.h>
+
+/* Copy from PAPI's cuda component (BSD License)
+ * @author  Asim YarKhan yarkhan@icl.utk.edu (updated in 2017 to support CUDA metrics)
+ * @author  Asim YarKhan yarkhan@icl.utk.edu (updated in 2015 for multiple CUDA contexts/devices)
+ * @author  Heike Jagode (First version, in collaboration with Robert Dietrich, TU Dresden) jagode@icl.utk.edu
+ */
+void (*_dl_non_dynamic_init) (void) __attribute__ ((weak));
+
+#define CU_CALL( call, handleerror )                                    \
+    do {                                                                \
+        CUresult _status = (call);                                      \
+        if (_status != CUDA_SUCCESS) {                                  \
+            fprintf(stderr, "Error: function %s failed with error %d.\n", #call, _status); \
+            handleerror;                                                \
+        }                                                               \
+    } while (0)
+
+#define CUPTI_CALL(call, handleerror)                                 \
+    do {                                                                \
+        CUptiResult _status = (call);                                   \
+        if (_status != CUPTI_SUCCESS) {                                 \
+            const char *errstr;                                         \
+            (*cuptiGetResultString)(_status, &errstr);               \
+            fprintf(stderr, "Error: function %s failed with error %s.\n", #call, errstr); \
+            handleerror;                                                \
+        }                                                               \
+    } while (0)
+
+#define CUDA_CALL( call, handleerror )                                \
+    do {                                                                \
+        cudaError_t _status = (call);                                   \
+        if (_status != cudaSuccess) {                                   \
+            fprintf(stderr, "Error: function %s failed with error %d.\n", #call, _status); \
+            handleerror;                                                \
+        }                                                               \
+    } while (0)
+
+#define CUAPIWEAK __attribute__( ( weak ) )
+#define DECLARECUFUNC(funcname, funcsig) CUresult CUAPIWEAK funcname funcsig;  CUresult( *funcname##Ptr ) funcsig;
+DECLARECUFUNC(cuCtxGetCurrent, (CUcontext *));
+DECLARECUFUNC(cuCtxSetCurrent, (CUcontext));
+DECLARECUFUNC(cuCtxDestroy, (CUcontext));
+DECLARECUFUNC(cuDeviceGet, (CUdevice *, int));
+DECLARECUFUNC(cuDeviceGetCount, (int *));
+DECLARECUFUNC(cuDeviceGetName, (char *, int, CUdevice));
+DECLARECUFUNC(cuInit, (unsigned int));
+DECLARECUFUNC(cuCtxPopCurrent, (CUcontext * pctx));
+DECLARECUFUNC(cuCtxPushCurrent, (CUcontext pctx));
+DECLARECUFUNC(cuCtxSynchronize, ());
+
+#define CUDAAPIWEAK __attribute__( ( weak ) )
+#define DECLARECUDAFUNC(funcname, funcsig) cudaError_t CUDAAPIWEAK funcname funcsig;  cudaError_t( *funcname##Ptr ) funcsig;
+DECLARECUDAFUNC(cudaGetDevice, (int *));
+DECLARECUDAFUNC(cudaSetDevice, (int));
+DECLARECUDAFUNC(cudaFree, (void *));
+
+#define CUPTIAPIWEAK __attribute__( ( weak ) )
+#define DECLARECUPTIFUNC(funcname, funcsig) CUptiResult CUPTIAPIWEAK funcname funcsig;  CUptiResult( *funcname##Ptr ) funcsig;
+DECLARECUPTIFUNC(cuptiEventGroupGetAttribute, (CUpti_EventGroup eventGroup, CUpti_EventGroupAttribute attrib, size_t * valueSize, void *value));
+DECLARECUPTIFUNC(cuptiDeviceGetEventDomainAttribute, (CUdevice device, CUpti_EventDomainID eventDomain, CUpti_EventDomainAttribute attrib, size_t * valueSize, void *value));
+DECLARECUPTIFUNC(cuptiEventGroupReadEvent, (CUpti_EventGroup eventGroup, CUpti_ReadEventFlags flags, CUpti_EventID event, size_t * eventValueBufferSizeBytes, uint64_t *        eventValueBuffer));
+DECLARECUPTIFUNC(cuptiEventGroupSetAttribute, (CUpti_EventGroup eventGroup, CUpti_EventGroupAttribute attrib, size_t valueSize, void *value));
+DECLARECUPTIFUNC(cuptiEventGroupSetDisable, (CUpti_EventGroupSet * eventGroupSet));
+DECLARECUPTIFUNC(cuptiEventGroupSetEnable, (CUpti_EventGroupSet * eventGroupSet));
+DECLARECUPTIFUNC(cuptiEventGroupSetsCreate, (CUcontext context, size_t eventIdArraySizeBytes, CUpti_EventID * eventIdArray, CUpti_EventGroupSets ** eventGroupPasses));
+DECLARECUPTIFUNC(cuptiEventGroupSetsDestroy, (CUpti_EventGroupSets * eventGroupSets));
+DECLARECUPTIFUNC(cuptiGetTimestamp, (uint64_t * timestamp));
+DECLARECUPTIFUNC(cuptiSetEventCollectionMode, (CUcontext context, CUpti_EventCollectionMode mode));
+DECLARECUPTIFUNC(cuptiDeviceEnumEventDomains, (CUdevice, size_t *, CUpti_EventDomainID *));
+DECLARECUPTIFUNC(cuptiDeviceGetNumEventDomains, (CUdevice, uint32_t *));
+DECLARECUPTIFUNC(cuptiEventDomainEnumEvents, (CUpti_EventDomainID, size_t *, CUpti_EventID *));
+DECLARECUPTIFUNC(cuptiEventDomainGetNumEvents, (CUpti_EventDomainID, uint32_t *));
+DECLARECUPTIFUNC(cuptiEventGetAttribute, (CUpti_EventID, CUpti_EventAttribute, size_t *, void *));
+DECLARECUPTIFUNC(cuptiGetResultString, (CUptiResult result, const char **str));
+
+
+static void *dl_libcuda = NULL;
+static void *dl_libcudart = NULL;
+static void *dl_libcupti = NULL;
+
+static int
+link_cputi_libraries(void)
+{
+#define DLSYM_AND_CHECK( dllib, name ) dlsym( dllib, name ); if ( dlerror() != NULL ) { return -1; }
+
+    /* Attempt to guess if we were statically linked to libc, if so bail */
+    if(_dl_non_dynamic_init != NULL) {
+        return -1;
+    }
+    /* Need to link in the cuda libraries, if not found disable the component */
+    dl_libcuda = dlopen("libcuda.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!dl_libcuda)
+    {
+        fprintf(stderr, "CUDA library libcuda.so not found.");
+        return -1;
+    }
+    cuCtxGetCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxGetCurrent");
+    cuCtxSetCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxSetCurrent");
+    cuDeviceGetPtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGet");
+    cuDeviceGetCountPtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGetCount");
+    cuDeviceGetNamePtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGetName");
+    cuInitPtr = DLSYM_AND_CHECK(dl_libcuda, "cuInit");
+    cuCtxPopCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxPopCurrent");
+    cuCtxPushCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxPushCurrent");
+    cuCtxSynchronizePtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxSynchronize");
+    cuCtxDestroyPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxDestroy");
+
+    dl_libcudart = dlopen("libcudart.so", RTLD_NOW | RTLD_GLOBAL | RTLD_NODELETE);
+    if (!dl_libcudart)
+    {
+        fprintf(stderr, "CUDA runtime library libcudart.so not found.");
+        return -1;
+    }
+    cudaGetDevicePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaGetDevice");
+    cudaSetDevicePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaSetDevice");
+    cudaFreePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaFree");
+
+    dl_libcupti = dlopen("libcupti.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!dl_libcupti)
+    {
+        fprintf(stderr, "CUDA runtime library libcupti.so not found.");
+        return -1;
+    }
+    cuptiDeviceGetEventDomainAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceGetEventDomainAttribute");
+    cuptiEventGroupGetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupGetAttribute");
+    cuptiEventGroupReadEventPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupReadEvent");
+    cuptiEventGroupSetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetAttribute");
+    cuptiEventGroupSetDisablePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetDisable");
+    cuptiEventGroupSetEnablePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetEnable");
+    cuptiEventGroupSetsCreatePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetsCreate");
+    cuptiEventGroupSetsDestroyPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetsDestroy");
+    cuptiGetTimestampPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiGetTimestamp");
+    cuptiSetEventCollectionModePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiSetEventCollectionMode");
+    cuptiDeviceEnumEventDomainsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceEnumEventDomains");
+    cuptiDeviceGetNumEventDomainsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceGetNumEventDomains");
+    cuptiEventDomainEnumEventsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventDomainEnumEvents");
+    cuptiEventDomainGetNumEventsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventDomainGetNumEvents");
+    cuptiEventGetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGetAttribute");
+    cuptiGetResultStringPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiGetResultString");
+    return 0;
+}
+
+
+static int check_nv_context(NvmonDevice_t device, CUcontext currentContext)
+{
+    int j = 0;
+    int need_pop = 0;
+    GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Current context %ld DevContext %ld, currentContext, device->context);
+    if (!device->context)
+    {
+        int context_of_dev = -1;
+        for (j = 0; j < nvGroupSet->numberOfGPUs; j++)
+        {
+            NvmonDevice_t dev = &nvGroupSet->gpus[j];
+            if (dev->context == currentContext)
+            {
+                context_of_dev = j;
+                break;
+            }
+        }
+        if (context_of_dev < 0 && !device->context)
+        {
+            device->context = currentContext;
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Reuse context %ld for device %d, device->context, device->deviceId);
+        }
+        else
+        {
+            CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
+            CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
+            CU_CALL((*cuCtxGetCurrentPtr)(&device->context), return -EFAULT);
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, New context %ld for device %d, device->context, device->deviceId);
+        }
+    }
+    else if (device->context != currentContext)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Use context %ld for device %d, device->context, device->deviceId);
+        CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
+        need_pop = 1;
+    }
+    else
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Context %ld fits for device %d, device->context, device->deviceId);
+    }
+    return need_pop;
+}
+
+static int
+init_cuda(void)
+{
+    CUresult cuErr = (*cuInitPtr)(0);
+    if (cuErr != CUDA_SUCCESS)
+    {
+        fprintf(stderr, "CUDA cannot be found and initialized (cuInit failed).\n");
+        return -ENODEV;
+    }
+    return 0;
+}
+
+static int
+get_numDevices(void)
+{
+    CUresult cuErr;
+    int count = 0;
+    cuErr = (*cuDeviceGetCountPtr)(&count);
+    if(cuErr == CUDA_ERROR_NOT_INITIALIZED)
+    {
+        int ret = init_cuda();
+        if (ret == 0)
+        {
+            cuErr = (*cuDeviceGetCountPtr)(&count);
+        }
+        else
+        {
+            return ret;
+        }
+    }
+    return count;
+}
+
+void nvmon_cupti_freeDevice(NvmonDevice_t dev)
+{
+    int j = 0;
+    GHashTableIter iter;
+    char* name = NULL;
+    uint32_t *id = NULL;
+    NvmonEvent_t event = NULL;
+    if (dev)
+    {
+        if (dev->context)
+        {
+            CU_CALL((*cuCtxDestroyPtr)(dev->context), j++);
+            dev->context = NULL;
+        }
+        if (dev->cuEventSets)
+        {
+            CUPTI_CALL((*cuptiEventGroupSetsDestroyPtr)(dev->cuEventSets), j++);
+            dev->cuEventSets = NULL;
+        }
+
+        // The event objects are freed here
+
+        if (dev->activeEvents)
+        {
+            free(dev->activeEvents);
+            dev->activeEvents = NULL;
+            dev->numActiveEvents = 0;
+        }
+
+        if (dev->nvEventSets)
+        {
+            for (j = 0; j < dev->numNvEventSets; j++)
+            {
+                free(dev->nvEventSets[j].results);
+            }
+            free(dev->nvEventSets);
+            dev->nvEventSets = NULL;
+            dev->numNvEventSets = 0;
+        }
+        if (dev->allevents)
+        {
+            for (j = 0; j < dev->numAllEvents; j++)
+            {
+                if (dev->allevents[j])
+                {
+                    free(dev->allevents[j]);
+                }
+            }
+        }
+        memset(dev, 0, sizeof(NvmonDevice));
+    }
+    return;
+}
+
+int
+nvmon_cupti_createDevice(int id, NvmonDevice *dev)
+{
+    int j = 0, k = 0, c = 0;
+    int numDomains = 0;
+    CUpti_EventDomainID* eventDomainIds = NULL;
+    int eventIdx = 0;
+    uint32_t totalEvents = 0;
+
+
+    dev->deviceId = id;
+    dev->cuEventSets = NULL;
+    dev->context = 0UL;
+    dev->activeEvents = NULL;
+    dev->numActiveEvents = 0;
+    dev->numNvEventSets = 0;
+    dev->nvEventSets = NULL;
+
+    if ((!dl_libcuda) || (!dl_libcudart) || (!dl_libcupti))
+    {
+        int err = link_cputi_libraries();
+        if (err < 0)
+        {
+            return -1;
+        }
+    }
+
+    // Assign device ID and get cuDevice from CUDA
+    CU_CALL((*cuDeviceGetPtr)(&dev->cuDevice, id), return -1);
+
+    // Get the number of event domains of the device
+    CUPTI_CALL((*cuptiDeviceGetNumEventDomainsPtr)(dev->cuDevice, &numDomains), return -1);
+    GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Nvmon: Dev %d Domains %d, id, numDomains);
+
+    // Get the domain IDs for the device
+    size_t domainarraysize = numDomains * sizeof(CUpti_EventDomainID);
+    eventDomainIds = malloc(domainarraysize);
+    if (!eventDomainIds) return -ENOMEM;
+    CUPTI_CALL((*cuptiDeviceEnumEventDomainsPtr)(dev->cuDevice, &domainarraysize, eventDomainIds), return -1);
+
+
+    // Count the events in all domains to allocate the event list
+    dev->numAllEvents = 0;
+    for (j = 0; j < numDomains; j++)
+    {
+        uint32_t domainNumEvents = 0;
+        CUpti_EventDomainID domainID = eventDomainIds[j];
+        CUPTI_CALL((*cuptiEventDomainGetNumEventsPtr)(domainID, &domainNumEvents), return -1);
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Nvmon: Dev %d Domain %d Events %d, id, j, domainNumEvents);
+        totalEvents += domainNumEvents;
+    }
+    // Now we now how many events are provided by the device, so allocate a big enough event list
+    dev->allevents = malloc(totalEvents * sizeof(NvmonEvent_t));
+    if (!dev->allevents)
+    {
+        free(eventDomainIds);
+        return -ENOMEM;
+    }
+
+    // We use hash tables for faster access
+    dev->eventHash = g_hash_table_new(g_str_hash, g_str_equal);
+    dev->evIdHash = g_hash_table_new(g_int64_hash, g_int64_equal);
+
+    for (j = 0; j < numDomains; j++)
+    {
+        uint32_t domainNumEvents = 0;
+        CUpti_EventDomainID domainID = eventDomainIds[j];
+        // How many events are provided by the domain
+        CUPTI_CALL((*cuptiEventDomainGetNumEventsPtr)(domainID, &domainNumEvents), return -1);
+        size_t tmpSize = domainNumEvents * sizeof(CUpti_EventID);
+        // Allocate space for all CUPTI event IDs in the domain
+        CUpti_EventID* cuEventIds = malloc(tmpSize);
+        // Get the CUPTI events
+        CUPTI_CALL((*cuptiEventDomainEnumEventsPtr)(domainID, &tmpSize, cuEventIds), return -1);
+        for (k = 0; k < domainNumEvents; k++)
+        {
+            CUpti_EventID eventId = cuEventIds[k];
+            // Name and description are limited in length
+            size_t tmpSizeBytesName = (NVMON_DEFAULT_STR_LEN-1) * sizeof(char);
+            size_t tmpSizeBytesDesc = (NVMON_DEFAULT_STR_LEN-1) * sizeof(char);
+
+            NvmonEvent_t event = malloc(sizeof(NvmonEvent));
+            if (!event)
+            {
+                free(cuEventIds);
+                free(dev->allevents);
+                free(eventDomainIds);
+                return -ENOMEM;
+            }
+
+            // Get event name and description
+            CUPTI_CALL((*cuptiEventGetAttributePtr)(eventId, CUPTI_EVENT_ATTR_NAME, &tmpSizeBytesName, event->name), return -1);
+            CUPTI_CALL((*cuptiEventGetAttributePtr)(eventId, CUPTI_EVENT_ATTR_LONG_DESCRIPTION, &tmpSizeBytesDesc, event->description), return -1);
+            event->name[tmpSizeBytesName/sizeof(char)] = '\0';
+            event->description[tmpSizeBytesDesc/sizeof(char)] = '\0';
+            // LIKWID events are all uppercase, so transform the event names
+            c = 0;
+            while (event->name[c] != '\0')
+            {
+                event->name[c] = toupper(event->name[c]);
+                c++;
+            }
+            // Save all gathered information in a NvmonEvent object
+
+            event->cuEventId = eventId;
+            event->eventId = dev->numAllEvents;
+            event->cuDomainId = domainID;
+            event->domainId = j;
+            event->type = NVMON_CUPTI_EVENT;
+            event->active = 0;
+            //GPUDEBUG_PRINT(DEBUGLEV_DETAIL, New Event %d CuEvent %d Domain %d CuDomain %d Name %s, event->eventId, (int)event->cuEventId, event->domainId, (int)event->cuDomainId, event->name);
+            // Add the object to the event list
+            dev->allevents[dev->numAllEvents] = event;
+            dev->numAllEvents++;
+            // Add the object to the hash tables
+            char* nameKey = g_strdup(event->name);
+
+            CUpti_EventID* idKey = malloc(sizeof(CUpti_EventID));
+            if (idKey)
+            {
+                *idKey = event->cuEventId;
+                // Key is event name
+                g_hash_table_insert(dev->eventHash, (gpointer)nameKey, (gpointer)event);
+                // Key is CUPTI event ID
+                g_hash_table_insert(dev->evIdHash, (gpointer)idKey, (gpointer)event);
+            }
+        }
+    }
+
+/*    CUDA_CALL((*cudaSetDevicePtr)(dev->deviceId), return -EFAULT);*/
+/*    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);*/
+/*    CU_CALL((*cuCtxGetCurrentPtr)(&dev->context), return -EFAULT);*/
+    return 0;
+}
+
+
+
+int nvmon_cupti_getEventsOfGpu(int gpuId, NvmonEventList_t* list)
+{
+    int ret = 0;
+    NvmonDevice device;
+    int err = nvmon_cupti_createDevice(gpuId, &device);
+    if (!err)
+    {
+        NvmonEventList_t l = malloc(sizeof(NvmonEventList));
+        if (l)
+        {
+            l->events = malloc(sizeof(NvmonEventListEntry) * device.numAllEvents);
+            if (l->events)
+            {
+                for (int i = 0; i < device.numAllEvents; i++)
+                {
+                    NvmonEventListEntry* out = &l->events[i];
+                    NvmonEvent_t event = device.allevents[i];
+                    out->name = malloc(strlen(event->name)+2);
+                    if (out->name)
+                    {
+                        ret = snprintf(out->name, strlen(event->name)+1, "%s", event->name);
+                        if (ret > 0)
+                        {
+                            out->name[ret] = '\0';
+                        }
+                    }
+                    out->desc = malloc(strlen(event->description)+2);
+                    if (out->desc)
+                    {
+                        ret = snprintf(out->desc, strlen(event->description)+1, "%s", event->description);
+                        if (ret > 0)
+                        {
+                            out->desc[ret] = '\0';
+                        }
+                    }
+                    out->limit = malloc(10*sizeof(char));
+                    if (out->limit)
+                    {
+                        switch (event->type)
+                        {
+                            case NVMON_CUPTI_EVENT:
+                                ret = snprintf(out->limit, 9, "GPU");
+                                if (ret > 0) out->limit[ret] = '\0';
+                                break;
+                        }
+                    }
+                }
+                l->numEvents = device.numAllEvents;
+                *list = l;
+            }
+            else
+            {
+                free(l);
+                nvmon_cupti_freeDevice(&device);
+                return -ENOMEM;
+            }
+        }
+    }
+    else
+    {
+        ERROR_PRINT(No such device %d, gpuId);
+    }
+    return 0;
+}
+
+
+
+int
+nvmon_cupti_addEventSets(NvmonDevice_t device, const char* eventString)
+{
+    int i = 0;
+    int err = 0;
+    int curDeviceId = -1;
+    CUcontext curContext;
+    struct bstrList* eventtokens = NULL;
+
+    CUDA_CALL((*cudaGetDevicePtr)(&curDeviceId), return -EFAULT);
+    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+
+    bstring eventBString = bfromcstr(eventString);
+    eventtokens = bsplit(eventBString, ',');
+    bdestroy(eventBString);
+
+    if (curDeviceId != device->deviceId)
+    {
+        CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
+    }
+
+    size_t sizeBytes = (eventtokens->qty) * sizeof(CUpti_EventID);
+
+    int popContext = 0;
+    CUpti_EventGroupSets * cuEventSets = NULL;
+
+    NvmonEventSet* tmpEventSet = realloc(device->nvEventSets, (device->numNvEventSets+1)*sizeof(NvmonEventSet));
+    if (!tmpEventSet)
+    {
+        ERROR_PRINT(Cannot enlarge GPU %d eventSet list, device->deviceId);
+        return -ENOMEM;
+    }
+    device->nvEventSets = tmpEventSet;
+    NvmonEventSet* devEventSet = &device->nvEventSets[device->numNvEventSets];
+
+    devEventSet->nvEvents = (NvmonEvent_t*) malloc(eventtokens->qty * sizeof(NvmonEvent_t));
+    if (devEventSet->nvEvents == NULL)
+    {
+        ERROR_PRINT(Cannot allocate event list for group %d\n, groupSet->numberOfActiveGroups);
+        return -ENOMEM;
+    }
+    devEventSet->cuEventIDs = (CUpti_EventID*) malloc(eventtokens->qty * sizeof(CUpti_EventID));
+    if (devEventSet->cuEventIDs == NULL)
+    {
+        ERROR_PRINT(Cannot allocate event ID list for group %d\n, groupSet->numberOfActiveGroups);
+        free(devEventSet->nvEvents);
+        return -ENOMEM;
+    }
+    devEventSet->results = malloc(eventtokens->qty * sizeof(NvmonEventResult));
+    if (devEventSet->cuEventIDs == NULL)
+    {
+        ERROR_PRINT(Cannot allocate result list for group %d\n, groupSet->numberOfActiveGroups);
+        free(devEventSet->cuEventIDs);
+        free(devEventSet->nvEvents);
+        return -ENOMEM;
+    }
+    memset(devEventSet->results, 0, eventtokens->qty * sizeof(NvmonEventResult));
+    devEventSet->numberOfEvents = 0;
+
+    // If the device has no context, check whether the current context is
+    // used already by another device
+    popContext = check_nv_context(device, curContext);
+
+
+    CUPTI_CALL((*cuptiSetEventCollectionModePtr)(device->context, CUPTI_EVENT_COLLECTION_MODE_KERNEL), return -EFAULT);
+
+    for (i = 0; i < eventtokens->qty; i++)
+    {
+        struct bstrList* evset = bsplit(eventtokens->entry[i], ':');
+        if (evset->qty != 2)
+        {
+            ERROR_PRINT(NVMON: Event %s invalid: Format <event>:<gpucounter>, bdata(eventtokens->entry[i]));
+        }
+        if (blength(evset->entry[0]) == 0 || blength(evset->entry[1]) == 0)
+        {
+            ERROR_PRINT(NVMON: Event %s invalid: Format <event>:<gpucounter>, bdata(eventtokens->entry[i]));
+        }
+        NvmonEvent_t event = g_hash_table_lookup(device->eventHash, (gpointer)bdata(evset->entry[0]));
+        if (!event)
+        {
+            GPUDEBUG_PRINT(DEBUGLEV_INFO, NVMON: Event %s unknown. Skipping..., bdata(evset->entry[0]));
+            continue; //unknown event
+        }
+        else
+        {
+            devEventSet->cuEventIDs[devEventSet->numberOfEvents] = event->cuEventId;
+            devEventSet->nvEvents[devEventSet->numberOfEvents] = event;
+            devEventSet->numberOfEvents++;
+            size_t s = devEventSet->numberOfEvents*sizeof(CUpti_EventID);
+
+            CUPTI_CALL((*cuptiEventGroupSetsCreatePtr)(device->context, s, devEventSet->cuEventIDs, &cuEventSets), devEventSet->numberOfEvents--;);
+            if (cuEventSets->numSets > 1)
+            {
+                ERROR_PRINT(Error adding event %s. Multiple measurement runs are required. skipping event ..., bdata(evset->entry[i]));
+                continue;
+            }
+        }
+    }
+    devEventSet->id = nvGroupSet->numberOfActiveGroups;
+    // Create an eventset with the currently configured and current event
+    if(popContext)
+    {
+        CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+    }
+    device->numNvEventSets++;
+}
+
+int nvmon_cupti_setupCounters(NvmonDevice_t device, NvmonEventSet* eventSet)
+{
+    int err = 0;
+    int popContext = 0;
+    int oldDevId = -1;
+    CUpti_EventGroupSets * cuEventSets = NULL;
+    CUcontext curContext;
+
+    if (eventSet->numberOfEvents == 0)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, eventSet->id);
+        return -EINVAL;
+    }
+    // Currently we are on which device?
+    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    // This is a workaround to (eventually create and) get the current context
+    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+
+    popContext = check_nv_context(device, curContext);
+
+
+    size_t grpEventIdsSize = eventSet->numberOfEvents * sizeof(CUpti_EventID);
+    CUPTI_CALL((*cuptiEventGroupSetsCreatePtr)(device->context, grpEventIdsSize, eventSet->cuEventIDs, &cuEventSets), err = -1;);
+    // Allocate temporary array to hold the group event IDs
+    CUpti_EventID *grpEventIds = malloc(grpEventIdsSize);
+    if (!grpEventIds)
+    {
+        return -ENOMEM;
+    }
+    // Delete current activeEvent list
+    if (device->activeEvents)
+    {
+        free(device->activeEvents);
+        device->activeEvents = NULL;
+        device->numActiveEvents = 0;
+    }
+    // Delete current activeCuGroups list
+    if (device->activeCuGroups)
+    {
+        free(device->activeCuGroups);
+        device->activeCuGroups = NULL;
+        device->numActiveCuGroups = 0;
+    }
+
+    // Create a new activeEvent list
+    device->activeEvents = malloc(eventSet->numberOfEvents * sizeof(NvmonActiveEvent));
+    if (!device->activeEvents)
+    {
+        free(grpEventIds);
+        grpEventIds = NULL;
+        return -ENOMEM;
+    }
+    // Create a new activeCuGroups list
+    device->activeCuGroups = malloc(cuEventSets->numSets * sizeof(CUpti_EventGroupSet**));
+    if (!device->activeCuGroups)
+    {
+        free(device->activeEvents);
+        device->activeEvents = NULL;
+        free(grpEventIds);
+        grpEventIds = NULL;
+        return -ENOMEM;
+    }
+
+
+    // Run over eventset and store all information we need for start/stop/reads in NvmonActiveEvent_t
+    CUpti_EventGroup curGroup;
+    uint32_t curNumInstances = 0, curNumTotalInstances = 0;
+    CUpti_EventGroupSet *curEventGroupSet;
+    for (int j = 0; j < cuEventSets->numSets; j++)
+    {
+        size_t sizeofuint32t = sizeof(uint32_t);
+        uint32_t numGEvents = 0, numGInstances = 0, numTotalGInstances = 0;
+        CUpti_EventGroupSet* groupset = &cuEventSets->sets[j];
+
+        for (int k = 0; k < groupset->numEventGroups; k++)
+        {
+            uint32_t one = 1;
+            CUpti_EventGroup group = groupset->eventGroups[k];
+            // Get the number of events in the group
+            CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_NUM_EVENTS, &sizeofuint32t, &numGEvents), return -EFAULT);
+            // Get the CUPTI event IDs in the group
+            CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_EVENTS, &grpEventIdsSize, grpEventIds), return -EFAULT);
+            // If we don't set this, each event has only a single instance but we want to measure all instances
+            CUPTI_CALL((*cuptiEventGroupSetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES, sizeof(uint32_t), &one), return -EFAULT);
+            // Get instance count for a group
+            CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_INSTANCE_COUNT, &sizeofuint32t, &numGInstances), return -EFAULT);
+
+            for (int l = 0; l < numGEvents; l++)
+            {
+                for (int m = 0; m <  eventSet->numberOfEvents; m++)
+                {
+                    if (eventSet->cuEventIDs[m] == grpEventIds[l])
+                    {
+                        CUpti_EventDomainID did = eventSet->nvEvents[m]->cuDomainId;
+                        // Get total instance count for a group
+                        CUPTI_CALL((*cuptiDeviceGetEventDomainAttributePtr)(device->cuDevice, did, CUPTI_EVENT_DOMAIN_ATTR_TOTAL_INSTANCE_COUNT, &sizeofuint32t, &numTotalGInstances), return -EFAULT);
+                        device->activeEvents[m].eventId = eventSet->nvEvents[m]->eventId;
+                        device->activeEvents[m].idxInSet = m;
+                        device->activeEvents[m].groupId = eventSet->id;
+                        device->activeEvents[m].cuEventId = eventSet->nvEvents[m]->cuEventId;
+                        device->activeEvents[m].cuDomainId = did;
+                        device->activeEvents[m].numTotalInstances = numTotalGInstances;
+                        device->activeEvents[m].cuGroup = group;
+                        device->activeEvents[m].cuGroupSet = groupset;
+                        device->activeEvents[m].numInstances = numGInstances;
+                        device->activeEvents[m].deviceId = device->deviceId;
+
+                        int found = 0;
+                        for (int i = 0; i < device->numActiveCuGroups; i++)
+                        {
+                            if (device->activeCuGroups[i] == groupset)
+                            {
+                                found = 1;
+                            }
+                        }
+                        if (!found)
+                        {
+                            device->activeCuGroups[device->numActiveCuGroups] = groupset;
+                            device->numActiveCuGroups++;
+                        }
+                        // Mark event as active. This is used to avoid measuring the same event on the same device twice
+                        eventSet->nvEvents[m]->active = 1;
+                        GPUDEBUG_PRINT(DEBUGLEV_INFO, Setup event %s (%d) for GPU %d, eventSet->nvEvents[m]->name, device->activeEvents[m].cuEventId, device->deviceId);
+                        device->numActiveEvents++;
+                    }
+                }
+            }
+        }
+    }
+    free(grpEventIds);
+    if(popContext)
+    {
+        CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+    }
+    return 0;
+}
+
+int nvmon_cupti_startCounters(NvmonDevice_t device)
+{
+    int j = 0;
+    CUcontext curContext;
+    int popContext = 0;
+    uint32_t one = 1;
+    uint64_t timestamp = 0;
+    int oldDevId = -1;
+
+    // Currently we are on which device?
+    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    // Take the timestamp, we assign it later for all devices
+    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
+    // This is a workaround to (eventually create and) get the current context
+    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+
+
+    //NvmonDevice_t device = &nvGroupSet->gpus[i];
+    if (device->numActiveCuGroups == 0)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
+        return 0;
+    }
+    if (device->deviceId != oldDevId)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
+        CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
+        CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    }
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+    device->timeStart = timestamp;
+    device->timeRead = timestamp;
+
+    // Are we in the proper context?
+    popContext = check_nv_context(device, curContext);
+
+    NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
+    for (j = 0; j < nvEventSet->numberOfEvents; j++)
+    {
+        NvmonEventResult* res = &nvEventSet->results[j];
+        res->startValue = 0.0;
+        res->stopValue = 0.0;
+        res->currentValue = 0.0;
+        res->fullValue = 0.0;
+        res->overflows = 0;
+    }
+
+    for (j = 0; j < device->numActiveCuGroups; j++)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Enable group %ld on Dev %d, device->activeCuGroups[j], device->deviceId);
+        CUPTI_CALL((*cuptiEventGroupSetEnablePtr)(device->activeCuGroups[j]), return -EFAULT);
+    }
+
+    // If we added the device context to the stack, pop it again
+    if (popContext)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
+        CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+    }
+}
+
+
+int nvmon_cupti_stopCounters(NvmonDevice_t device)
+{
+    int i = 0, j = 0, k = 0;
+    int oldDevId = -1;
+    uint64_t timestamp = 0;
+    CUcontext curContext;
+
+
+    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    // Take the timestamp, we assign it later for all devices
+    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
+
+
+    for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
+    {
+        int popContext = 0;
+        uint32_t one = 1;
+        int maxTotalInstances = 0;
+        size_t valuesSize = 0;
+        NvmonDevice_t device = &nvGroupSet->gpus[i];
+        if (device->deviceId != oldDevId)
+        {
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
+            CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
+            CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+        }
+        CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+        NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
+        if (device->numActiveCuGroups == 0)
+        {
+            GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
+            continue;
+        }
+        device->timeStop = timestamp;
+
+        // Are we in the proper context?
+        popContext = check_nv_context(device, curContext);
+
+        for (j = 0; j < device->numActiveEvents; j++)
+        {
+            maxTotalInstances = MAX(maxTotalInstances, device->activeEvents[j].numTotalInstances);
+        }
+        uint64_t *tmpValues = (uint64_t *) malloc(maxTotalInstances * sizeof(uint64_t));
+
+        for (j = 0; j < device->numActiveEvents; j++)
+        {
+            NvmonActiveEvent_t event = &device->activeEvents[j];
+            valuesSize = sizeof(uint64_t) * event->numTotalInstances;
+            memset(tmpValues, 0, valuesSize);
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Read Grp %ld Ev %ld for device %d, event->cuGroup, event->cuEventId, device->deviceId);
+            CUPTI_CALL((*cuptiEventGroupReadEventPtr)(event->cuGroup, CUPTI_EVENT_READ_FLAG_NONE, event->cuEventId, &valuesSize, tmpValues), free(tmpValues); return -EFAULT);
+            uint64_t valuesum = 0;
+            for (k = 0; k < event->numInstances; k++)
+            {
+                valuesum += tmpValues[k];
+            }
+            NvmonEventResult* res = &nvEventSet->results[event->idxInSet];
+            res->stopValue = (double)valuesum;
+            res->lastValue = res->currentValue;
+            res->fullValue += res->stopValue - res->startValue;
+            res->lastValue += res->stopValue - res->currentValue;
+            res->currentValue = (double)valuesum;
+        }
+        for (j = 0; j < device->numActiveCuGroups; j++)
+        {
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Disable group %ld, device->activeCuGroups[j]);
+            CUPTI_CALL((*cuptiEventGroupSetDisablePtr)(device->activeCuGroups[j]), return -EFAULT);
+        }
+        // If we added the device context to the stack, pop it again
+        if (popContext)
+        {
+            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
+            CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+        }
+    }
+}
+
+
+int nvmon_cupti_readCounters(NvmonDevice_t device)
+{
+    int j = 0, k = 0;
+    int oldDevId = -1;
+    uint64_t timestamp = 0;
+    CUcontext curContext;
+    size_t sizeofuint32num = sizeof(uint32_t);
+    int maxTotalInstances = 0;
+
+
+    for (j = 0; j < device->numActiveEvents; j++)
+    {
+        maxTotalInstances = MAX(maxTotalInstances, device->activeEvents[j].numTotalInstances);
+    }
+
+
+    // In this array we collect the instance values of an events (summed up later)
+    size_t valuesSize = sizeof(uint64_t) * maxTotalInstances;
+    uint64_t *tmpValues = (uint64_t *) malloc(valuesSize);
+    if (!tmpValues)
+    {
+        ERROR_PRINT(Not enough memory to allocate space for instance values);
+        return -ENOMEM;
+    }
+
+    // Currently we are on which device?
+    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    // Take the timestamp, we assign it later for all devices
+    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
+/*    // This is a workaround to (eventually create and) get the current context*/
+/*    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);*/
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+
+
+    int popContext = 0;
+    if (device->numActiveEvents == 0)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
+        return 0;
+    }
+    if (device->deviceId != oldDevId)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
+        CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
+        CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
+    }
+    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
+    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
+    device->timeRead = timestamp;
+
+    // Are we in the proper context?
+    popContext = check_nv_context(device, curContext);
+    // Synchronize devices. I'm not sure whether this is required as each
+    // device measures it's own events
+    CU_CALL((*cuCtxSynchronizePtr)(), return -EFAULT);
+    NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
+
+    for (j = 0; j < device->numActiveEvents; j++)
+    {
+        NvmonActiveEvent_t event = &device->activeEvents[j];
+        // Empty space for instance values
+        valuesSize = sizeof(uint64_t) * event->numTotalInstances;
+        memset(tmpValues, 0, valuesSize);
+        // Read all instance values
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Read Grp %ld Ev %ld for device %d, event->cuGroup, event->cuEventId, device->deviceId);
+        CUPTI_CALL((*cuptiEventGroupReadEventPtr)(event->cuGroup, CUPTI_EVENT_READ_FLAG_NONE, event->cuEventId, &valuesSize, tmpValues), return -EFAULT);
+        // Sum all instance values
+        uint64_t valuesum = 0;
+        for (k = 0; k < event->numInstances; k++)
+        {
+            valuesum += tmpValues[k];
+        }
+
+        NvmonEventResult* res = &nvEventSet->results[event->idxInSet];
+        res->lastValue = res->currentValue;
+        res->currentValue = (double)valuesum;
+        res->fullValue += res->currentValue - res->startValue;
+        res->lastValue += res->currentValue - res->lastValue;
+    }
+    if (popContext)
+    {
+        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
+        CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+    }
+    free(tmpValues);
+    return 0;
+}
+
+
+NvmonFunctions nvmon_cupti_functions = {
+    .freeDevice = nvmon_cupti_freeDevice,
+    .createDevice = nvmon_cupti_createDevice,
+    .getEventList = nvmon_cupti_getEventsOfGpu,
+    .addEvents = nvmon_cupti_addEventSets,
+    .setupCounters = nvmon_cupti_setupCounters,
+    .startCounters = nvmon_cupti_startCounters,
+    .readCounters = nvmon_cupti_readCounters,
+};
+
+#endif
diff --git a/src/includes/nvmon_nvml.h b/src/includes/nvmon_nvml.h
new file mode 100644
index 000000000..4ba63ece9
--- /dev/null
+++ b/src/includes/nvmon_nvml.h
@@ -0,0 +1,50 @@
+/*
+ * =======================================================================================
+ *
+ *      Filename:  nvmon_nvml.h
+ *
+ *      Description:  Header File of nvmon module (NVML backend).
+ *
+ *      Version:   <VERSION>
+ *      Released:  <DATE>
+ *
+ *      Author:   Thomas Gruber (tg), thomas.gruber@googlemail.com
+ *      Project:  likwid
+ *
+ *      Copyright (C) 2019 RRZE, University Erlangen-Nuremberg
+ *
+ *      This program is free software: you can redistribute it and/or modify it under
+ *      the terms of the GNU General Public License as published by the Free Software
+ *      Foundation, either version 3 of the License, or (at your option) any later
+ *      version.
+ *
+ *      This program is distributed in the hope that it will be useful, but WITHOUT ANY
+ *      WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A
+ *      PARTICULAR PURPOSE.  See the GNU General Public License for more details.
+ *
+ *      You should have received a copy of the GNU General Public License along with
+ *      this program.  If not, see <http://www.gnu.org/licenses/>.
+ *
+ * =======================================================================================
+ */
+#ifndef LIKWID_NVMON_NVML_H
+#define LIKWID_NVMON_NVML_H
+
+
+
+
+
+
+NvmonFunctions nvmon_nvml_functions = {
+    .freeDevice = NULL,
+    .createDevice = NULL,
+    .getEventList = NULL,
+    .addEvents = NULL,
+    .setupCounters = NULL,
+    .startCounters = NULL,
+    .readCounters = NULL,
+};
+
+
+
+#endif /* LIKWID_NVMON_NVML_H */
diff --git a/src/includes/nvmon_perfworks.h b/src/includes/nvmon_perfworks.h
new file mode 100644
index 000000000..f2e1b9d33
--- /dev/null
+++ b/src/includes/nvmon_perfworks.h
@@ -0,0 +1,274 @@
+/*
+ * =======================================================================================
+ *
+ *      Filename:  nvmon_cupti.h
+ *
+ *      Description:  Header File of nvmon module (PerfWorks backend).
+ *
+ *      Version:   <VERSION>
+ *      Released:  <DATE>
+ *
+ *      Author:   Thomas Gruber (tg), thomas.gruber@googlemail.com
+ *      Project:  likwid
+ *
+ *      Copyright (C) 2019 RRZE, University Erlangen-Nuremberg
+ *
+ *      This program is free software: you can redistribute it and/or modify it under
+ *      the terms of the GNU General Public License as published by the Free Software
+ *      Foundation, either version 3 of the License, or (at your option) any later
+ *      version.
+ *
+ *      This program is distributed in the hope that it will be useful, but WITHOUT ANY
+ *      WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR A
+ *      PARTICULAR PURPOSE.  See the GNU General Public License for more details.
+ *
+ *      You should have received a copy of the GNU General Public License along with
+ *      this program.  If not, see <http://www.gnu.org/licenses/>.
+ *
+ * =======================================================================================
+ */
+#ifndef LIKWID_NVMON_PERFWORKS_H
+#define LIKWID_NVMON_PERFWORKS_H
+
+#if __CUDA_API_VERSION >= 10000
+
+static void *perfw_dl_libnvperf = NULL;
+static void *perfw_dl_libnvperf_t = NULL;
+static void *perfw_dl_libcupti = NULL;
+static void *perfw_dl_libcuda = NULL;
+#include <cupti_target.h>
+#include <cupti_profiler_target.h>
+#include <cuda_runtime_api.h>
+
+
+#include <nvperf_host.h>
+#include <nvperf_cuda_host.h>
+#include<nvperf_target.h>
+
+#define CUPTI_API_CALL(apiFuncCall)                                            \
+do {                                                                           \
+    CUptiResult _status = apiFuncCall;                                         \
+    if (_status != CUPTI_SUCCESS) {                                            \
+        fprintf(stderr, "%s:%d: error: function %s failed with error %d.\n",   \
+                __FILE__, __LINE__, #apiFuncCall, _status);                    \
+        exit(-1);                                                              \
+    }                                                                          \
+} while (0)
+
+
+
+#define NVPERFAPIWEAK __attribute__( ( weak ) )
+
+#define DECLARENVPERFFUNC(funcname, funcsig) NVPA_Status NVPERFAPIWEAK funcname funcsig;  NVPA_Status( *funcname##Ptr ) funcsig;
+//DECLARENVPERFFUNC(NVPW_MetricsContext_GetMetricNames_Begin, (NVPW_GetSupportedChipNames_Params *));
+
+DECLARENVPERFFUNC(NVPW_CUDA_MetricsContext_Create, (NVPW_CUDA_MetricsContext_Create_Params *));
+DECLARENVPERFFUNC(NVPW_MetricsContext_Destroy, (NVPW_MetricsContext_Destroy_Params*));
+DECLARENVPERFFUNC(NVPW_MetricsContext_GetMetricNames_Begin, (NVPW_MetricsContext_GetMetricNames_Begin_Params*));
+DECLARENVPERFFUNC(NVPW_MetricsContext_GetMetricNames_End, (NVPW_MetricsContext_GetMetricNames_End_Params*));
+DECLARENVPERFFUNC(NVPW_GetSupportedChipNames, (NVPW_GetSupportedChipNames_Params*));
+DECLARENVPERFFUNC(NVPW_Device_GetNames, (NVPW_Device_GetNames_Params*))
+
+#ifndef DECLARECUPTIFUNC
+#define CUPTIAPIWEAK __attribute__( ( weak ) )
+#define DECLARECUPTIFUNC(funcname, funcsig) CUptiResult CUPTIAPIWEAK funcname funcsig;  CUptiResult( *funcname##Ptr ) funcsig;
+#endif
+
+DECLARECUPTIFUNC(cuptiDeviceGetChipName, (CUpti_Device_GetChipName_Params*));
+
+#ifndef DECLARECUFUNC
+#define CUAPIWEAK __attribute__( ( weak ) )
+#define DECLARECUFUNC(funcname, funcsig) CUresult CUAPIWEAK funcname funcsig;  CUresult( *funcname##Ptr ) funcsig;
+#endif
+
+DECLARECUFUNC(cuInit, (unsigned int));
+
+#ifndef DLSYM_AND_CHECK
+#define DLSYM_AND_CHECK( dllib, name ) dlsym( dllib, name ); if ( dlerror() != NULL ) { return -1; }
+#endif
+
+static int
+link_perfworks_libraries(void)
+{
+    if(_dl_non_dynamic_init != NULL) {
+        return -1;
+    }
+    perfw_dl_libcuda = dlopen("libcuda.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!perfw_dl_libcuda)
+    {
+        fprintf(stderr, "CUDA library libcuda.so not found.\n");
+        return -1;
+    }
+    perfw_dl_libnvperf = dlopen("libnvperf_host.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!perfw_dl_libnvperf)
+    {
+        fprintf(stderr, "CUDA library libnvperf_host.so not found.\n");
+        return -1;
+    }
+    perfw_dl_libnvperf_t = dlopen("libnvperf_target.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!perfw_dl_libnvperf_t)
+    {
+        fprintf(stderr, "CUDA library libnvperf_target.so not found.\n");
+        return -1;
+    }
+    perfw_dl_libcupti = dlopen("libcupti.so", RTLD_NOW | RTLD_GLOBAL);
+    if (!perfw_dl_libcupti)
+    {
+        fprintf(stderr, "CUDA runtime library libcupti.so not found.\n");
+        return -1;
+    }
+    cuInitPtr = DLSYM_AND_CHECK(perfw_dl_libcuda, "cuInit");
+    cuptiDeviceGetChipNamePtr = DLSYM_AND_CHECK(perfw_dl_libcupti, "cuptiDeviceGetChipName");
+    NVPW_CUDA_MetricsContext_CreatePtr = DLSYM_AND_CHECK(perfw_dl_libnvperf, "NVPW_CUDA_MetricsContext_Create");
+    //NVPW_GetSupportedChipNamesPtr = DLSYM_AND_CHECK(perfw_dl_libnvperf, "NVPW_GetSupportedChipNames");
+    NVPW_Device_GetNamesPtr = DLSYM_AND_CHECK(perfw_dl_libnvperf_t, "NVPW_Device_GetNames");
+
+
+    CUresult cuErr = (*cuInitPtr)(0);
+    if (cuErr != CUDA_SUCCESS)
+    {
+        fprintf(stderr, "CUDA cannot be found and initialized (cuInit failed).\n");
+        return -ENODEV;
+    }
+    return 0;
+}
+
+static void
+release_perfworks_libraries(void)
+{
+    if (perfw_dl_libnvperf)
+    {
+        dlclose(perfw_dl_libnvperf);
+        perfw_dl_libnvperf = NULL;
+    }
+    if (perfw_dl_libcupti)
+    {
+        dlclose(perfw_dl_libcupti);
+        perfw_dl_libcupti = NULL;
+    }
+}
+
+int nvmon_perfworks_getEventsOfGpu(int gpuId, NvmonEventList_t* list)
+{
+    NVPA_Status err = 0;
+    int ret = topology_gpu_init();
+    if (ret != EXIT_SUCCESS)
+    {
+        return -ENODEV;
+    }
+    GpuTopology_t gtopo = get_gpuTopology();
+    ret = link_perfworks_libraries();
+    if (ret < 0)
+    {
+        return ret;
+    }
+
+
+    NVPW_Device_GetNames_Params getNamesParams = {NVPW_Device_GetNames_Params_STRUCT_SIZE};
+    err = (*NVPW_Device_GetNamesPtr)(&getNamesParams);
+    if (err != NVPA_STATUS_SUCCESS)
+    {
+        fprintf(stderr, "Cannot get name of GPU chip %d\n", err);
+        return -ENODEV;
+    }
+    printf("GPU name %s\n", getNamesParams.pChipName);
+    printf("GPU name %s\n", getNamesParams.pDeviceName);
+
+/*    */
+/*    CUpti_Device_GetChipName_Params getChipNameParams = { CUpti_Device_GetChipName_Params_STRUCT_SIZE };*/
+/*    getChipNameParams.pPriv = NULL;*/
+/*    getChipNameParams.deviceIndex = 0;*/
+/*    */
+/*    */
+/*    */
+/*    NVPW_CUDA_MetricsContext_Create_Params metricsContextCreateParams = { NVPW_CUDA_MetricsContext_Create_Params_STRUCT_SIZE };*/
+/*    metricsContextCreateParams.pChipName = getChipNameParams.pChipName;*/
+/*    err = (*NVPW_CUDA_MetricsContext_CreatePtr)(&metricsContextCreateParams);*/
+/*    printf("Here\n");*/
+/*    if (err != NVPA_STATUS_SUCCESS)*/
+/*    {*/
+/*        printf("Failed to get event list of device %d\n", gpuId);*/
+/*        return -ENODEV;*/
+/*    }*/
+/*    NVPW_MetricsContext_Destroy_Params metricsContextDestroyParams = { NVPW_MetricsContext_Destroy_Params_STRUCT_SIZE };*/
+/*    metricsContextDestroyParams.pMetricsContext = metricsContextCreateParams.pMetricsContext;*/
+/*    err = myNVPW_MetricsContext_Destroy((NVPW_MetricsContext_Destroy_Params *)&metricsContextDestroyParams);*/
+/*    */
+/*    NVPW_MetricsContext_GetMetricNames_Begin_Params getMetricNameBeginParams = { NVPW_MetricsContext_GetMetricNames_Begin_Params_STRUCT_SIZE };*/
+/*    getMetricNameBeginParams.pMetricsContext = metricsContextCreateParams.pMetricsContext;*/
+/*    getMetricNameBeginParams.hidePeakSubMetrics = 0;*/
+/*    getMetricNameBeginParams.hidePerCycleSubMetrics = 0;*/
+/*    getMetricNameBeginParams.hidePctOfPeakSubMetrics = 0;*/
+/*    err = myNVPW_MetricsContext_GetMetricNames_Begin(&getMetricNameBeginParams);*/
+/*    if (err != NVPA_STATUS_SUCCESS)*/
+/*    {*/
+/*        printf("Failed to get event list of device %d\n", gpuId);*/
+/*        return -ENODEV;*/
+/*    }*/
+/*    NVPW_MetricsContext_GetMetricNames_End_Params getMetricNameEndParams = { NVPW_MetricsContext_GetMetricNames_End_Params_STRUCT_SIZE };*/
+/*    getMetricNameEndParams.pMetricsContext = metricsContextCreateParams.pMetricsContext;*/
+/*    err = myNVPW_MetricsContext_GetMetricNames_End((NVPW_MetricsContext_GetMetricNames_End_Params *)&getMetricNameEndParams);*/
+/*    */
+/*    */
+/*    NvmonEventList_t l = malloc(sizeof(NvmonEventList));*/
+/*    if (l)*/
+/*    {*/
+/*        */
+/*        l->events = malloc(l->numEvents*sizeof(NvmonEventListEntry));*/
+/*        if (l->events)*/
+/*        {*/
+/*            for (int i = 0; i < getMetricNameBeginParams.numMetrics; i++)*/
+/*            {*/
+/*                NvmonEventListEntry* out = &l->events[i];*/
+/*                const char* mname = getMetricNameBeginParams.ppMetricNames[i];*/
+/*                out->name = malloc(strlen(mname)+2);*/
+/*                if (out->name)*/
+/*                {*/
+/*                    ret = snprintf(out->name, strlen(mname)+1, "%s", mname);*/
+/*                    if (ret > 0)*/
+/*                    {*/
+/*                        out->name[ret] = '\0';*/
+/*                    }*/
+/*                }*/
+/*                out->limit = malloc(10*sizeof(char));*/
+/*                if (out->limit)*/
+/*                {*/
+/*                    ret = snprintf(out->limit, 9, "GPU");*/
+/*                    if (ret > 0) out->limit[ret] = '\0';*/
+/*                }*/
+/*            }*/
+/*            l->numEvents = getMetricNameBeginParams.numMetrics;*/
+/*            *list = l;*/
+/*            return 0;*/
+/*        }*/
+/*        free(l);*/
+/*    }*/
+/*    return -ENOMEM;*/
+    return 0;
+}
+
+
+
+NvmonFunctions nvmon_perfworks_functions = {
+    .freeDevice = NULL,
+    .createDevice = NULL,
+    .getEventList = nvmon_perfworks_getEventsOfGpu,
+    .addEvents = NULL,
+    .setupCounters = NULL,
+    .startCounters = NULL,
+    .readCounters = NULL,
+};
+#else
+NvmonFunctions nvmon_perfworks_functions = {
+    .freeDevice = NULL,
+    .createDevice = NULL,
+    .getEventList = NULL,
+    .addEvents = NULL,
+    .setupCounters = NULL,
+    .startCounters = NULL,
+    .readCounters = NULL,
+};
+
+#endif
+
+#endif /* LIKWID_NVMON_PERFWORKS_H */
diff --git a/src/includes/nvmon_types.h b/src/includes/nvmon_types.h
index 70dc60dff..0e967f902 100644
--- a/src/includes/nvmon_types.h
+++ b/src/includes/nvmon_types.h
@@ -37,6 +37,8 @@
 #include <inttypes.h>
 #include <perfgroup.h>
 
+#include <cupti.h>
+
 #define NVMON_DEFAULT_STR_LEN 1024
 
 typedef enum {
@@ -93,6 +95,7 @@ A NvmonEventSet holds a set of event and counter combinations and some global in
 \extends NvmonGroupSet
 */
 typedef struct {
+    int                   id; /*!< \brief ID of event set */
     int                   numberOfEvents; /*!< \brief Number of events in \a events */
     NvmonEvent_t*         nvEvents; /*!< \brief List of events with length numberOfEvents */
     NvmonEventResult*     results; /* \brief List of event results with length numberOfEvents */
@@ -102,6 +105,12 @@ typedef struct {
     double                runTime; /*!< \brief Sum of all time information in seconds that the group was running */
 } NvmonEventSet;
 
+
+typedef enum {
+    LIKWID_NVMON_CUPTI_BACKEND,
+    LIKWID_NVMON_PERFWORKS_BACKEND,
+} NvmonBackends;
+
 typedef struct {
     int deviceId;
     CUdevice cuDevice;
@@ -122,6 +131,7 @@ typedef struct {
     uint64_t timeStart;
     uint64_t timeRead;
     uint64_t timeStop;
+    NvmonBackends backend;
 } NvmonDevice;
 typedef NvmonDevice* NvmonDevice_t;
 
@@ -132,6 +142,17 @@ typedef NvmonDevice* NvmonDevice_t;
 /*typedef NvmonControl* NvmonControl_t;*/
 
 
+typedef struct {
+    int  (*getEventList)(int gpuId, NvmonEventList_t* list);
+    int (*createDevice)(int id, NvmonDevice *dev);
+    void (*freeDevice)(NvmonDevice* dev);
+    int (*addEvents)(NvmonDevice* device, const char* eventString);
+    int (*setupCounters)(NvmonDevice* device, NvmonEventSet* eventSet);
+    int (*startCounters)(NvmonDevice* device);
+    int (*stopCounters)(NvmonDevice* device);
+    int (*readCounters)(NvmonDevice* device);
+    int (*finalizeCounters)(NvmonDevice* device);
+} NvmonFunctions;
 
 
 /*! \brief Structure specifying all performance monitoring event groups
@@ -149,6 +170,9 @@ typedef struct {
     //NvmonEventSet*   groups; /*!< \brief List of eventSets */
     int              numberOfGPUs; /*!< \brief Amount of GPUs in \a gpus */
     NvmonDevice*     gpus; /*!< \brief List of GPUs */
+    int              numberOfBackends;
+    NvmonFunctions*  backends[3];
 } NvmonGroupSet;
 
+
 #endif
diff --git a/src/luawid.c b/src/luawid.c
index 71d43e8b4..bab79cad7 100644
--- a/src/luawid.c
+++ b/src/luawid.c
@@ -76,7 +76,6 @@ static int nvmon_initialized = 0;
 
 CpuInfo_t cpuinfo = NULL;
 CpuTopology_t cputopo = NULL;
-GpuTopology_t gputopo = NULL;
 NumaTopology_t numainfo = NULL;
 AffinityDomains_t affinity = NULL;
 PowerInfo_t power;
@@ -1453,40 +1452,6 @@ lua_likwid_sockstr_to_socklist(lua_State* L)
     return 2;
 }
 
-static int
-lua_likwid_gpustr_to_gpulist(lua_State* L)
-{
-    int ret = 0;
-    char* gpustr = (char *)luaL_checkstring(L, 1);
-    if (!gputopology_isInitialized)
-    {
-        topology_gpu_init();
-        gputopo = get_gpuTopology();
-        gputopology_isInitialized = 1;
-    }
-    int* gpulist = (int*) malloc(gputopo->numDevices * sizeof(int));
-    if (gpulist == NULL)
-    {
-        lua_pushstring(L,"Cannot allocate data for the GPU list");
-        lua_error(L);
-    }
-    ret = gpustr_to_gpulist(gpustr, gpulist, gputopo->numDevices);
-    if (ret <= 0)
-    {
-        lua_pushstring(L,"Cannot parse GPU string");
-        lua_error(L);
-    }
-    lua_pushnumber(L, ret);
-    lua_newtable(L);
-    for (int i=0;i<ret;i++)
-    {
-        lua_pushinteger(L, (lua_Integer)( i+1));
-        lua_pushinteger(L, (lua_Integer)( gpulist[i]));
-        lua_settable(L,-3);
-    }
-    free(gpulist);
-    return 2;
-}
 
 static int
 lua_likwid_putAffinityInfo(lua_State* L)
@@ -2224,7 +2189,9 @@ lua_likwid_setVerbosity(lua_State* L)
     luaL_argcheck(L, (verbosity >= 0 && verbosity <= DEBUGLEV_DEVELOP), -1,
                 "Verbosity must be between 0 (only errors) and 3 (developer)");
     perfmon_setVerbosity(verbosity);
+#ifdef LIKWID_WITH_NVMON
     nvmon_setVerbosity(verbosity);
+#endif /* LIKWID_WITH_NVMON */
     return 0;
 }
 
@@ -2817,7 +2784,9 @@ lua_likwid_setresuser(lua_State* L)
     return 1;
 }
 
+#ifdef LIKWID_WITH_NVMON
 
+GpuTopology_t gputopo = NULL;
 
 static int
 lua_likwid_getGpuTopology(lua_State* L)
@@ -2964,6 +2933,41 @@ lua_likwid_putGpuTopology(lua_State* L)
     return 0;
 }
 
+static int
+lua_likwid_gpustr_to_gpulist(lua_State* L)
+{
+    int ret = 0;
+    char* gpustr = (char *)luaL_checkstring(L, 1);
+    if (!gputopology_isInitialized)
+    {
+        topology_gpu_init();
+        gputopo = get_gpuTopology();
+        gputopology_isInitialized = 1;
+    }
+    int* gpulist = (int*) malloc(gputopo->numDevices * sizeof(int));
+    if (gpulist == NULL)
+    {
+        lua_pushstring(L,"Cannot allocate data for the GPU list");
+        lua_error(L);
+    }
+    ret = gpustr_to_gpulist(gpustr, gpulist, gputopo->numDevices);
+    if (ret <= 0)
+    {
+        lua_pushstring(L,"Cannot parse GPU string");
+        lua_error(L);
+    }
+    lua_pushnumber(L, ret);
+    lua_newtable(L);
+    for (int i=0;i<ret;i++)
+    {
+        lua_pushinteger(L, (lua_Integer)( i+1));
+        lua_pushinteger(L, (lua_Integer)( gpulist[i]));
+        lua_settable(L,-3);
+    }
+    free(gpulist);
+    return 2;
+}
+
 static int
 lua_likwid_getGpuEventsAndCounters(lua_State* L)
 {
@@ -3327,6 +3331,21 @@ lua_likwid_gpuFinalize(lua_State *L)
     return 0;
 }
 
+static int
+lua_likwid_gpuSupported(lua_State *L)
+{
+    lua_pushboolean(L, 1);
+    return 1;
+}
+#else
+static int
+lua_likwid_gpuSupported(lua_State *L)
+{
+    lua_pushboolean(L, 1);
+    return 0;
+}
+#endif /* LIKWID_WITH_NVMON */
+
 
 /* #####   FUNCTION DEFINITIONS  -  EXPORTED FUNCTIONS   ################## */
 
@@ -3474,7 +3493,9 @@ luaopen_liblikwid(lua_State* L){
     lua_register(L, "likwid_seteuid", lua_likwid_seteuid);
     lua_register(L, "likwid_setresuid", lua_likwid_setresuid);
     lua_register(L, "likwid_setresuser", lua_likwid_setresuser);
-    // gpu functions
+    // Nvidia GPU functions
+    lua_register(L, "likwid_gpuSupported", lua_likwid_gpuSupported);
+#ifdef LIKWID_WITH_NVMON
     lua_register(L, "likwid_getGpuTopology", lua_likwid_getGpuTopology);
     lua_register(L, "likwid_putGpuTopology", lua_likwid_putGpuTopology);
     lua_register(L, "likwid_getGpuEventsAndCounters", lua_likwid_getGpuEventsAndCounters);
@@ -3500,9 +3521,10 @@ luaopen_liblikwid(lua_State* L){
     lua_register(L, "likwid_gpuGetNameOfCounter",lua_likwid_gpuGetNameOfCounter);
     lua_register(L, "likwid_gpuGetNameOfMetric",lua_likwid_gpuGetNameOfMetric);
     lua_register(L, "likwid_gpuGetNameOfGroup",lua_likwid_gpuGetNameOfGroup);
+#endif /* LIKWID_WITH_NVMON */
 #ifdef __MIC__
     setuid(0);
     seteuid(0);
-#endif
+#endif /* __MIC__ */
     return 0;
 }
diff --git a/src/nvmon.c b/src/nvmon.c
index 1f5709382..498d92036 100644
--- a/src/nvmon.c
+++ b/src/nvmon.c
@@ -46,8 +46,8 @@
 #include <sys/types.h>
 
 #include <dlfcn.h>
-#include <cupti.h>
-#include <cuda_runtime_api.h>
+//#include <cupti.h>
+#include <cuda.h>
 
 
 #include <types.h>
@@ -64,496 +64,28 @@
 
 #include <likwid.h>
 
+
+
 static int nvmon_initialized = 0;
 NvmonGroupSet* nvGroupSet = NULL;
 int likwid_nvmon_verbosity = DEBUGLEV_ONLY_ERROR;
 
-LikwidNvResults* gMarkerResults = NULL;
-int gMarkerRegions = 0;
-
-
-//static NvmonControl_t likwid_nvmon_control = NULL;
-/* Copy from PAPI's cuda component (BSD License)
- * @author  Asim YarKhan yarkhan@icl.utk.edu (updated in 2017 to support CUDA metrics)
- * @author  Asim YarKhan yarkhan@icl.utk.edu (updated in 2015 for multiple CUDA contexts/devices)
- * @author  Heike Jagode (First version, in collaboration with Robert Dietrich, TU Dresden) jagode@icl.utk.edu
- */
-void (*_dl_non_dynamic_init) (void) __attribute__ ((weak));
-
-#define CU_CALL( call, handleerror )                                    \
-    do {                                                                \
-        CUresult _status = (call);                                      \
-        if (_status != CUDA_SUCCESS) {                                  \
-            fprintf(stderr, "Error: function %s failed with error %d.\n", #call, _status); \
-            handleerror;                                                \
-        }                                                               \
-    } while (0)
-
-#define CUPTI_CALL(call, handleerror)                                 \
-    do {                                                                \
-        CUptiResult _status = (call);                                   \
-        if (_status != CUPTI_SUCCESS) {                                 \
-            const char *errstr;                                         \
-            (*cuptiGetResultString)(_status, &errstr);               \
-            fprintf(stderr, "Error: function %s failed with error %s.\n", #call, errstr); \
-            handleerror;                                                \
-        }                                                               \
-    } while (0)
-
-#define CUDA_CALL( call, handleerror )                                \
-    do {                                                                \
-        cudaError_t _status = (call);                                   \
-        if (_status != cudaSuccess) {                                   \
-            fprintf(stderr, "Error: function %s failed with error %d.\n", #call, _status); \
-            handleerror;                                                \
-        }                                                               \
-    } while (0)
-
-#define CUAPIWEAK __attribute__( ( weak ) )
-#define DECLARECUFUNC(funcname, funcsig) CUresult CUAPIWEAK funcname funcsig;  CUresult( *funcname##Ptr ) funcsig;
-DECLARECUFUNC(cuCtxGetCurrent, (CUcontext *));
-DECLARECUFUNC(cuCtxSetCurrent, (CUcontext));
-DECLARECUFUNC(cuDeviceGet, (CUdevice *, int));
-DECLARECUFUNC(cuDeviceGetCount, (int *));
-DECLARECUFUNC(cuDeviceGetName, (char *, int, CUdevice));
-DECLARECUFUNC(cuInit, (unsigned int));
-DECLARECUFUNC(cuCtxPopCurrent, (CUcontext * pctx));
-DECLARECUFUNC(cuCtxPushCurrent, (CUcontext pctx));
-DECLARECUFUNC(cuCtxSynchronize, ());
-
-#define CUDAAPIWEAK __attribute__( ( weak ) )
-#define DECLARECUDAFUNC(funcname, funcsig) cudaError_t CUDAAPIWEAK funcname funcsig;  cudaError_t( *funcname##Ptr ) funcsig;
-DECLARECUDAFUNC(cudaGetDevice, (int *));
-DECLARECUDAFUNC(cudaSetDevice, (int));
-DECLARECUDAFUNC(cudaFree, (void *));
-
-#define CUPTIAPIWEAK __attribute__( ( weak ) )
-#define DECLARECUPTIFUNC(funcname, funcsig) CUptiResult CUPTIAPIWEAK funcname funcsig;  CUptiResult( *funcname##Ptr ) funcsig;
-DECLARECUPTIFUNC(cuptiEventGroupGetAttribute, (CUpti_EventGroup eventGroup, CUpti_EventGroupAttribute attrib, size_t * valueSize, void *value));
-DECLARECUPTIFUNC(cuptiDeviceGetEventDomainAttribute, (CUdevice device, CUpti_EventDomainID eventDomain, CUpti_EventDomainAttribute attrib, size_t * valueSize, void *value));
-DECLARECUPTIFUNC(cuptiEventGroupReadEvent, (CUpti_EventGroup eventGroup, CUpti_ReadEventFlags flags, CUpti_EventID event, size_t * eventValueBufferSizeBytes, uint64_t *        eventValueBuffer));
-DECLARECUPTIFUNC(cuptiEventGroupSetAttribute, (CUpti_EventGroup eventGroup, CUpti_EventGroupAttribute attrib, size_t valueSize, void *value));
-DECLARECUPTIFUNC(cuptiEventGroupSetDisable, (CUpti_EventGroupSet * eventGroupSet));
-DECLARECUPTIFUNC(cuptiEventGroupSetEnable, (CUpti_EventGroupSet * eventGroupSet));
-DECLARECUPTIFUNC(cuptiEventGroupSetsCreate, (CUcontext context, size_t eventIdArraySizeBytes, CUpti_EventID * eventIdArray, CUpti_EventGroupSets ** eventGroupPasses));
-DECLARECUPTIFUNC(cuptiEventGroupSetsDestroy, (CUpti_EventGroupSets * eventGroupSets));
-DECLARECUPTIFUNC(cuptiGetTimestamp, (uint64_t * timestamp));
-DECLARECUPTIFUNC(cuptiSetEventCollectionMode, (CUcontext context, CUpti_EventCollectionMode mode));
-DECLARECUPTIFUNC(cuptiDeviceEnumEventDomains, (CUdevice, size_t *, CUpti_EventDomainID *));
-DECLARECUPTIFUNC(cuptiDeviceGetNumEventDomains, (CUdevice, uint32_t *));
-DECLARECUPTIFUNC(cuptiEventDomainEnumEvents, (CUpti_EventDomainID, size_t *, CUpti_EventID *));
-DECLARECUPTIFUNC(cuptiEventDomainGetNumEvents, (CUpti_EventDomainID, uint32_t *));
-DECLARECUPTIFUNC(cuptiEventGetAttribute, (CUpti_EventID, CUpti_EventAttribute, size_t *, void *));
-DECLARECUPTIFUNC(cuptiGetResultString, (CUptiResult result, const char **str));
-
-
-static void *dl_libcuda = NULL;
-static void *dl_libcudart = NULL;
-static void *dl_libcupti = NULL;
-
-static int
-link_libraries(void)
-{
-#define DLSYM_AND_CHECK( dllib, name ) dlsym( dllib, name ); if ( dlerror() != NULL ) { return -1; }
-
-    /* Attempt to guess if we were statically linked to libc, if so bail */
-    if(_dl_non_dynamic_init != NULL) {
-        return -1;
-    }
-    /* Need to link in the cuda libraries, if not found disable the component */
-    dl_libcuda = dlopen("libcuda.so", RTLD_NOW | RTLD_GLOBAL);
-    if (!dl_libcuda)
-    {
-        fprintf(stderr, "CUDA library libcuda.so not found.");
-        return -1;
-    }
-    cuCtxGetCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxGetCurrent");
-    cuCtxSetCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxSetCurrent");
-    cuDeviceGetPtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGet");
-    cuDeviceGetCountPtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGetCount");
-    cuDeviceGetNamePtr = DLSYM_AND_CHECK(dl_libcuda, "cuDeviceGetName");
-    cuInitPtr = DLSYM_AND_CHECK(dl_libcuda, "cuInit");
-    cuCtxPopCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxPopCurrent");
-    cuCtxPushCurrentPtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxPushCurrent");
-    cuCtxSynchronizePtr = DLSYM_AND_CHECK(dl_libcuda, "cuCtxSynchronize");
-
-    dl_libcudart = dlopen("libcudart.so", RTLD_NOW | RTLD_GLOBAL | RTLD_NODELETE);
-    if (!dl_libcudart)
-    {
-        fprintf(stderr, "CUDA runtime library libcudart.so not found.");
-        return -1;
-    }
-    cudaGetDevicePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaGetDevice");
-    cudaSetDevicePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaSetDevice");
-    cudaFreePtr = DLSYM_AND_CHECK(dl_libcudart, "cudaFree");
-
-    dl_libcupti = dlopen("libcupti.so", RTLD_NOW | RTLD_GLOBAL);
-    if (!dl_libcupti)
-    {
-        fprintf(stderr, "CUDA runtime library libcupti.so not found.");
-        return -1;
-    }
-    cuptiDeviceGetEventDomainAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceGetEventDomainAttribute");
-    cuptiEventGroupGetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupGetAttribute");
-    cuptiEventGroupReadEventPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupReadEvent");
-    cuptiEventGroupSetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetAttribute");
-    cuptiEventGroupSetDisablePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetDisable");
-    cuptiEventGroupSetEnablePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetEnable");
-    cuptiEventGroupSetsCreatePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetsCreate");
-    cuptiEventGroupSetsDestroyPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGroupSetsDestroy");
-    cuptiGetTimestampPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiGetTimestamp");
-    cuptiSetEventCollectionModePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiSetEventCollectionMode");
-    cuptiDeviceEnumEventDomainsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceEnumEventDomains");
-    cuptiDeviceGetNumEventDomainsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiDeviceGetNumEventDomains");
-    cuptiEventDomainEnumEventsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventDomainEnumEvents");
-    cuptiEventDomainGetNumEventsPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventDomainGetNumEvents");
-    cuptiEventGetAttributePtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiEventGetAttribute");
-    cuptiGetResultStringPtr = DLSYM_AND_CHECK(dl_libcupti, "cuptiGetResultString");
-    return 0;
-}
+#include <nvmon_cupti.h>
+#include <nvmon_nvml.h>
+#include <nvmon_perfworks.h>
 
-static int
-init_cuda(void)
-{
-    CUresult cuErr = (*cuInitPtr)(0);
-    if (cuErr != CUDA_SUCCESS)
-    {
-        fprintf(stderr, "CUDA cannot be found and initialized (cuInit failed).\n");
-        return -ENODEV;
-    }
-    return 0;
-}
-
-static int
-get_numDevices(void)
-{
-    CUresult cuErr;
-    int count = 0;
-    cuErr = (*cuDeviceGetCountPtr)(&count);
-    if(cuErr == CUDA_ERROR_NOT_INITIALIZED)
-    {
-        int ret = init_cuda();
-        if (ret == 0)
-        {
-            cuErr = (*cuDeviceGetCountPtr)(&count);
-        }
-        else
-        {
-            return ret;
-        }
-    }
-    return count;
-}
-
-/*static int*/
-/*nvmon_create_control(void)*/
-/*{*/
-/*    if (!likwid_nvmon_control)*/
-/*    {*/
-/*        likwid_nvmon_control = malloc(sizeof(NvmonControl));*/
-/*        likwid_nvmon_control->numDevices = 0;*/
-/*        likwid_nvmon_control->devices = NULL;*/
-/*    }*/
-/*}*/
-
-static NvmonDevice_t
-free_device(NvmonDevice_t dev)
-{
-    int j = 0;
-    GHashTableIter iter;
-    char* name = NULL;
-    uint32_t *id = NULL;
-    NvmonEvent_t event = NULL;
-    if (dev)
-    {
-        if (dev->context)
-        {
-            cuCtxDestroy(dev->context);
-            dev->context = NULL;
-        }
-        if (dev->cuEventSets)
-        {
-            CUPTI_CALL((*cuptiEventGroupSetsDestroyPtr)(dev->cuEventSets), return NULL);
-            dev->cuEventSets = NULL;
-        }
-        // We just need to free the keys in the hash tables.
-/*        g_hash_table_iter_init(&iter, dev->eventHash);*/
-/*        while (g_hash_table_iter_next(&iter, (gpointer)name, (gpointer)&event))*/
-/*        {*/
-/*            free(name);*/
-/*        }*/
-/*        g_hash_table_iter_init(&iter, dev->evIdHash);*/
-/*        while (g_hash_table_iter_next(&iter, (gpointer)id, (gpointer)&event))*/
-/*        {*/
-/*            free(id);*/
-/*        }*/
-        // The event objects are freed here
-        if (dev->allevents)
-        {
-            for (j = 0; j < dev->numAllEvents; j++)
-                free(dev->allevents[j]);
-        }
-        if (dev->activeEvents)
-        {
-            free(dev->activeEvents);
-            dev->activeEvents = NULL;
-            dev->numActiveEvents = 0;
-        }
-/*        if (dev->activeCuGroups)*/
-/*        {*/
-/*            printf("%p\n", dev->activeCuGroups);*/
-/*            free(dev->activeCuGroups);*/
-/*            dev->activeCuGroups = NULL;*/
-/*            dev->numActiveCuGroups = 0;*/
-/*        }*/
-        if (dev->nvEventSets)
-        {
-            for (j = 0; j < dev->numNvEventSets; j++)
-            {
-                free(dev->nvEventSets[j].results);
-            }
-            free(dev->nvEventSets);
-            dev->nvEventSets = NULL;
-            dev->numNvEventSets = 0;
-        }
-        memset(dev, 0, sizeof(NvmonDevice));
-        //free(dev);
-    }
-    return NULL;
-}
-
-static void print_event(NvmonEvent_t event)
-{
-    printf("Event %s: %s\n", event->name, event->description);
-}
 
-static int
-create_device(int id, NvmonDevice *dev)
-{
-    int j = 0, k = 0, c = 0;
-    int numDomains = 0;
-    CUpti_EventDomainID* eventDomainIds = NULL;
-    int eventIdx = 0;
-    uint32_t totalEvents = 0;
-
-    // Allocate new device space
-/*    NvmonDevice_t dev = malloc(sizeof(NvmonDevice));*/
-/*    if (!dev) return free_device(dev);*/
-/*    memset(dev, 0, sizeof(NvmonDevice));*/
-    dev->deviceId = id;
-    dev->cuEventSets = NULL;
-    dev->context = NULL;
-    dev->activeEvents = NULL;
-    dev->numActiveEvents = 0;
-    dev->numNvEventSets = 0;
-    dev->nvEventSets = NULL;
-
-    // Assign device ID and get cuDevice from CUDA
-    dev->deviceId = id;
-    CU_CALL((*cuDeviceGetPtr)(&dev->cuDevice, id), return -1);
-    // Get the name of the device
-/*    CU_CALL((*cuDeviceGetNamePtr)(dev->name, NVMON_DEFAULT_STR_LEN-1, dev->cuDevice), return free_device(dev));*/
-/*    dev->name[NVMON_DEFAULT_STR_LEN-1] = '\0';*/
-    // Get the number of event domains of the device
-    CUPTI_CALL((*cuptiDeviceGetNumEventDomainsPtr)(dev->cuDevice, &numDomains), return -1);
-    GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Nvmon: Dev %d Domains %d, id, numDomains);
-
-    // Get the domain IDs for the device
-    size_t domainarraysize = numDomains * sizeof(CUpti_EventDomainID);
-    eventDomainIds = malloc(domainarraysize);
-    if (!eventDomainIds) return -ENOMEM;
-    CUPTI_CALL((*cuptiDeviceEnumEventDomainsPtr)(dev->cuDevice, &domainarraysize, eventDomainIds), return -1);
-
-
-    // Count the events in all domains to allocate the event list
-    dev->numAllEvents = 0;
-    for (j = 0; j < numDomains; j++)
-    {
-        uint32_t domainNumEvents = 0;
-        CUpti_EventDomainID domainID = eventDomainIds[j];
-        CUPTI_CALL((*cuptiEventDomainGetNumEventsPtr)(domainID, &domainNumEvents), return -1);
-        GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Nvmon: Dev %d Domain %d Events %d, id, j, domainNumEvents);
-        totalEvents += domainNumEvents;
-    }
-    // Now we now how many events are provided by the device, so allocate a big enough event list
-    dev->allevents = malloc(totalEvents * sizeof(NvmonEvent_t));
-    if (!dev->allevents)
-    {
-        free(eventDomainIds);
-        return -ENOMEM;
-    }
-
-    // We use hash tables for faster access
-    dev->eventHash = g_hash_table_new(g_str_hash, g_str_equal);
-    dev->evIdHash = g_hash_table_new(g_int64_hash, g_int64_equal);
-
-    for (j = 0; j < numDomains; j++)
-    {
-        uint32_t domainNumEvents = 0;
-        CUpti_EventDomainID domainID = eventDomainIds[j];
-        // How many events are provided by the domain
-        CUPTI_CALL((*cuptiEventDomainGetNumEventsPtr)(domainID, &domainNumEvents), return -1);
-        size_t tmpSize = domainNumEvents * sizeof(CUpti_EventID);
-        // Allocate space for all CUPTI event IDs in the domain
-        CUpti_EventID* cuEventIds = malloc(tmpSize);
-        // Get the CUPTI events
-        CUPTI_CALL((*cuptiEventDomainEnumEventsPtr)(domainID, &tmpSize, cuEventIds), return -1);
-        for (k = 0; k < domainNumEvents; k++)
-        {
-            CUpti_EventID eventId = cuEventIds[k];
-            // Name and description are limited in length
-            size_t tmpSizeBytesName = (NVMON_DEFAULT_STR_LEN-1) * sizeof(char);
-            size_t tmpSizeBytesDesc = (NVMON_DEFAULT_STR_LEN-1) * sizeof(char);
-
-            NvmonEvent_t event = malloc(sizeof(NvmonEvent));
-            if (!event)
-            {
-                free(cuEventIds);
-                free(dev->allevents);
-                free(eventDomainIds);
-                return -ENOMEM;
-            }
-            // Get event name and description
-            CUPTI_CALL((*cuptiEventGetAttributePtr)(eventId, CUPTI_EVENT_ATTR_NAME, &tmpSizeBytesName, event->name), return -1);
-            CUPTI_CALL((*cuptiEventGetAttributePtr)(eventId, CUPTI_EVENT_ATTR_LONG_DESCRIPTION, &tmpSizeBytesDesc, event->description), return -1);
-            event->name[tmpSizeBytesName/sizeof(char)] = '\0';
-            event->description[tmpSizeBytesDesc/sizeof(char)] = '\0';
-            // LIKWID events are all uppercase, so transform the event names
-            c = 0;
-            while (event->name[c] != '\0')
-            {
-                event->name[c] = toupper(event->name[c]);
-                c++;
-            }
-            // Save all gathered information in a NvmonEvent object
-
-            event->cuEventId = eventId;
-            event->eventId = dev->numAllEvents;
-            event->cuDomainId = domainID;
-            event->domainId = j;
-            event->type = NVMON_CUPTI_EVENT;
-            event->active = 0;
-            //GPUDEBUG_PRINT(DEBUGLEV_DETAIL, New Event %d CuEvent %d Domain %d CuDomain %d Name %s, event->eventId, (int)event->cuEventId, event->domainId, (int)event->cuDomainId, event->name);
-            // Add the object to the event list
-            dev->allevents[dev->numAllEvents] = event;
-            dev->numAllEvents++;
-            // Add the object to the hash tables
-            char* nameKey = g_strdup(event->name);
-            //printf("Add event %s\n", nameKey);
-            CUpti_EventID* idKey = malloc(sizeof(CUpti_EventID));
-            if (idKey)
-            {
-                *idKey = event->cuEventId;
-                // Key is event name
-                g_hash_table_insert(dev->eventHash, (gpointer)nameKey, (gpointer)event);
-                // Key is CUPTI event ID
-                g_hash_table_insert(dev->evIdHash, (gpointer)idKey, (gpointer)event);
-                //print_event(event);
-            }
-        }
-    }
-    return 0;
-}
-
-
-
-static void print_active_event(NvmonActiveEvent_t event)
-{
-    printf("Event %d\n", event->idxInSet);
-}
+LikwidNvResults* gMarkerResults = NULL;
+int gMarkerRegions = 0;
 
 
-int nvmon_getEventsOfGpu(int gpuId, NvmonEventList_t* list)
-{
-    int ret = 0;
-    NvmonDevice device;
-    ret = link_libraries();
-    if (ret < 0)
-    {
-        return -ENODEV;
-    }
-    int err = create_device(gpuId, &device);
-    if (!err)
-    {
-        NvmonEventList_t l = malloc(sizeof(NvmonEventList));
-        if (l)
-        {
-            l->events = malloc(sizeof(NvmonEventListEntry) * device.numAllEvents);
-            if (l->events)
-            {
-                for (int i = 0; i < device.numAllEvents; i++)
-                {
-                    NvmonEventListEntry* out = &l->events[i];
-                    NvmonEvent_t event = device.allevents[i];
-                    out->name = malloc(strlen(event->name)+2);
-                    if (out->name)
-                    {
-                        ret = snprintf(out->name, strlen(event->name)+1, "%s", event->name);
-                        if (ret > 0)
-                        {
-                            out->name[ret] = '\0';
-                        }
-                    }
-                    out->desc = malloc(strlen(event->description)+2);
-                    if (out->name)
-                    {
-                        ret = snprintf(out->desc, strlen(event->description)+1, "%s", event->description);
-                        if (ret > 0)
-                        {
-                            out->desc[ret] = '\0';
-                        }
-                    }
-                    out->limit = malloc(10*sizeof(char));
-                    if (out->limit)
-                    {
-                        switch (event->type)
-                        {
-                            case NVMON_CUPTI_EVENT:
-                                ret = snprintf(out->limit, 9, "GPU");
-                                if (ret > 0) out->limit[ret] = '\0';
-                                break;
-                        }
-                    }
-                }
-                l->numEvents = device.numAllEvents;
-                *list = l;
-            }
-            else
-            {
-                free(l);
-                free_device(&device);
-                return -ENOMEM;
-            }
-        }
-    }
-    else
-    {
-        ERROR_PRINT(No such device %d, gpuId);
-    }
-    return 0;
-}
 
-void nvmon_returnEventsOfGpu(NvmonEventList_t list)
-{
-    if (list)
-    {
-        if (list->numEvents > 0 && list->events)
-        {
-            for (int i = 0; i < list->numEvents; i++)
-            {
-                NvmonEventListEntry* out = &list->events[i];
-                free(out->name);
-                free(out->desc);
-                free(out->limit);
-            }
-        }
-    }
-}
 
 int
 nvmon_init(int nrGpus, const int* gpuIds)
 {
     int idx = 0;
+    int ret = 0;
     GpuTopology_t gtopo = NULL;
     if (nvmon_initialized == 1)
     {
@@ -576,11 +108,6 @@ nvmon_init(int nrGpus, const int* gpuIds)
     {
         return -EEXIST;
     }
-    int ret = link_libraries();
-    if (ret < 0)
-    {
-        return -ENODEV;
-    }
 
     ret = topology_gpu_init();
     if (ret != EXIT_SUCCESS)
@@ -608,6 +135,16 @@ nvmon_init(int nrGpus, const int* gpuIds)
     nvGroupSet->numberOfActiveGroups = 0;
     nvGroupSet->groups = NULL;
     nvGroupSet->activeGroup = -1;
+    nvGroupSet->numberOfBackends = 0;
+
+#ifdef LIKWID_NVMON_CUPTI_H
+    nvGroupSet->backends[LIKWID_NVMON_CUPTI_BACKEND] = &nvmon_cupti_functions;
+    nvGroupSet->numberOfBackends++;
+#endif
+#ifdef LIKWID_NVMON_PERFWORKS_H
+    nvGroupSet->backends[LIKWID_NVMON_PERFWORKS_BACKEND] = &nvmon_perfworks_functions;
+    nvGroupSet->numberOfBackends++;
+#endif
 
     for (int i = 0; i < nrGpus; i++)
     {
@@ -629,8 +166,34 @@ nvmon_init(int nrGpus, const int* gpuIds)
             nvGroupSet = NULL;
             return -ENOMEM;
         }
-        ret = create_device(gpuIds[i], &nvGroupSet->gpus[idx]);
-        idx++;
+        NvmonDevice_t device = &nvGroupSet->gpus[idx];
+        if (gtopo->devices[gpuIds[i]].ccapMajor < 7)
+        {
+            device->backend = LIKWID_NVMON_CUPTI_BACKEND;
+        }
+        else
+        {
+            ERROR_PRINT(NVIDIA PerfWorks API current not supported);
+            free(nvGroupSet->gpus);
+            free(nvGroupSet);
+            nvGroupSet = NULL;
+            return -ENOMEM;
+        }
+
+        NvmonFunctions* funcs = nvGroupSet->backends[device->backend];
+        if (funcs)
+        {
+            if (funcs->createDevice)
+                ret = funcs->createDevice(gpuIds[i], device);
+        }
+        if (ret < 0)
+        {
+            ERROR_PRINT(Cannot create device %d, gpuIds[i]);
+            free(nvGroupSet->gpus);
+            free(nvGroupSet);
+            nvGroupSet = NULL;
+            return -ENOMEM;
+        }
     }
 
     nvmon_initialized = 1;
@@ -646,7 +209,10 @@ nvmon_finalize(void)
     {
         for (int i = 0; i < nvGroupSet->numberOfGPUs; i++)
         {
-            free_device(&nvGroupSet->gpus[i]);
+            NvmonDevice_t device = &nvGroupSet->gpus[i];
+            NvmonFunctions* funcs = nvGroupSet->backends[device->backend];
+            if (funcs)
+                funcs->freeDevice(&nvGroupSet->gpus[i]);
         }
         free(nvGroupSet->gpus);
         free(nvGroupSet);
@@ -655,6 +221,104 @@ nvmon_finalize(void)
     }
 }
 
+static int concatNvmonEventLists(NvmonEventList_t base, NvmonEventList_t new)
+{
+    if ((!new) || (!base))
+    {
+        return -EINVAL;
+    }
+
+    int totalevents = base->numEvents + new->numEvents;
+
+    NvmonEventListEntry* tmp = realloc(base->events, totalevents*sizeof(NvmonEventListEntry));
+    if (!tmp)
+    {
+        return -ENOMEM;
+    }
+    base->events = tmp;
+
+    for (int i = 0; i < new->numEvents; i++)
+    {
+        int bidx = base->numEvents + i;
+        base->events[bidx].name = malloc((strlen(new->events[i].name)+1)*sizeof(char));
+        if (base->events[bidx].name)
+        {
+            strncpy(base->events[bidx].name, new->events[i].name, strlen(new->events[i].name));
+        }
+        base->events[bidx].desc = malloc((strlen(new->events[i].desc)+1)*sizeof(char));
+        if (base->events[bidx].desc)
+        {
+            strncpy(base->events[bidx].desc, new->events[i].desc, strlen(new->events[i].desc));
+        }
+        base->events[bidx].limit = malloc((strlen(new->events[i].limit)+1)*sizeof(char));
+        if (base->events[bidx].limit)
+        {
+            strncpy(base->events[bidx].limit, new->events[i].limit, strlen(new->events[i].limit));
+        }
+    }
+    return 0;
+}
+
+
+
+int
+nvmon_getEventsOfGpu(int gpuId, NvmonEventList_t* list)
+{
+    int err = 0;
+    int ret = topology_gpu_init();
+    if (ret != EXIT_SUCCESS)
+    {
+        return -ENODEV;
+    }
+    GpuTopology_t gtopo = get_gpuTopology();
+    printf("Here\n");
+    int available = -1;
+    for (int i = 0; i < gtopo->numDevices; i++)
+    {
+        if (gtopo->devices[i].devid == gpuId)
+        {
+            available = i;
+            break;
+        }
+    }
+    printf("Device %d\n", available);
+    if (available >= 0)
+    {
+        if (gtopo->devices[available].ccapMajor < 7)
+        {
+            err = nvmon_cupti_functions.getEventList(available, list);
+        }
+        else
+        {
+            ERROR_PRINT(NVIDIA PerfWorks API current not supported);
+            return -ENODEV;
+            //err = nvmon_perfworks_functions.getEventList(available, list);
+        }
+    }
+    return err;
+
+}
+
+void nvmon_returnEventsOfGpu(NvmonEventList_t list)
+{
+    if (list)
+    {
+        if (list->numEvents > 0 && list->events)
+        {
+            for (int i = 0; i < list->numEvents; i++)
+            {
+                NvmonEventListEntry* out = &list->events[i];
+                free(out->name);
+                free(out->desc);
+                free(out->limit);
+            }
+            free(list->events);
+            list->events = NULL;
+            list->numEvents = 0;
+        }
+    }
+}
+
 
 
 int
@@ -662,15 +326,15 @@ nvmon_addEventSet(const char* eventCString)
 {
     int i = 0, j = 0, k = 0, l = 0, m = 0, err = 0;
     int isPerfGroup = 0;
-    int curDeviceId = -1;
+
     int devId = 0;
     int configuredEvents = 0;
-    CUcontext curContext;
+
     NvmonDevice_t device = NULL;
     Configuration_t config = NULL;
     GroupInfo ginfo;
 
-    struct bstrList* eventtokens;
+
     if (!eventCString)
     {
         return -EINVAL;
@@ -734,137 +398,20 @@ nvmon_addEventSet(const char* eventCString)
     }
     bdestroy(eventBString);
     char * evstr = get_eventStr(&nvGroupSet->groups[nvGroupSet->numberOfGroups-1]);
-    eventBString = bfromcstr(evstr);
+/*    eventBString = bfromcstr(evstr);*/
     GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, EventStr %s, evstr);
-    eventtokens = bsplit(eventBString, ',');
-    bdestroy(eventBString);
-
-
-
-    CUDA_CALL((*cudaGetDevicePtr)(&curDeviceId), return -EFAULT);
-    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-
+/*    eventtokens = bsplit(eventBString, ',');*/
+/*    bdestroy(eventBString);*/
 
     for (devId = 0; devId < nvGroupSet->numberOfGPUs; devId++)
     {
+        int err = 0;
         device = &nvGroupSet->gpus[devId];
-        size_t sizeBytes = (eventtokens->qty) * sizeof(CUpti_EventID);
-
-        int popContext = 0;
-        CUpti_EventGroupSets * cuEventSets = NULL;
-
-        NvmonEventSet* tmpEventSet = realloc(device->nvEventSets, (device->numNvEventSets+1)*sizeof(NvmonEventSet));
-        if (!tmpEventSet)
+        NvmonFunctions* funcs = nvGroupSet->backends[device->backend];
+        if (funcs->addEvents)
         {
-            ERROR_PRINT(Cannot enlarge GPU %d eventSet list, device->deviceId);
-            err = -ENOMEM;
-            continue;
+            err = funcs->addEvents(device, evstr);
         }
-        device->nvEventSets = tmpEventSet;
-        NvmonEventSet* devEventSet = &device->nvEventSets[device->numNvEventSets];
-
-        devEventSet->nvEvents = (NvmonEvent_t*) malloc(eventtokens->qty * sizeof(NvmonEvent_t));
-        if (devEventSet->nvEvents == NULL)
-        {
-            ERROR_PRINT(Cannot allocate event list for group %d\n, groupSet->numberOfActiveGroups);
-            return -ENOMEM;
-        }
-        devEventSet->cuEventIDs = (CUpti_EventID*) malloc(eventtokens->qty * sizeof(CUpti_EventID));
-        if (devEventSet->cuEventIDs == NULL)
-        {
-            ERROR_PRINT(Cannot allocate event ID list for group %d\n, groupSet->numberOfActiveGroups);
-            free(devEventSet->nvEvents);
-            return -ENOMEM;
-        }
-        devEventSet->results = malloc(eventtokens->qty * sizeof(NvmonEventResult));
-        if (devEventSet->cuEventIDs == NULL)
-        {
-            ERROR_PRINT(Cannot allocate result list for group %d\n, groupSet->numberOfActiveGroups);
-            free(devEventSet->cuEventIDs);
-            free(devEventSet->nvEvents);
-            return -ENOMEM;
-        }
-        memset(devEventSet->results, 0, eventtokens->qty * sizeof(NvmonEventResult));
-        devEventSet->numberOfEvents = 0;
-
-        // If the device has no context, check whether the current context is
-        // used already by another device
-        if (!device->context)
-        {
-            int contextDevice = -1;
-            for (j = 0; j < nvGroupSet->numberOfGPUs; j++)
-            {
-                NvmonDevice_t dev = &nvGroupSet->gpus[j];
-                if (dev->context == curContext)
-                {
-                    contextDevice = j;
-                    break;
-                }
-            }
-            if (contextDevice < 0)
-            {
-                // Current context is _not_ used by another device, use it for the
-                // current device
-                device->context = curContext;
-            }
-            else
-            {
-                // Create a new context and assign it to the current event device
-                CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
-                CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-                CU_CALL((*cuCtxGetCurrentPtr)(&device->context), return -EFAULT);
-            }
-        }
-        else if (device->context != curContext)
-        {
-            // The device context is not the current one, so add it to the
-            // context stack
-            CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
-            popContext = 1;
-        }
-
-
-        CUPTI_CALL((*cuptiSetEventCollectionModePtr)(device->context, CUPTI_EVENT_COLLECTION_MODE_KERNEL), return -EFAULT);
-
-        for (i = 0; i < eventtokens->qty; i++)
-        {
-            struct bstrList* evset = bsplit(eventtokens->entry[i], ':');
-            if (evset->qty != 2)
-            {
-                ERROR_PRINT(NVMON: Event %s invalid: Format <event>:<gpucounter>, bdata(eventtokens->entry[i]));
-            }
-            if (blength(evset->entry[0]) == 0 || blength(evset->entry[1]) == 0)
-            {
-                ERROR_PRINT(NVMON: Event %s invalid: Format <event>:<gpucounter>, bdata(eventtokens->entry[i]));
-            }
-            NvmonEvent_t event = g_hash_table_lookup(device->eventHash, (gpointer)bdata(evset->entry[0]));
-            if (!event)
-            {
-                GPUDEBUG_PRINT(DEBUGLEV_INFO, NVMON: Event %s unknown. Skipping..., bdata(evset->entry[0]));
-                continue; //unknown event
-            }
-            else
-            {
-                devEventSet->cuEventIDs[devEventSet->numberOfEvents] = event->cuEventId;
-                devEventSet->nvEvents[devEventSet->numberOfEvents] = event;
-                devEventSet->numberOfEvents++;
-                size_t s = devEventSet->numberOfEvents*sizeof(CUpti_EventID);
-
-                CUPTI_CALL((*cuptiEventGroupSetsCreatePtr)(device->context, s, devEventSet->cuEventIDs, &cuEventSets), devEventSet->numberOfEvents--;);
-                if (cuEventSets->numSets > 1)
-                {
-                    ERROR_PRINT(Error adding event %s. Multiple measurement runs are required. skipping event ..., bdata(evset->entry[i]));
-                    continue;
-                }
-            }
-        }
-        // Create an eventset with the currently configured and current event
-        if(popContext)
-        {
-            CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
-        }
-        device->numNvEventSets++;
     }
 
     // Check whether group has any event in any device
@@ -872,195 +419,35 @@ nvmon_addEventSet(const char* eventCString)
     return (nvGroupSet->numberOfActiveGroups-1);
 }
 
+
 int nvmon_setupCounters(int gid)
 {
-    int j = 0, k = 0, l = 0, m = 0;
-    int devId = 0, eid = 0;
+    int i = 0;
     int oldDevId;
     CUcontext curContext;
-    NvmonDevice_t device = NULL;
+
     if ((!nvGroupSet) || (!nvmon_initialized) || (gid < 0))
     {
         return -EFAULT;
     }
 
-    // Currently we are on which device?
-    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-    // This is a workaround to (eventually create and) get the current context
-    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-
     if (gid >= 0 && gid < nvGroupSet->numberOfGroups)
     {
-        for (devId = 0; devId < nvGroupSet->numberOfGPUs; devId++)
+        for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
         {
-            int popContext = 0;
-            device = &nvGroupSet->gpus[devId];
-            NvmonEventSet* devEventSet = &device->nvEventSets[gid];
-            CUpti_EventGroupSets * cuEventSets = NULL;
-            if (devEventSet->numberOfEvents == 0)
-            {
-                GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, gid);
-                continue;
-            }
-
-            // If the device has no context, check whether the current context is
-            // used already by another device
-            if (!device->context)
-            {
-                int contextDevice = -1;
-                for (j = 0; j < nvGroupSet->numberOfGPUs; j++)
-                {
-                    NvmonDevice_t dev = &nvGroupSet->gpus[j];
-                    if (dev->context == curContext)
-                    {
-                        contextDevice = j;
-                        break;
-                    }
-                }
-                if (contextDevice < 0)
-                {
-                    // Current context is _not_ used by another device, use it for the
-                    // current device
-                    device->context = curContext;
-                    printf("Reuse context %ld for device %d\n", device->context, device->deviceId);
-                }
-                else
-                {
-                    // Create a new context and assign it to the current event device
-                    CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
-                    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-                    CU_CALL((*cuCtxGetCurrentPtr)(&device->context), return -EFAULT);
-                    printf("New context %ld for device %d\n", device->context, device->deviceId);
-                }
-            }
-            else if (device->context != curContext)
-            {
-                // The device context is not the current one, so add it to the
-                // context stack
-                CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
-                popContext = 1;
-            }
-
-
-            size_t grpEventIdsSize = devEventSet->numberOfEvents * sizeof(CUpti_EventID);
-            CUPTI_CALL((*cuptiEventGroupSetsCreatePtr)(device->context, grpEventIdsSize, devEventSet->cuEventIDs, &cuEventSets), m++;);
-
-            // Allocate temporary array to hold the group event IDs
-            CUpti_EventID *grpEventIds = malloc(grpEventIdsSize);
-            if (!grpEventIds)
-            {
-                return -ENOMEM;
-            }
-
-            // Delete current activeEvent list
-            if (device->activeEvents)
-            {
-                free(device->activeEvents);
-                device->activeEvents = NULL;
-                device->numActiveEvents = 0;
-            }
-            // Delete current activeCuGroups list
-            if (device->activeCuGroups)
-            {
-                free(device->activeCuGroups);
-                device->activeCuGroups = NULL;
-                device->numActiveCuGroups = 0;
-            }
-
-            // Create a new activeEvent list
-            device->activeEvents = malloc(devEventSet->numberOfEvents * sizeof(NvmonActiveEvent));
-            if (!device->activeEvents)
-            {
-                free(grpEventIds);
-                grpEventIds = NULL;
-                return -ENOMEM;
-            }
-            // Create a new activeCuGroups list
-            device->activeCuGroups = malloc(cuEventSets->numSets * sizeof(CUpti_EventGroupSet**));
-            if (!device->activeCuGroups)
-            {
-                free(device->activeEvents);
-                device->activeEvents = NULL;
-                free(grpEventIds);
-                grpEventIds = NULL;
-                return -ENOMEM;
-            }
-
-
+            NvmonDevice_t device = &nvGroupSet->gpus[i];
 
+            NvmonEventSet* devEventSet = &device->nvEventSets[gid];
 
-            // Run over eventset and store all information we need for start/stop/reads in NvmonActiveEvent_t
-            CUpti_EventGroup curGroup;
-            uint32_t curNumInstances = 0, curNumTotalInstances = 0;
-            CUpti_EventGroupSet *curEventGroupSet;
-            for (j = 0; j < cuEventSets->numSets; j++)
+            for (int j = 0; j < nvGroupSet->numberOfBackends; j++)
             {
-                size_t sizeofuint32t = sizeof(uint32_t);
-                uint32_t numGEvents = 0, numGInstances = 0, numTotalGInstances = 0;
-                CUpti_EventGroupSet* groupset = &cuEventSets->sets[j];
-
-                for (k = 0; k < groupset->numEventGroups; k++)
+                int err = 0;
+                NvmonFunctions* funcs = nvGroupSet->backends[j];
+                if (funcs->setupCounters)
                 {
-                    uint32_t one = 1;
-                    CUpti_EventGroup group = groupset->eventGroups[k];
-                    // Get the number of events in the group
-                    CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_NUM_EVENTS, &sizeofuint32t, &numGEvents), return -EFAULT);
-                    // Get the CUPTI event IDs in the group
-                    CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_EVENTS, &grpEventIdsSize, grpEventIds), return -EFAULT);
-                    // If we don't set this, each event has only a single instance but we want to measure all instances
-                    CUPTI_CALL((*cuptiEventGroupSetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_PROFILE_ALL_DOMAIN_INSTANCES, sizeof(uint32_t), &one), return -EFAULT);
-                    // Get instance count for a group
-                    CUPTI_CALL((*cuptiEventGroupGetAttributePtr)(group, CUPTI_EVENT_GROUP_ATTR_INSTANCE_COUNT, &sizeofuint32t, &numGInstances), return -EFAULT);
-
-                    for (l = 0; l < numGEvents; l++)
-                    {
-                        for (m = 0; m <  devEventSet->numberOfEvents; m++)
-                        {
-                            if (devEventSet->cuEventIDs[m] == grpEventIds[l])
-                            {
-                                CUpti_EventDomainID did = devEventSet->nvEvents[m]->cuDomainId;
-                                // Get total instance count for a group
-                                CUPTI_CALL((*cuptiDeviceGetEventDomainAttributePtr)(device->cuDevice, did, CUPTI_EVENT_DOMAIN_ATTR_TOTAL_INSTANCE_COUNT, &sizeofuint32t, &numTotalGInstances), return -EFAULT);
-                                device->activeEvents[m].eventId = devEventSet->nvEvents[m]->eventId;
-                                device->activeEvents[m].idxInSet = m;
-                                device->activeEvents[m].groupId = gid;
-                                device->activeEvents[m].cuEventId = devEventSet->nvEvents[m]->cuEventId;
-                                device->activeEvents[m].cuDomainId = did;
-                                device->activeEvents[m].numTotalInstances = numTotalGInstances;
-                                device->activeEvents[m].cuGroup = group;
-                                device->activeEvents[m].cuGroupSet = groupset;
-                                device->activeEvents[m].numInstances = numGInstances;
-                                device->activeEvents[m].deviceId = device->deviceId;
-
-                                int found = 0;
-                                for (int i = 0; i < device->numActiveCuGroups; i++)
-                                {
-                                    if (device->activeCuGroups[i] == groupset)
-                                    {
-                                        found = 1;
-                                    }
-                                }
-                                if (!found)
-                                {
-                                    printf("activeCuGroups %d %d = %ld\n", device->numActiveCuGroups, devId, groupset);
-                                    device->activeCuGroups[device->numActiveCuGroups] = groupset;
-                                    device->numActiveCuGroups++;
-                                }
-                                // Mark event as active. This is used to avoid measuring the same event on the same device twice
-                                devEventSet->nvEvents[m]->active = 1;
-                                GPUDEBUG_PRINT(DEBUGLEV_INFO, Setup event %s (%d) for GPU %d, devEventSet->nvEvents[m]->name, device->activeEvents[m].cuEventId, device->deviceId);
-                                device->numActiveEvents++;
-                            }
-                        }
-                    }
+                    err = funcs->setupCounters(device, devEventSet);
                 }
             }
-            free(grpEventIds);
-            if(popContext)
-            {
-                CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
-            }
         }
     }
     nvGroupSet->activeGroup = gid;
@@ -1072,81 +459,23 @@ int nvmon_setupCounters(int gid)
 int
 nvmon_startCounters(void)
 {
-    int i = 0, j = 0, k = 0;
-    int oldDevId = -1;
-    uint64_t timestamp = 0;
-    CUcontext curContext;
+    int i = 0;
+    int err = 0;
     if ((!nvGroupSet) || (!nvmon_initialized) || (nvGroupSet->activeGroup < 0))
     {
         return -EFAULT;
     }
 
-    // Currently we are on which device?
-    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-    // Take the timestamp, we assign it later for all devices
-    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
-    // This is a workaround to (eventually create and) get the current context
-    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-
     for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
     {
-        int popContext = 0;
-        uint32_t one = 1;
         NvmonDevice_t device = &nvGroupSet->gpus[i];
-        if (device->numActiveCuGroups == 0)
+        for (int j = 0; j < nvGroupSet->numberOfBackends; j++)
         {
-            GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
-            continue;
-        }
-        if (device->deviceId != oldDevId)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
-            CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
-            CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-        }
-        CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-        device->timeStart = timestamp;
-        device->timeRead = timestamp;
-
-        // Are we in the proper context?
-        if (device->context != curContext)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Push Context %ld -> %ld for device %d, curContext, device->context, device->deviceId);
-            CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
-            popContext = 1;
-        }
-        else
-        {
-            // Although we are already in the right context, we set it again
-            // to be sure
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Set Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxSetCurrentPtr)(device->context), return -EFAULT);
-        }
-
-        NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
-        for (j = 0; j < nvEventSet->numberOfEvents; j++)
-        {
-            NvmonEventResult* res = &nvEventSet->results[j];
-            res->startValue = 0.0;
-            res->stopValue = 0.0;
-            res->currentValue = 0.0;
-            res->fullValue = 0.0;
-            res->overflows = 0;
-        }
-
-        for (j = 0; j < device->numActiveCuGroups; j++)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Enable group %ld on Dev %d (%d), device->activeCuGroups[j], i, device->deviceId);
-            printf("Enable group %ld on Dev %d (%d)\n", device->activeCuGroups[j], i, device->deviceId);
-            CUPTI_CALL((*cuptiEventGroupSetEnablePtr)(device->activeCuGroups[j]), return -EFAULT);
-        }
-
-        // If we added the device context to the stack, pop it again
-        if (popContext)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
+            NvmonFunctions* funcs = nvGroupSet->backends[j];
+            if (funcs->startCounters)
+            {
+                err = funcs->startCounters(device);
+            }
         }
     }
 
@@ -1156,93 +485,23 @@ nvmon_startCounters(void)
 int
 nvmon_stopCounters(void)
 {
-    int i = 0, j = 0, k = 0;
-    int oldDevId = -1;
-    uint64_t timestamp = 0;
-    CUcontext curContext;
+    int i = 0;
+    int err = 0;
     if ((!nvGroupSet) || (!nvmon_initialized) || (nvGroupSet->activeGroup < 0))
     {
         return -EFAULT;
     }
 
-    // Currently we are on which device?
-    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-    // Take the timestamp, we assign it later for all devices
-    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
-
-
     for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
     {
-        int popContext = 0;
-        uint32_t one = 1;
-        int maxTotalInstances = 0;
-        size_t valuesSize = 0;
         NvmonDevice_t device = &nvGroupSet->gpus[i];
-        if (device->deviceId != oldDevId)
+        for (int j = 0; j < nvGroupSet->numberOfBackends; j++)
         {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
-            CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
-            CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-        }
-        CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-        NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
-        if (device->numActiveCuGroups == 0)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
-            continue;
-        }
-        device->timeStop = timestamp;
-
-        // Are we in the proper context?
-        if (device->context != curContext)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Push Context %ld -> %ld for device %d, curContext, device->context, device->deviceId);
-            CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
-            popContext = 1;
-        }
-        else
-        {
-            // Although we are already in the right context, we set it again
-            // to be sure
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Set Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxSetCurrentPtr)(device->context), return -EFAULT);
-        }
-
-        for (j = 0; j < device->numActiveEvents; j++)
-        {
-            maxTotalInstances = MAX(maxTotalInstances, device->activeEvents[j].numTotalInstances);
-        }
-        uint64_t *tmpValues = (uint64_t *) malloc(maxTotalInstances * sizeof(uint64_t));
-
-        for (j = 0; j < device->numActiveEvents; j++)
-        {
-            NvmonActiveEvent_t event = &device->activeEvents[j];
-            valuesSize = sizeof(uint64_t) * event->numTotalInstances;
-            memset(tmpValues, 0, valuesSize);
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Read Grp %ld Ev %ld for device %d, event->cuGroup, event->cuEventId, device->deviceId);
-            CUPTI_CALL((*cuptiEventGroupReadEventPtr)(event->cuGroup, CUPTI_EVENT_READ_FLAG_NONE, event->cuEventId, &valuesSize, tmpValues), return -EFAULT);
-            uint64_t valuesum = 0;
-            for (k = 0; k < event->numInstances; k++)
+            NvmonFunctions* funcs = nvGroupSet->backends[j];
+            if (funcs->stopCounters)
             {
-                valuesum += tmpValues[k];
+                err = funcs->stopCounters(device);
             }
-            NvmonEventResult* res = &nvEventSet->results[event->idxInSet];
-            res->stopValue = (double)valuesum;
-            res->lastValue = res->currentValue;
-            res->fullValue += res->stopValue - res->startValue;
-            res->lastValue += res->stopValue - res->currentValue;
-            res->currentValue = (double)valuesum;
-        }
-        for (j = 0; j < device->numActiveCuGroups; j++)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Disable group %ld, device->activeCuGroups[j]);
-            CUPTI_CALL((*cuptiEventGroupSetDisablePtr)(device->activeCuGroups[j]), return -EFAULT);
-        }
-        // If we added the device context to the stack, pop it again
-        if (popContext)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
         }
     }
 
@@ -1252,12 +511,8 @@ nvmon_stopCounters(void)
 int nvmon_readCounters(void)
 {
 
-    int i = 0, j = 0, k = 0, l = 0, m = 0;
-    int oldDevId = -1;
-    uint64_t timestamp = 0;
-    CUcontext curContext;
-    size_t sizeofuint32num = sizeof(uint32_t);
-    int maxTotalInstances = 0;
+    int i = 0;
+    int err = 0;
     if ((!nvGroupSet) || (!nvmon_initialized) || (nvGroupSet->activeGroup < 0))
     {
         return -EFAULT;
@@ -1266,97 +521,16 @@ int nvmon_readCounters(void)
     for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
     {
         NvmonDevice_t device = &nvGroupSet->gpus[i];
-        for (j = 0; j < device->numActiveEvents; j++)
-        {
-            maxTotalInstances = MAX(maxTotalInstances, device->activeEvents[j].numTotalInstances);
-        }
-    }
-    // In this array we collect the instance values of an events (summed up later)
-    size_t valuesSize = sizeof(uint64_t) * maxTotalInstances;
-    uint64_t *tmpValues = (uint64_t *) malloc(valuesSize);
-    if (!tmpValues)
-    {
-        ERROR_PRINT(Not enough memory to allocate space for instance values);
-        return -ENOMEM;
-    }
-
-    // Currently we are on which device?
-    CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-    // Take the timestamp, we assign it later for all devices
-    CUPTI_CALL((*cuptiGetTimestampPtr)(&timestamp), return -EFAULT);
-/*    // This is a workaround to (eventually create and) get the current context*/
-/*    CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);*/
-/*    CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);*/
-
-    for (i = 0; i < nvGroupSet->numberOfGPUs; i++)
-    {
-        int popContext = 0;
-        NvmonDevice_t device = &nvGroupSet->gpus[i];
-        if (device->numActiveEvents == 0)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DETAIL, Skipping GPU%d it has no events in group %d, device->deviceId, nvGroupSet->activeGroup);
-            continue;
-        }
-        if (device->deviceId != oldDevId)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Change GPU device %d -> %d, oldDevId, device->deviceId);
-            CUDA_CALL((*cudaSetDevicePtr)(device->deviceId), return -EFAULT);
-            CUDA_CALL((*cudaGetDevicePtr)(&oldDevId), return -EFAULT);
-        }
-        CUDA_CALL((*cudaFreePtr)(NULL), return -EFAULT);
-        CU_CALL((*cuCtxGetCurrentPtr)(&curContext), return -EFAULT);
-        device->timeRead = timestamp;
-
-        // Are we in the proper context?
-        if (curContext != device->context)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Push Context %ld -> %ld for device %d, curContext, device->context, device->deviceId);
-            CU_CALL((*cuCtxPushCurrentPtr)(device->context), return -EFAULT);
-            popContext = 1;
-        }
-        else
+        for (int j = 0; j < nvGroupSet->numberOfBackends; j++)
         {
-            // Although we are already in the right context, we set it again
-            // to be sure
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Set Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxSetCurrentPtr)(device->context), return -EFAULT);
-        }
-        // Synchronize devices. I'm not sure whether this is required as each
-        // device measures it's own events
-        CU_CALL((*cuCtxSynchronizePtr)(), return -EFAULT);
-        NvmonEventSet* nvEventSet = &device->nvEventSets[nvGroupSet->activeGroup];
-
-        for (j = 0; j < device->numActiveEvents; j++)
-        {
-            NvmonActiveEvent_t event = &device->activeEvents[j];
-            // Empty space for instance values
-            valuesSize = sizeof(uint64_t) * event->numTotalInstances;
-            memset(tmpValues, 0, valuesSize);
-            // Read all instance values
-            //printf("%d %d %d\n", event->cuGroup, event->cuEventId, valuesSize);
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Read Grp %ld Ev %ld for device %d, event->cuGroup, event->cuEventId, device->deviceId);
-            CUPTI_CALL((*cuptiEventGroupReadEventPtr)(event->cuGroup, CUPTI_EVENT_READ_FLAG_NONE, event->cuEventId, &valuesSize, tmpValues), return -EFAULT);
-            // Sum all instance values
-            uint64_t valuesum = 0;
-            for (k = 0; k < event->numInstances; k++)
+            NvmonFunctions* funcs = nvGroupSet->backends[j];
+            if (funcs->readCounters)
             {
-                valuesum += tmpValues[k];
+                err = funcs->readCounters(device);
             }
-
-            NvmonEventResult* res = &nvEventSet->results[event->idxInSet];
-            res->lastValue = res->currentValue;
-            res->currentValue = (double)valuesum;
-            res->fullValue += res->currentValue - res->startValue;
-            res->lastValue += res->currentValue - res->lastValue;
-        }
-        if (popContext)
-        {
-            GPUDEBUG_PRINT(DEBUGLEV_DEVELOP, Pop Context %ld for device %d, device->context, device->deviceId);
-            CU_CALL((*cuCtxPopCurrentPtr)(&device->context), return -EFAULT);
         }
 
     }
-    free(tmpValues);
     return 0;
 }
 
diff --git a/src/topology_gpu.c b/src/topology_gpu.c
index a474bbd25..b004e6836 100644
--- a/src/topology_gpu.c
+++ b/src/topology_gpu.c
@@ -41,7 +41,7 @@
 
 #include <cupti.h>
 #include <dlfcn.h>
-#include <cuda_runtime_api.h>
+#include <cuda.h>
 
 #include <error.h>
 #include <likwid.h>
@@ -69,6 +69,8 @@ static void *topo_dl_libcuda = NULL;
 static int topology_gpu_initialized = 0;
 GpuTopology gpuTopology = {0, NULL};
 
+#ifdef LIKWID_WITH_NVMON
+
 DECLARECUFUNC(cuDeviceGet, (CUdevice *, int));
 DECLARECUFUNC(cuDeviceGetCount, (int *));
 DECLARECUFUNC(cuDeviceGetName, (char *, int, CUdevice));
@@ -77,6 +79,7 @@ DECLARECUFUNC(cuDeviceComputeCapability, (int*, int*, CUdevice));
 DECLARECUFUNC(cuDeviceGetAttribute, (int*, CUdevice_attribute, CUdevice));
 DECLARECUFUNC(cuDeviceGetProperties, (CUdevprop* prop, CUdevice));
 DECLARECUFUNC(cuDeviceTotalMem, (size_t*, CUdevice));
+DECLARECUFUNC(cuDeviceTotalMem_v2, (size_t*, CUdevice));
 
 
 static int
@@ -103,6 +106,7 @@ topo_link_libraries(void)
     cuDeviceGetAttributePtr = DLSYM_AND_CHECK(topo_dl_libcuda, "cuDeviceGetAttribute");
     cuDeviceGetPropertiesPtr = DLSYM_AND_CHECK(topo_dl_libcuda, "cuDeviceGetProperties");
     cuDeviceTotalMemPtr = DLSYM_AND_CHECK(topo_dl_libcuda, "cuDeviceTotalMem");
+    cuDeviceTotalMem_v2Ptr = DLSYM_AND_CHECK(topo_dl_libcuda, "cuDeviceTotalMem_v2");
     return 0;
 }
 
@@ -180,13 +184,13 @@ topology_gpu_init()
     ret = topo_link_libraries();
     if (ret != 0)
     {
-        ERROR_PLAIN_PRINT(Cannot open CUPA library to fill GPU topology);
+        ERROR_PLAIN_PRINT(Cannot open CUDA library to fill GPU topology);
         return EXIT_FAILURE;
     }
     int num_devs = topo_get_numDevices();
     if (num_devs < 0)
     {
-        ERROR_PLAIN_PRINT(Cannot get number of devices from CUPA library);
+        ERROR_PLAIN_PRINT(Cannot get number of devices from CUDA library);
         return EXIT_FAILURE;
     }
     if (num_devs > 0)
@@ -200,6 +204,17 @@ topology_gpu_init()
         {
             CUdevice dev;
             CU_CALL((*cuDeviceGetPtr)(&dev, i), return topology_gpu_cleanup(i-1, -ENODEV));
+            size_t s = 0;
+#if __CUDA_API_VERSION >= 10000
+            CU_CALL((*cuDeviceTotalMem_v2Ptr)(&s, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
+            if (s == 0)
+            {
+                CU_CALL((*cuDeviceTotalMemPtr)(&s, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
+            }
+#else
+            CU_CALL((*cuDeviceTotalMemPtr)(&s, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
+#endif
+            gpuTopology.devices[i].mem = (unsigned long long)s;
             gpuTopology.devices[i].name = malloc(1024 * sizeof(char));
             if (!gpuTopology.devices[i].name)
             {
@@ -209,7 +224,7 @@ topology_gpu_init()
             CU_CALL((*cuDeviceGetNamePtr)(gpuTopology.devices[i].name, 1023, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
             gpuTopology.devices[i].name[1024] = '\0';
             gpuTopology.devices[i].devid = i;
-            CU_CALL((*cuDeviceTotalMemPtr)(&gpuTopology.devices[i].mem, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
+
             CU_CALL((*cuDeviceComputeCapabilityPtr)(&gpuTopology.devices[i].ccapMajor, &gpuTopology.devices[i].ccapMinor, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
             CUdevprop props;
             CU_CALL((*cuDeviceGetPropertiesPtr)(&props, dev), return topology_gpu_cleanup(i-1, -ENOMEM));
@@ -269,3 +284,5 @@ get_gpuTopology(void)
         return &gpuTopology;
     }
 }
+
+#endif /* LIKWID_WITH_NVMON */
diff --git a/test/Makefile b/test/Makefile
index a1860399b..10c398660 100644
--- a/test/Makefile
+++ b/test/Makefile
@@ -64,9 +64,10 @@ testTBBGCC:
 testTBBICC:
 	@if [ $(TBB_AVAILABLE) -ne 0 -a $(ICPC_AVAILABLE) -ne 0 ]; then icpc -O3 $(LIKWID_DEFINES) $(LIKWID_INCLUDES) -o $@ testTBB.cc -ltbb $(LIKWID_LIB) -llikwid; else echo "Either TBB or ICPC missing"; fi
 
-.PHONY: clean streamGCC streamICC streamGCC_C11 streamICC_C11 testmarker-cnt testmarker-omp testmarkerF90 test-mpi stream_cilk serial test-likwidAPI streamAPIGCC test-msr-access testTBBGCC testTBBICC jacobi-2D-5pt-icc jacobi-2D-5pt-gcc
-
-clean:
-	rm -f streamGCC streamICC streamGCC_C11 streamICC_C11 stream_cilk testmarker-cnt testmarkerF90 test-mpi testmarker-omp serial test-likwidAPI streamAPIGCC test-msr-access testTBBGCC testTBBICC jacobi-2D-5pt-icc jacobi-2D-5pt-gcc
+matmul_marker: cupti/matmul_marker.c cupti/matmul.cu
+	nvcc -O3 -I. $(LIKWID_INCLUDE) -I./cupti-DLIKWID_NVMON cupti/matmul.cu cupti/matmul_marker.c $(LIKWID_LIB) -o $@ -lm -llikwid
 
+.PHONY: clean streamGCC streamICC streamGCC_C11 streamICC_C11 testmarker-cnt testmarker-omp testmarkerF90 test-mpi stream_cilk serial test-likwidAPI streamAPIGCC test-msr-access testTBBGCC testTBBICC jacobi-2D-5pt-icc jacobi-2D-5pt-gcc matmul_marker matmul
 
+clean:
+	rm -f streamGCC streamICC streamGCC_C11 streamICC_C11 stream_cilk testmarker-cnt testmarkerF90 test-mpi testmarker-omp serial test-likwidAPI streamAPIGCC test-msr-access testTBBGCC testTBBICC jacobi-2D-5pt-icc jacobi-2D-5pt-gcc matmul_marker matmul
diff --git a/test/cupti/matmul.cu b/test/cupti/matmul.cu
new file mode 100644
index 000000000..f2bd8c9fa
--- /dev/null
+++ b/test/cupti/matmul.cu
@@ -0,0 +1,370 @@
+#include <stdio.h>
+#include "realtime.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+
+#ifdef SET_DOUBLE
+typedef double FLOAT_TYPE;
+#else
+typedef float FLOAT_TYPE;
+#endif
+#define BLOCK_SIZE   32
+#define GET_MATRIX(m, row, col) (m.elements[row*m.cols + col])
+// Versions of matrix computations
+enum EMultVersions {SIMPLE, SHARED};
+
+typedef struct{
+    int cols;
+    int rows;
+    FLOAT_TYPE *elements;
+} Matrix;
+
+__device__
+FLOAT_TYPE getElement(Matrix m, int row, int col)
+{
+    if (row < m.rows && col < m.cols)
+    {
+        return m.elements[m.cols*row + col];
+    }
+    else
+    {
+        return 0;
+    }
+}
+
+__device__
+void setElement(Matrix m, int row, int col, FLOAT_TYPE val)
+{
+    if (row < m.rows && col < m.cols)
+    {
+        m.elements[m.cols*row + col] = val;
+    }
+}
+
+__global__
+void sharedMult(Matrix a, Matrix b, Matrix res)
+{
+    // block ids
+    int brow = blockIdx.y;
+    int bcol = blockIdx.x;
+    // id of thread in his block
+    int trow = threadIdx.y;
+    int tcol = threadIdx.x;
+
+    FLOAT_TYPE sum = 0;
+    for (int m = 0; m < a.cols; m += BLOCK_SIZE)
+    {
+        __shared__ FLOAT_TYPE aTile[BLOCK_SIZE][BLOCK_SIZE];
+        __shared__ FLOAT_TYPE bTile[BLOCK_SIZE][BLOCK_SIZE];
+
+        aTile[trow][tcol] = getElement(a, brow*BLOCK_SIZE + trow, tcol + m); 
+        bTile[trow][tcol] = getElement(b, trow + m, bcol*BLOCK_SIZE + tcol);
+        __syncthreads();
+
+        for(int k = 0;  k < BLOCK_SIZE; ++k)
+        {
+            sum += aTile[trow][k] * bTile[k][tcol];
+        }
+        __syncthreads();
+    }
+    setElement(res, 
+               blockIdx.y * blockDim.y + threadIdx.y,
+               blockIdx.x * blockDim.x + threadIdx.x,
+               sum);
+}
+
+__global__
+void simpleMult(Matrix a, Matrix b, Matrix res)
+{
+    int row = blockIdx.y * blockDim.y + threadIdx.y;
+    int col = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= a.rows || col >= b.cols)
+    {
+        return;
+    }
+
+    FLOAT_TYPE sum = 0;
+    for (int k = 0; k < a.cols; ++k)
+    {
+        sum += getElement(a, row, k) * getElement(b,  k, col);
+    }
+    setElement(res, row, col, sum);
+}
+
+/// <summary>checkes the last cuda command. if it wasn't successful
+/// stops the program with error description</summary>
+/// <param name="line">optional paramer with line number(where 
+/// the function was called</param>
+inline void checkErr(int line = -1)
+{
+    cudaError_t err = cudaGetLastError();
+    if(cudaSuccess != err) 
+    {
+        if (line != -1)
+            fprintf(stderr, "line nr. %d:", line );
+        fprintf(stderr, "Cuda error: %s.\n", cudaGetErrorString( err) );
+        exit(-1);
+    }                         
+}
+
+/// <summary>computes depending on the size of the problem matrix </summary>
+/// <returns>min number of blocks necessary to solve the problem</returns>
+dim3 numberOfBlocks(dim3 threadsPerBlock, int rows, int cols)
+{
+    dim3 numberOfBlocks(1,1,1);
+    
+    if(cols%threadsPerBlock.x)
+    {
+        numberOfBlocks.x = cols/threadsPerBlock.x + 1;
+    }
+    else
+    {
+        numberOfBlocks.x = cols/threadsPerBlock.x;
+    }
+
+    if(rows%threadsPerBlock.y)
+    {
+        numberOfBlocks.y = rows/threadsPerBlock.y + 1;
+    }
+    else
+    {
+        numberOfBlocks.y = rows/threadsPerBlock.y;
+    }
+    return numberOfBlocks;
+}
+
+/// <summary>computes a * b on GPU and stores the result in res</summary>
+/// <param name="version">version of kernel </param>
+/// <param name="threadsPerBlock">size of block. Global size would be adjust 
+/// automatically</param>
+/// <param name="a">input Matrix a</param>
+/// <param name="b">input Matrix b</param>
+/// <param name="res">output Matrix</param>
+void matirxMultOnGPU(EMultVersions version,         // kernel version
+                     const dim3 &threadsPerBlock,   // blocksize
+                     const Matrix &a,              // Matrix A
+                     const Matrix &b,              // Matrix B
+                     Matrix &res)                  // Matrix res = A*B
+{
+    // compute the size of the result matrix
+    res.rows = a.rows;
+    res.cols = b.cols;
+    res.elements = (FLOAT_TYPE*)malloc(res.cols*res.rows*sizeof(FLOAT_TYPE));
+    // allocate memory on GPU for matrices A, B and Res
+    size_t sizeOfA = a.rows*a.cols * sizeof(FLOAT_TYPE);
+    Matrix deviceA;
+    deviceA.cols = a.cols;
+    deviceA.rows = a.rows;
+    cudaMalloc((void**)&deviceA.elements, sizeOfA);
+    checkErr(__LINE__);
+
+    size_t sizeOfB = b.rows*b.cols * sizeof(FLOAT_TYPE);
+    Matrix deviceB;
+    deviceB.cols = b.cols;
+    deviceB.rows = b.rows;
+    cudaMalloc((void**)&deviceB.elements, sizeOfB);
+    checkErr(__LINE__);
+    size_t sizeOfRes = res.rows*res.cols * sizeof(FLOAT_TYPE);
+    Matrix deviceRes;
+    deviceRes.cols = res.cols;
+    deviceRes.rows = res.rows;
+    cudaMalloc((void**)&deviceRes.elements, sizeOfRes);
+    checkErr(__LINE__);
+    // copy data on GPU
+    cudaMemcpy(deviceA.elements, a.elements, sizeOfA,
+        cudaMemcpyHostToDevice);
+    checkErr(__LINE__);
+    cudaMemcpy(deviceB.elements, b.elements, sizeOfB, 
+        cudaMemcpyHostToDevice);
+    checkErr(__LINE__);
+    dim3 blocksPerGrid =
+        numberOfBlocks(threadsPerBlock, res.rows, res.cols);
+    // data size in GB witch has to be written/read on/from GPU
+    double dataSize = (a.rows*a.cols+b.cols*b.rows + a.rows*b.cols) * 
+        sizeof(FLOAT_TYPE)/(1e9);
+    double numberOfOperations = b.cols*a.rows*(2*a.cols-1.0)/(1e9);
+
+    printf("Input information\n");
+    printf("Matrix A: %dx%d\n",a.cols, a.cols);
+    printf("Matrix B: %dx%d\n",b.rows, b.cols);
+    printf("Memory transport(GB):%f\n",dataSize);
+    printf("Number of operations(*1e-9): %f\n\n", numberOfOperations);
+
+    double s, e;
+    s = GetRealTime();
+    // size of sharedMemory array from shredMult kernel
+    switch(version)
+    {
+    case SIMPLE:
+        simpleMult<<<blocksPerGrid, threadsPerBlock>>>
+            (deviceA, deviceB, deviceRes);
+        break;
+    case SHARED:
+        sharedMult<<<blocksPerGrid, threadsPerBlock>>>
+            (deviceA, deviceB, deviceRes);
+        break;
+    }
+    // makes sure, that computation is already ready
+    cudaThreadSynchronize();
+    e = GetRealTime();
+    printf("time on GPU (without memory transport and init):%f\n", e - s);
+    printf("GFlops: %f\n\n", numberOfOperations / (e-s));
+    //copy computing result on host
+    cudaMemcpy(res.elements, deviceRes.elements, sizeOfRes, 
+        cudaMemcpyDeviceToHost);
+    checkErr(__LINE__);
+
+    cudaFree(deviceA.elements);
+    cudaFree(deviceB.elements);
+    cudaFree(deviceRes.elements);
+}
+
+///<summary>version consists of text description of EMultVersion</summary>
+///<return>SHARED or SIMPLE. If strVersion contains wrong description, then
+/// SIMPLE</return>
+EMultVersions getAlgoVersion(const char* strVersion)
+{
+    const char *shared = "SHARED";
+    EMultVersions version = SIMPLE;
+    if (strcmp(shared, strVersion) == 0)
+        version = SHARED;
+    return version;
+}
+
+/// <summary>reserves memory for m.cols*m.rows elements of Type FLOAT_TYPE
+/// fills m with random values from [0,..,1]</summary>
+/// <param name="m">pointer to matrix which should be inialized</param>
+/// <param name="rows">number of rows in m</param>
+/// <param name="cols">number of columns in m</param>
+void fill(Matrix* m, int rows, int cols)
+{
+    m->cols = cols;
+    m->rows = rows;
+
+    size_t numberOfElements = cols * rows;
+    m->elements = (FLOAT_TYPE*)malloc(numberOfElements * sizeof(FLOAT_TYPE));
+    for (size_t i = 0; i < numberOfElements; ++i)
+    {
+        m->elements[i] = ((FLOAT_TYPE) (rand() % 101))/100;
+    }
+}
+
+/// <summary>frees memory referenced by m.elements</summary>
+/// <param name="m">Matrix with allocated memory</param>
+void freeMatrix(Matrix m)
+{
+    free(m.elements);
+}
+
+/// <summary>compares res with a*b. prints first fail</summaru>
+/// <param name="a">matrix a</param>
+/// <param name="b">matrix b</param> 
+/// <param name="res">the matrix, which should be controled</param>
+bool compare(Matrix a, Matrix b, Matrix res)
+{
+    if (res.cols != b.cols || res.rows != a.rows)
+        return false;
+    for (int row = 0; row < a.rows; ++ row)
+    {
+        for (int col = 0; col < b.cols; ++ col)
+        {
+            FLOAT_TYPE sum = 0;
+            for (int k = 0; k < a.cols; ++k)
+            {
+                sum += GET_MATRIX(a, row, k) * GET_MATRIX(b, k, col);
+            }
+            if (fabs(sum-GET_MATRIX(res, row, col))/fabs(GET_MATRIX(res, row, col)) > 0.0001)
+            {
+                printf("from CPU: %.10f, from GPU %.10f\n", 
+                       sum, 
+                       GET_MATRIX(res, row, col));
+                return false;
+            }
+        }
+    }
+    return true;
+}
+
+/// <summary>usage executable MATRIXSIZE VERSION RESALTTEST</summary>
+int matmul(int deviceId, size_t size)
+{
+    // dummy function in order to create context
+    cudaFree(0);
+
+    EMultVersions version = SHARED;
+    Matrix a;
+    Matrix b;
+    int testResult = 0;
+    printf("start matrix mult computation\n");
+#ifdef SET_DOUBLE
+    printf("set double(FloatType = double) precision\n");
+#else
+    printf("set single(FloatType = float) precision with CUDA\n");
+#endif
+    // create test data
+//    if (argc == 1) // default behavior
+//    {
+        fill(&a, size, size);
+        fill(&b, size, size);
+//    }    
+//    else if (argc == 2) // choose dimension of the matrix
+//    {
+//        int matrixDim = atoi(argv[1]);
+//        fill(&a, matrixDim, matrixDim);
+//        fill(&b, matrixDim, matrixDim);
+//    }
+//    else if (argc == 3) // choose algorithm and dimension of the matrix
+//    {
+//        int matrixDim = atoi(argv[1]);
+//        fill(&a, matrixDim, matrixDim);
+//        fill(&b, matrixDim, matrixDim);
+//        version = getAlgoVersion(argv[2]);
+//    }
+//    else // choose algorithm and dimension of the matrix, and test behavior
+//    {
+//        int matrixDim = atoi(argv[1]);
+//        fill(&a, matrixDim, matrixDim);
+//        fill(&b, matrixDim, matrixDim);
+//        version = getAlgoVersion(argv[2]);
+//        testResult = atoi(argv[3]);
+//    }
+    cudaSetDevice(deviceId);
+    Matrix res = {0,0,0};
+    dim3 threadsPerBlock(BLOCK_SIZE, BLOCK_SIZE);
+
+    // variables for time measurement
+    double s, e;
+    s = GetRealTime();
+    switch (version)
+    {
+    case SIMPLE:
+        printf("simple version.\n");
+        matirxMultOnGPU(SIMPLE, threadsPerBlock, a, b, res);
+        break;
+    case SHARED:
+        printf("shared version.\n");
+        matirxMultOnGPU(SHARED, threadsPerBlock, a, b, res);
+        break;
+    }
+    e = GetRealTime();  
+    double numberOfOperations = b.cols*a.rows*(2*a.cols - 1.0)/(1e9);
+    printf("total time (GPU init, memory reservation and computation): %f\n", e - s);
+    printf("GFlops(GPU init, memory reservation and computation)): %f\n\n", numberOfOperations / (e-s));
+    if (testResult)
+    {
+        printf("start test: ");
+        printf((compare(a, b, res)?"passed\n":"fail\n"));
+    }
+    freeMatrix(a);
+    freeMatrix(b);
+    freeMatrix(res);
+    printf("end matrix mult computation\n\n");
+    return 0;
+}
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/test/cupti/matmul.h b/test/cupti/matmul.h
new file mode 100644
index 000000000..67432f3f8
--- /dev/null
+++ b/test/cupti/matmul.h
@@ -0,0 +1,17 @@
+#ifndef TEST_CUPTI_H
+#define TEST_CUPTI_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+
+int matmul(int deviceId, int size);
+
+
+#ifdef __cplusplus
+}
+#endif
+
+
+#endif
diff --git a/test/cupti/matmul_marker.c b/test/cupti/matmul_marker.c
new file mode 100644
index 000000000..ba066b215
--- /dev/null
+++ b/test/cupti/matmul_marker.c
@@ -0,0 +1,27 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include <unistd.h>
+#include <stdint.h>
+#include <matmul.h>
+
+#include <likwid-gpumarker.h>
+
+
+int main(int argc, char* argv[])
+{
+    int count = 0;
+    uint64_t *run1 = NULL;
+    uint64_t *run2 = NULL;
+    int err = 0;
+    int gid = 0;
+    LIKWID_GPUMARKER_INIT;
+
+    LIKWID_GPUMARKER_START("matmul");
+    matmul(0, 3200);
+    //matmul(2, 3200);
+    LIKWID_GPUMARKER_STOP("matmul");
+
+    LIKWID_GPUMARKER_CLOSE;
+
+    return 0;
+}
diff --git a/test/cupti/realtime.h b/test/cupti/realtime.h
new file mode 100644
index 000000000..26aea8377
--- /dev/null
+++ b/test/cupti/realtime.h
@@ -0,0 +1,74 @@
+#ifndef _REALTIME_H_
+#define _REALTIME_H_
+
+#if defined _OPENMP
+#include <omp.h>
+#elif defined MPI_INT
+#include <mpi.h>
+#elif defined WIN32|WIN64
+#include <Windows.h>
+#include <stdio.h>
+#include <stdlib.h>
+#define Li2Double(x) \
+    ((double)((x).HighPart) * 4.294967296E9 + (double)((x).LowPart))
+#else
+#include <sys/time.h>
+#include <time.h>
+#endif
+
+double GetRealTime (void)
+{
+/* is there a common MPI define if MPI is active? MPI does not work */
+/* crude fix: look for MPI_INT */
+#if defined _OPENMP
+    return omp_get_wtime();
+
+#elif defined MPI_INT
+    return MPI_Wtime();
+
+#elif defined WIN32
+    LARGE_INTEGER time, freq;
+
+    if (QueryPerformanceCounter(&time) == 0)
+    {
+        DWORD err = GetLastError();
+        LPVOID buf;
+        FormatMessage(FORMAT_MESSAGE_ALLOCATE_BUFFER|FORMAT_MESSAGE_FROM_SYSTEM,
+                      NULL,
+                      err,
+                      MAKELANGID(LANG_NEUTRAL, SUBLANG_DEFAULT),
+                      (LPTSTR) &buf,
+                      0, NULL);
+        printf("QueryPerformanceCounter() failed with error %d: %s \n",err, buf);
+        exit(1);
+    }
+
+    if (QueryPerformanceFrequency(&freq) == 0)
+    {
+        DWORD err = GetLastError();
+        LPVOID buf;
+        FormatMessage(FORMAT_MESSAGE_ALLOCATE_BUFFER|FORMAT_MESSAGE_FROM_SYSTEM,
+                      NULL,
+                      err,
+                      MAKELANGID(LANG_NEUTRAL, SUBLANG_DEFAULT),
+                      (LPTSTR) &buf,
+                      0, NULL);
+        printf("QueryPerformanceFrequency() failed with error %d: %s \n",err, buf);
+        exit(1);
+    }
+    
+    return Li2Double(time) / Li2Double(freq);
+
+#elif defined SUNOS
+    hrtime_t t;
+    t = gethrtime();
+    return (((double)t)*1.0E-9);
+
+#else
+    struct timeval tv;
+    gettimeofday(&tv, (struct timezone*)0);
+    return ((double)tv.tv_sec + (double)tv.tv_usec / 1000000.0 );
+
+#endif
+}
+#endif