[OpenMP] [OMPT] [8/8] Added lit tests for OMPT target callbacks

[dhruvachak]

Summary

Added a new target ompt mode that depends on libomptarget OMPT support.
Added tests that verify callbacks for target regions, kernel launch,
and data transfer operations. All of them should pass on amdgpu using
make check-libomptarget.

[dhruvachak]

Rebased

[mhalk]

Rebased

[dhruvachak]

@mhalk Can this be landed on top of trunk today? I would be interested in landing this one on top of what we have today so that we get some test coverage as part of "make check". You will need to update the tests to match the current code and update the tests as patches 4-7 land.

I see this patch still depends on D127367. Can you remove this dependency, update the CHECK lines if not done already, and verify?

Thanks.

[mhalk]

Removed dependency on previous patches.
Note: Tests will now check for failed callback registrations.

[mhalk]

Fixed update.

[dhruvachak]

@jhuber6 Is this patch ok to land? This will provide some test coverage to the existing OMPT target functionality, the device-independent part. As plugin-specific changes land, these tests will have to be updated.

[dhruvachak]

ping. @jhuber6 @dreachem

[jplehr]

Just to reiterate what I think this does is: It adds tests which currently ensure that the call backs can not be registered (as it should).
The tests would then be updated for the other patches, when these call backs can be registered, right?

[dhruvachak]

In D127372#4179027, @jplehr wrote:

Just to reiterate what I think this does is: It adds tests which currently ensure that the call backs can not be registered (as it should).
The tests would then be updated for the other patches, when these call backs can be registered, right?

That's right. So landing this would test existing OMPT target functionality.

[jplehr]

LGTM

openmp/libomptarget/test/lit.cfg

 # Allow REQUIRES / UNSUPPORTED / XFAIL to work
 config.target_triple = [ ]
 for feature in config.test_compiler_features:
     config.available_features.add(feature)
 
 if config.libomptarget_debug:
   config.available_features.add('libomptarget-debug')
 
+if config.has_libomptarget_ompt:
+  config.available_features.add('ompt')
+
 config.available_features.add(config.libomptarget_current_target)
 
 # Determine whether the test system supports unified memory.
 # For CUDA, this is the case with compute capability 70 (Volta) or higher.
 # For all other targets, we currently assume it is.
 supports_unified_shared_memory = True
 if config.libomptarget_current_target.startswith('nvptx'):
   try:

openmp/libomptarget/test/lit.site.cfg.in

 config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
 config.llvm_lib_directory = "@LIBOMPTARGET_LLVM_LIBRARY_DIR@"
 config.operating_system = "@CMAKE_SYSTEM_NAME@"
 config.libomptarget_all_targets = "@LIBOMPTARGET_ALL_TARGETS@".split()
 config.libomptarget_current_target = "@CURRENT_TARGET@"
 config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
 config.libomptarget_not = "@OPENMP_NOT_EXECUTABLE@"
 config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
+config.has_libomptarget_ompt = @LIBOMPTARGET_OMPT_SUPPORT@
 
 # Let the main config do the real work.
 lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

openmp/libomptarget/test/ompt/callbacks.h

