diff --git a/openmp/libomptarget/test/lit.cfg b/openmp/libomptarget/test/lit.cfg --- a/openmp/libomptarget/test/lit.cfg +++ b/openmp/libomptarget/test/lit.cfg @@ -80,6 +80,9 @@ 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. diff --git a/openmp/libomptarget/test/lit.site.cfg.in b/openmp/libomptarget/test/lit.site.cfg.in --- a/openmp/libomptarget/test/lit.site.cfg.in +++ b/openmp/libomptarget/test/lit.site.cfg.in @@ -19,6 +19,7 @@ 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") diff --git a/openmp/libomptarget/test/ompt/callbacks.h b/openmp/libomptarget/test/ompt/callbacks.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/callbacks.h @@ -0,0 +1,129 @@ +#include +#include + +// Tool related code below +#include + +// 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(); +} diff --git a/openmp/libomptarget/test/ompt/register_both.h b/openmp/libomptarget/test/ompt/register_both.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_both.h @@ -0,0 +1,49 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_emi.h b/openmp/libomptarget/test/ompt/register_emi.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_emi.h @@ -0,0 +1,47 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_emi_map.h b/openmp/libomptarget/test/ompt/register_emi_map.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_emi_map.h @@ -0,0 +1,48 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_no_device_init.h b/openmp/libomptarget/test/ompt/register_no_device_init.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_no_device_init.h @@ -0,0 +1,47 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_non_emi.h b/openmp/libomptarget/test/ompt/register_non_emi.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_non_emi.h @@ -0,0 +1,47 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_non_emi_map.h b/openmp/libomptarget/test/ompt/register_non_emi_map.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_non_emi_map.h @@ -0,0 +1,48 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/register_wrong_return.h b/openmp/libomptarget/test/ompt/register_wrong_return.h new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/register_wrong_return.h @@ -0,0 +1,47 @@ +#include + +// 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 diff --git a/openmp/libomptarget/test/ompt/veccopy.c b/openmp/libomptarget/test/ompt/veccopy.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy.c @@ -0,0 +1,84 @@ +// 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 +#include + +#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: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 + +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_disallow_both.c b/openmp/libomptarget/test/ompt/veccopy_disallow_both.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_disallow_both.c @@ -0,0 +1,102 @@ +// 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 +#include + +#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: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_emi.c b/openmp/libomptarget/test/ompt/veccopy_emi.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_emi.c @@ -0,0 +1,102 @@ +// 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 +#include +#include + +#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: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1 +/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=0 +/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=0 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_emi_map.c b/openmp/libomptarget/test/ompt/veccopy_emi_map.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_emi_map.c @@ -0,0 +1,103 @@ +// 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 +#include +#include + +#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: 0: Could not register callback 'ompt_callback_target_map_emi' +/// CHECK: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=1 +/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Target EMI: kind=1 endpoint=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=1 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=2 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=2 +/// CHECK: Callback Submit EMI: endpoint=1 req_num_teams=0 +/// CHECK: Callback Submit EMI: endpoint=2 req_num_teams=0 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=3 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=1 optype=4 +/// CHECK: Callback DataOp EMI: endpoint=2 optype=4 +/// CHECK: Callback Target EMI: kind=1 endpoint=2 +/// CHECK: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_map.c b/openmp/libomptarget/test/ompt/veccopy_map.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_map.c @@ -0,0 +1,85 @@ +// 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 +#include + +#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: 0: Could not register callback 'ompt_callback_target_map' +/// CHECK: Callback Init: +/// CHECK: Callback Load: +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 + +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_no_device_init.c b/openmp/libomptarget/test/ompt/veccopy_no_device_init.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_no_device_init.c @@ -0,0 +1,85 @@ +// 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 +#include + +#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-NOT: Callback Init: +/// CHECK-NOT: Callback Load: +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 + +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK-NOT: Callback Fini: diff --git a/openmp/libomptarget/test/ompt/veccopy_wrong_return.c b/openmp/libomptarget/test/ompt/veccopy_wrong_return.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/ompt/veccopy_wrong_return.c @@ -0,0 +1,85 @@ +// 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 +#include + +#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-NOT: Callback Init: +/// CHECK-NOT: Callback Load: +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 + +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=1 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=2 +/// CHECK-NOT: Callback Submit: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] req_num_teams=0 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=3 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback DataOp: target_id=[[TARGET_ID:[0-9]+]] host_op_id=[[HOST_OP_ID:[0-9]+]] optype=4 +/// CHECK-NOT: Callback Target: target_id=[[TARGET_ID:[0-9]+]] kind=1 endpoint=2 +/// CHECK-NOT: Callback Fini: