OMPD 5 and limited OMPD support for cuda devices

CMakeLists.txt

@@ -18,6 +18,8 @@ if (OPENMP_STANDALONE_BUILD OR "${CMAKE_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_S
     "Enable -Werror flags to turn warnings into errors for supporting compilers.")
   set(OPENMP_LIBDIR_SUFFIX "" CACHE STRING
     "Suffix of lib installation directory, e.g. 64 => lib64")
+  # Do not use OPENMP_LIBDIR_SUFFIX directly, use OPENMP_INSTALL_LIBDIR.
+  set(OPENMP_INSTALL_LIBDIR "lib${OPENMP_LIBDIR_SUFFIX}")
 
   # Group test settings.
   set(OPENMP_TEST_C_COMPILER ${CMAKE_C_COMPILER} CACHE STRING
@@ -28,7 +30,7 @@ if (OPENMP_STANDALONE_BUILD OR "${CMAKE_SOURCE_DIR}" STREQUAL "${CMAKE_CURRENT_S
 else()
   set(OPENMP_ENABLE_WERROR ${LLVM_ENABLE_WERROR})
   # If building in tree, we honor the same install suffix LLVM uses.
-  set(OPENMP_LIBDIR_SUFFIX ${LLVM_LIBDIR_SUFFIX})
+  set(OPENMP_INSTALL_LIBDIR "lib${LLVM_LIBDIR_SUFFIX}")
 
   if (NOT MSVC)
     set(OPENMP_TEST_C_COMPILER ${LLVM_RUNTIME_OUTPUT_INTDIR}/clang)

README.rst

@@ -257,9 +257,11 @@ Options for ``libomptarget``
 Options for ``NVPTX device RTL``
 --------------------------------
 
-**LIBOMPTARGET_NVPTX_ENABLE_BCLIB** = ``OFF|ON``
+**LIBOMPTARGET_NVPTX_ENABLE_BCLIB** = ``ON|OFF``
   Enable CUDA LLVM bitcode offloading device RTL. This is used for link time
-  optimization of the OMP runtime and application code.
+  optimization of the OMP runtime and application code. This option is enabled
+  by default if the build system determines that `CMAKE_C_COMPILER` is able to
+  compile and link the library.
 
 **LIBOMPTARGET_NVPTX_CUDA_COMPILER** = ``""``
   Location of a CUDA compiler capable of emitting LLVM bitcode. Currently only

libompd/cuda_examples/test_target_generic.c

@@ -0,0 +1,33 @@
+// Testing generic mode of nvptx devRtl
+#include <stdio.h>
+
+#pragma omp declare target
+void test_breakpoint() {
+  asm("");
+}
+#pragma omp end declare target
+
+void vec_mult(int N)
+{
+  int i;
+  float p[N], v1[N], v2[N];
+  //init(v1, v2, N);
+  #pragma omp target map(v1, v2, p)
+  {
+    test_breakpoint();
+  #pragma omp parallel for 
+  for (i=0; i<N; i++)
+  {
+    test_breakpoint();
+    p[i] = v1[i] * v2[i];
+  }
+    test_breakpoint();
+  }
+//output(p, N);
+}
+int main() {
+  printf("calling vec_mul...\n");
+  vec_mult(64);
+  printf("done\n");
+  return 0;
+}

libompd/cuda_examples/test_target_multilevel.c

@@ -0,0 +1,43 @@
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp declare target
+void test_breakpoint() {
+  asm("");
+}
+#pragma omp end declare target
+
+void vec_mult(int N)
+{
+  int i;
+  float p[N], v1[N], v2[N];
+  omp_set_nested(1);
+  #pragma omp target map(v1, v2, p)
+  {
+    omp_set_nested(1);
+  #pragma omp parallel shared(v1, v2, p, N) num_threads(4)
+  {
+    printf("Outer region - thread ID: %d\n", omp_get_thread_num());
+    #pragma omp for
+    for (int i = 0; i < N; ++i)
+    {
+      float acc = 0;
+      #pragma omp parallel shared(v1, v2, p, N) num_threads(4)
+      #pragma omp for
+      for(int j = 0; j < N; ++j)
+      {
+        test_breakpoint();
+        p[i] += v1[i] + v2[i];
+      }
+    }
+  }
+    printf("End of target region\n");
+  }
+//output(p, N);
+}
+int main() {
+  printf("calling vec_mul...\n");
+  vec_mult(64);
+  printf("done\n");
+  return 0;
+}

libompd/cuda_examples/test_target_noparallel.c

@@ -0,0 +1,24 @@
+#include <stdio.h>
+
+#pragma omp declare target
+void test_breakpoint() {
+  asm("");
+}
+#pragma omp end declare target
+
+void vec_mult(int N)
+{
+  int i;
+  float p[N], v1[N], v2[N];
+  #pragma omp target map(v1, v2, p)
+  {
+    test_breakpoint();
+    p[0] = v[0] * v[0];
+  }
+}
+int main() {
+  printf("calling vec_mul...\n");
+  vec_mult(64);
+  printf("done\n");
+  return 0;
+}

libompd/cuda_examples/test_target_single.c

@@ -0,0 +1,29 @@
+#include <stdio.h>
+
+#pragma omp declare target
+float mult(float u, float v) {
+  return u * v;
+}
+#pragma omp end declare target
+
+void vec_mult(int N)
+{
+  int i;
+  float p[N], v1[N], v2[N];
+  //init(v1, v2, N);
+  #pragma omp target map(v1, v2, p)
+  {
+    #pragma omp parallel for
+    for (i=0; i<N; i++)
+    {
+      p[i] = mult(v1[i], v2[i]);
+    }
+  }
+//output(p, N);
+}
+int main() {
+  printf("calling vec_mul...\n");
+  vec_mult(64);
+  printf("done\n");
+  return 0;
+}

libompd/cuda_examples/test_target_spmd.c

@@ -0,0 +1,31 @@
+// Testing spmd mode
+#include <stdio.h>
+
+#pragma omp declare target
+void test_breakpoint() {
+  asm("");
+}
+#pragma omp end declare target
+
+void vec_mult(int N)
+{
+  int i;
+  float p[N], v1[N], v2[N];
+  //init(v1, v2, N);
+  #pragma omp target map(v1, v2, p)
+  {
+  #pragma omp parallel for
+  for (i=0; i<N; i++)
+  {
+    test_breakpoint();
+    p[i] = v1[i] * v2[i];
+  }
+  }
+//output(p, N);
+}
+int main() {
+  printf("calling vec_mul...\n");
+  vec_mult(2048);
+  printf("done\n");
+  return 0;
+}

libompd/cuda_examples/test_target_task.c

@@ -0,0 +1,33 @@
+#include <stdio.h>
+#include <stdint.h>
+
+#pragma omp declare target
+void task1() {
+  printf("Hello from Task 1\n");
+  uint32_t enter_frame = 0;
+  for(;1;) {
+  }
+}
+void task2() {
+  printf("Hello from Task 2\n");
+  for(;1;) {
+  }
+}
+#pragma omp end declare target
+
+int main() {
+  #pragma omp target
+  {
+  #pragma omp parallel num_threads(4)
+  {
+    #pragma omp single
+    {
+      #pragma omp task
+      task1();
+      #pragma omp task
+      task2();
+    }
+  }
+  }
+  return 0;
+}

libompd/gdb-wrapper/CMakeLists.txt

@@ -1,5 +1,9 @@
 project (odb)
 
+cmake_minimum_required(VERSION 2.8)
+
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${CMAKE_SOURCE_DIR}/cmake/Modules/")
+
 set (cppfiles
   InputOutputManager.cpp
   ChildProcess.cpp
@@ -30,14 +34,12 @@ add_executable (odb-bin ${cppfiles} odb.cpp)
 set_target_properties (odb-bin PROPERTIES OUTPUT_NAME odb)
 add_library (odb ${cppfiles})
 
-if (ODB_LINUX)
-target_link_libraries (odb-bin dl) 
-target_link_libraries (odb dl) 
-endif (ODB_LINUX)
+target_link_libraries (odb-bin dl)
+target_link_libraries (odb dl)
 
 include_directories (
 	${CMAKE_CURRENT_SOURCE_DIR}
-#	${CMAKE_CURRENT_SOURCE_DIR}/../src/
+	${CMAKE_CURRENT_SOURCE_DIR}/../src/
     ${CMAKE_BINARY_DIR}/include
 )
 

libompd/gdb-wrapper/Callbacks.cpp

@@ -37,17 +37,16 @@ void initializeCallbacks(const GdbProcessPtr &proc)
   gdb = proc;
 
   // Initialize static table
-  cb.dmemory_alloc  = CB_dmemory_alloc;
-  cb.dmemory_free   = CB_dmemory_free;
-  cb.print_string   = CB_print_string;
-  cb.get_thread_context_for_osthread = CB_thread_context;
-  cb.get_containing_process_context = CB_process_context;
-  cb.tsizeof_prim   = CB_tsizeof_prim;
-  cb.tsymbol_addr   = CB_tsymbol_addr;
-  cb.read_tmemory   = CB_read_tmemory;
-  cb.write_tmemory  = CB_write_tmemory;
-  cb.host_to_target  = CB_host_to_target;
-  cb.target_to_host  = CB_target_to_host;
+  cb.memory_alloc       = CB_dmemory_alloc;
+  cb.memory_free        = CB_dmemory_free;
+  cb.print_string       = CB_print_string;
+  cb.get_thread_context_for_thread_id = CB_thread_context;
+  cb.sizeof_types       = CB_tsizeof_prim;
+  cb.symbol_addr_lookup = CB_tsymbol_addr;
+  cb.read_memory        = CB_read_tmemory;
+  cb.write_memory       = CB_write_tmemory;
+  cb.host_to_device     = CB_host_to_target;
+  cb.device_to_host     = CB_target_to_host;
 }
 
 ompd_callbacks_t * getCallbacksTable()
@@ -78,14 +77,14 @@ ompd_rc_t CB_dmemory_free (
 
 ompd_rc_t CB_thread_context (
     ompd_address_space_context_t *context,
-    ompd_osthread_kind_t          kind,
+    ompd_thread_id_t             kind,
     ompd_size_t                   sizeof_osthread,
     const void*                   osthread,
     ompd_thread_context_t **tcontext
     )
 {
   ompd_rc_t ret = context ? ompd_rc_ok : ompd_rc_stale_handle;
-  if (kind == ompd_osthread_cudalogical) {
+  if (kind == OMPD_THREAD_ID_CUDALOGICAL) {
     *tcontext = ((OMPDContext*)context)->getContextForThread((CudaThread*)osthread);
   }
   else {
@@ -126,7 +125,7 @@ void init_sizes(){
 
 ompd_rc_t CB_tsizeof_prim(
     ompd_address_space_context_t *context,
-    ompd_target_type_sizes_t *sizes)
+    ompd_device_type_sizes_t *sizes)
 {
   ompd_rc_t ret = context ? ompd_rc_ok : ompd_rc_stale_handle;
   static int inited = 0;
@@ -135,7 +134,12 @@ ompd_rc_t CB_tsizeof_prim(
     inited=1;
     init_sizes();
   }
-  memcpy(sizes, prim_sizes, sizeof(prim_sizes[0])*ompd_type_max);
+  sizes->sizeof_char = prim_sizes[ompd_type_char];
+  sizes->sizeof_short = prim_sizes[ompd_type_short];
+  sizes->sizeof_int = prim_sizes[ompd_type_int];
+  sizes->sizeof_long = prim_sizes[ompd_type_long];
+  sizes->sizeof_long_long = prim_sizes[ompd_type_long_long];
+  sizes->sizeof_pointer = prim_sizes[ompd_type_pointer];  
 
   return ret;
 }
@@ -175,7 +179,7 @@ ompd_rc_t CB_tsymbol_addr(
   parser.matchAddressValue(gdb->readOutput().c_str(), addr);
 
   if (strlen(addr) > 0)
-    symbol_addr->address = (ompd_taddr_t) strtoull (addr, NULL, 0);
+    symbol_addr->address = (ompd_addr_t) strtoull (addr, NULL, 0);
   else if (strlen(addr) == 0)
     ret = ompd_rc_error;
 
@@ -267,7 +271,7 @@ ompd_rc_t CB_write_tmemory (
     ompd_address_space_context_t *context,
     ompd_thread_context_t *tcontext,
     ompd_address_t addr,
-    ompd_tword_t nbytes,
+    ompd_word_t nbytes,
     const void *buffer)
 {
   return ompd_rc_unsupported;
@@ -277,7 +281,7 @@ ompd_rc_t CB_read_tmemory (
     ompd_address_space_context_t *context,
     ompd_thread_context_t *tcontext,
     ompd_address_t addr,
-    ompd_tword_t nbytes,
+    ompd_word_t nbytes,
     void *buffer)
 {
   if (!context)

libompd/gdb-wrapper/Callbacks.h

@@ -48,7 +48,7 @@ ompd_rc_t CB_dmemory_free (
 
 ompd_rc_t CB_thread_context (
     ompd_address_space_context_t *context,
-    ompd_osthread_kind_t          kind,
+    ompd_thread_id_t          kind,
     ompd_size_t                   sizeof_osthread,
     const void*                   osthread,
     ompd_thread_context_t **tcontext);
@@ -59,7 +59,7 @@ ompd_rc_t CB_process_context (
 
 ompd_rc_t CB_tsizeof_prim (
     ompd_address_space_context_t *context,
-    ompd_target_type_sizes_t *sizes);
+    ompd_device_type_sizes_t *sizes);
 
 ompd_rc_t CB_tsymbol_addr (
     ompd_address_space_context_t *context,
@@ -71,15 +71,15 @@ ompd_rc_t CB_read_tmemory (
     ompd_address_space_context_t *context,
     ompd_thread_context_t *tcontext,
     const ompd_address_t addr,
-    ompd_tword_t nbytes,
+    ompd_word_t nbytes,
     void *buffer
     );
 
 ompd_rc_t CB_write_tmemory (
     ompd_address_space_context_t *context,
     ompd_thread_context_t *tcontext,
     const ompd_address_t addr,
-    ompd_tword_t nbytes,
+    ompd_word_t nbytes,
     const void *buffer
     );
 

libompd/gdb-wrapper/CudaGdb.h

@@ -13,6 +13,7 @@
 #include <vector>
 #include <map>
 #include "ompd.h"
+#include "../src/ompd-private.h"
 
 struct CudaThread {
   ompd_cudathread_coord_t coord;