+#include <assert.h>
+#include <stdlib.h>
+
+// Tool related code below
+#include <omp-tools.h>
+
+// For EMI callbacks
+ompt_id_t next_op_id = 0x8000000000000001;
+
+// OMPT callbacks
+
+// Synchronous callbacks
+static void on_ompt_callback_device_initialize(int device_num, const char *type,
+                                               ompt_device_t *device,
+                                               ompt_function_lookup_t lookup,
+                                               const char *documentation) {
+  printf("Callback Init: device_num=%d type=%s device=%p lookup=%p doc=%p\n",
+         device_num, type, device, lookup, documentation);
+}
+
+static void on_ompt_callback_device_finalize(int device_num) {
+  printf("Callback Fini: device_num=%d\n", device_num);
+}
+
+static void on_ompt_callback_device_load(int device_num, const char *filename,
+                                         int64_t offset_in_file,
+                                         void *vma_in_file, size_t bytes,
+                                         void *host_addr, void *device_addr,
+                                         uint64_t module_id) {
+  printf("Callback Load: device_num:%d module_id:%lu filename:%s host_adddr:%p "
+         "device_addr:%p bytes:%lu\n",
+         device_num, module_id, filename, host_addr, device_addr, bytes);
+}
+
+static void on_ompt_callback_target_data_op(
+    ompt_id_t target_id, ompt_id_t host_op_id, ompt_target_data_op_t optype,
+    void *src_addr, int src_device_num, void *dest_addr, int dest_device_num,
+    size_t bytes, const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  // Both src and dest must not be null
+  assert((src_addr != 0 || dest_addr != 0) && "Both src and dest addr null");
+  printf("  Callback DataOp: target_id=%lu host_op_id=%lu optype=%d src=%p "
+         "src_device_num=%d "
+         "dest=%p dest_device_num=%d bytes=%lu code=%p\n",
+         target_id, host_op_id, optype, src_addr, src_device_num, dest_addr,
+         dest_device_num, bytes, codeptr_ra);
+}
+
+static void on_ompt_callback_target(ompt_target_t kind,
+                                    ompt_scope_endpoint_t endpoint,
+                                    int device_num, ompt_data_t *task_data,
+                                    ompt_id_t target_id,
+                                    const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  printf("Callback Target: target_id=%lu kind=%d endpoint=%d device_num=%d "
+         "code=%p\n",
+         target_id, kind, endpoint, device_num, codeptr_ra);
+}
+
+static void on_ompt_callback_target_submit(ompt_id_t target_id,
+                                           ompt_id_t host_op_id,
+                                           unsigned int requested_num_teams) {
+  printf("  Callback Submit: target_id=%lu host_op_id=%lu req_num_teams=%d\n",
+         target_id, host_op_id, requested_num_teams);
+}
+
+static void on_ompt_callback_target_map(ompt_id_t target_id,
+                                        unsigned int nitems, void **host_addr,
+                                        void **device_addr, size_t *bytes,
+                                        unsigned int *mapping_flags,
+                                        const void *codeptr_ra) {
+  printf("Target map callback is unimplemented\n");
+  abort();
+}
+
+static void on_ompt_callback_target_data_op_emi(
+    ompt_scope_endpoint_t endpoint, ompt_data_t *target_task_data,
+    ompt_data_t *target_data, ompt_id_t *host_op_id,
+    ompt_target_data_op_t optype, void *src_addr, int src_device_num,
+    void *dest_addr, int dest_device_num, size_t bytes,
+    const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  // Both src and dest must not be null
+  assert((src_addr != 0 || dest_addr != 0) && "Both src and dest addr null");
+  if (endpoint == ompt_scope_begin)
+    *host_op_id = next_op_id++;
+  printf("  Callback DataOp EMI: endpoint=%d optype=%d target_task_data=%p "
+         "(0x%lx) target_data=%p (0x%lx) host_op_id=%p (0x%lx) src=%p "
+         "src_device_num=%d "
+         "dest=%p dest_device_num=%d bytes=%lu code=%p\n",
+         endpoint, optype, target_task_data, target_task_data->value,
+         target_data, target_data->value, host_op_id, *host_op_id, src_addr,
+         src_device_num, dest_addr, dest_device_num, bytes, codeptr_ra);
+}
+
+static void on_ompt_callback_target_emi(ompt_target_t kind,
+                                        ompt_scope_endpoint_t endpoint,
+                                        int device_num, ompt_data_t *task_data,
+                                        ompt_data_t *target_task_data,
+                                        ompt_data_t *target_data,
+                                        const void *codeptr_ra) {
+  assert(codeptr_ra != 0 && "Unexpected null codeptr");
+  if (endpoint == ompt_scope_begin)
+    target_data->value = next_op_id++;
+  printf("Callback Target EMI: kind=%d endpoint=%d device_num=%d task_data=%p "
+         "(0x%lx) target_task_data=%p (0x%lx) target_data=%p (0x%lx) code=%p\n",
+         kind, endpoint, device_num, task_data, task_data->value,
+         target_task_data, target_task_data->value, target_data,
+         target_data->value, codeptr_ra);
+}
+
+static void on_ompt_callback_target_submit_emi(
+    ompt_scope_endpoint_t endpoint, ompt_data_t *target_data,
+    ompt_id_t *host_op_id, unsigned int requested_num_teams) {
+  printf("  Callback Submit EMI: endpoint=%d  req_num_teams=%d target_data=%p "
+         "(0x%lx) host_op_id=%p (0x%lx)\n",
+         endpoint, requested_num_teams, target_data, target_data->value,
+         host_op_id, *host_op_id);
+}
+
+static void on_ompt_callback_target_map_emi(ompt_data_t *target_data,
+                                            unsigned int nitems,
+                                            void **host_addr,
+                                            void **device_addr, size_t *bytes,
+                                            unsigned int *mapping_flags,
+                                            const void *codeptr_ra) {
+  printf("Target map emi callback is unimplemented\n");
+  abort();
+}

openmp/libomptarget/test/ompt/register_both.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_emi.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit_emi);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_emi_map.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op_emi);
+  register_ompt_callback(ompt_callback_target_emi);
+  register_ompt_callback(ompt_callback_target_submit_emi);
+  register_ompt_callback(ompt_callback_target_map_emi);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_no_device_init.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  // If no device init callback is registered, the other callbacks won't be
+  // activated.
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_non_emi.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_non_emi_map.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 0; // failed
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+  register_ompt_callback(ompt_callback_target_map);
+
+  return 1; // success
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/register_wrong_return.h

+#include <omp-tools.h>
+
+// From openmp/runtime/test/ompt/callback.h
+#define register_ompt_callback_t(name, type)                                   \
+  do {                                                                         \
+    type f_##name = &on_##name;                                                \
+    if (ompt_set_callback(name, (ompt_callback_t)f_##name) == ompt_set_never)  \
+      printf("0: Could not register callback '" #name "'\n");                  \
+  } while (0)
+
+#define register_ompt_callback(name) register_ompt_callback_t(name, name##_t)
+
+// OMPT entry point handles
+static ompt_set_callback_t ompt_set_callback = 0;
+
+// Init functions
+int ompt_initialize(ompt_function_lookup_t lookup, int initial_device_num,
+                    ompt_data_t *tool_data) {
+  ompt_set_callback = (ompt_set_callback_t)lookup("ompt_set_callback");
+
+  if (!ompt_set_callback)
+    return 1; // failed but wrongly returning 1
+
+  register_ompt_callback(ompt_callback_device_initialize);
+  register_ompt_callback(ompt_callback_device_finalize);
+  register_ompt_callback(ompt_callback_device_load);
+  register_ompt_callback(ompt_callback_target_data_op);
+  register_ompt_callback(ompt_callback_target);
+  register_ompt_callback(ompt_callback_target_submit);
+
+  return 0; // success but should return 1 according to the spec
+}
+
+void ompt_finalize(ompt_data_t *tool_data) {}
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+ompt_start_tool_result_t *ompt_start_tool(unsigned int omp_version,
+                                          const char *runtime_version) {
+  static ompt_start_tool_result_t ompt_start_tool_result = {&ompt_initialize,
+                                                            &ompt_finalize, 0};
+  return &ompt_start_tool_result;
+}
+#ifdef __cplusplus
+}
+#endif

openmp/libomptarget/test/ompt/veccopy.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that registers non-EMI callbacks
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_non_emi.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op'
+/// CHECK: Could not register callback 'ompt_callback_target'
+/// CHECK: Could not register callback 'ompt_callback_target_submit'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_disallow_both.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that shows that both EMI and non-EMI
+ * callbacks cannot be registered for the same type. In the
+ * current implementation, the EMI callback overrides the non-EMI
+ * callback.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_both.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op'
+/// CHECK: Could not register callback 'ompt_callback_target'
+/// CHECK: Could not register callback 'ompt_callback_target_submit'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_emi.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that registers EMI callbacks
+ */
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_emi.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op_emi'
+/// CHECK: Could not register callback 'ompt_callback_target_emi'
+/// CHECK: Could not register callback 'ompt_callback_target_submit_emi'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_emi_map.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that shows that map-EMI callbacks are not supported.
+ */
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_emi_map.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op_emi'
+/// CHECK: Could not register callback 'ompt_callback_target_emi'
+/// CHECK: Could not register callback 'ompt_callback_target_submit_emi'
+/// CHECK: Could not register callback 'ompt_callback_target_map_emi'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_map.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that shows that map callbacks are not supported.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_non_emi_map.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op'
+/// CHECK: Could not register callback 'ompt_callback_target'
+/// CHECK: Could not register callback 'ompt_callback_target_submit'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_no_device_init.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that shows that if no device init callback
+ * is registered, the other callbacks won't be activated.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_no_device_init.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op'
+/// CHECK: Could not register callback 'ompt_callback_target'
+/// CHECK: Could not register callback 'ompt_callback_target_submit'
+
+/// CHECK: Success

openmp/libomptarget/test/ompt/veccopy_wrong_return.c

+// RUN: %libomptarget-compile-run-and-check-generic
+// REQUIRES: ompt
+// UNSUPPORTED: nvptx64-nvidia-cuda
+// UNSUPPORTED: nvptx64-nvidia-cuda-oldDriver
+// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
+// UNSUPPORTED: x86_64-pc-linux-gnu
+// UNSUPPORTED: x86_64-pc-linux-gnu-oldDriver
+// UNSUPPORTED: x86_64-pc-linux-gnu-LTO
+
+/*
+ * Example OpenMP program that shows that if the initialize function
+ * returns the wrong status code, the callbacks won't be activated.
+ */
+
+#include <omp.h>
+#include <stdio.h>
+
+#include "callbacks.h"
+#include "register_wrong_return.h"
+
+int main() {
+  int N = 100000;
+
+  int a[N];
+  int b[N];
+
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = 0;
+
+  for (i = 0; i < N; i++)
+    b[i] = i;
+
+#pragma omp target parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+#pragma omp target teams distribute parallel for
+  {
+    for (int j = 0; j < N; j++)
+      a[j] = b[j];
+  }
+
+  int rc = 0;
+  for (i = 0; i < N; i++)
+    if (a[i] != b[i]) {
+      rc++;
+      printf("Wrong value: a[%d]=%d\n", i, a[i]);
+    }
+
+  if (!rc)
+    printf("Success\n");
+
+  return rc;
+}
+
+/// CHECK: Could not register callback 'ompt_callback_device_initialize'
+/// CHECK: Could not register callback 'ompt_callback_device_finalize'
+/// CHECK: Could not register callback 'ompt_callback_device_load'
+/// CHECK: Could not register callback 'ompt_callback_target_data_op'
+/// CHECK: Could not register callback 'ompt_callback_target'
+/// CHECK: Could not register callback 'ompt_callback_target_submit'
+
+/// CHECK: Success