[Offload] Move /openmp/libomptarget to /offload by jdoerfert · Pull Request #75125 · llvm/llvm-project (original) (raw)

⚠️ C/C++ code formatter, clang-format found issues in your code. ⚠️

You can test this locally with the following command:

git-clang-format --diff 446371f5ccbd81aa7707b4f1a1baa0ea0955dc67 c6a443db06e19ee114f74d847f52b1411df95833 -- offload/DeviceRTL/include/Allocator.h offload/DeviceRTL/include/Configuration.h offload/DeviceRTL/include/Debug.h offload/DeviceRTL/include/Interface.h offload/DeviceRTL/include/LibC.h offload/DeviceRTL/include/Mapping.h offload/DeviceRTL/include/State.h offload/DeviceRTL/include/Synchronization.h offload/DeviceRTL/include/Types.h offload/DeviceRTL/include/Utils.h offload/DeviceRTL/src/Allocator.cpp offload/DeviceRTL/src/Configuration.cpp offload/DeviceRTL/src/Debug.cpp offload/DeviceRTL/src/Kernel.cpp offload/DeviceRTL/src/LibC.cpp offload/DeviceRTL/src/Mapping.cpp offload/DeviceRTL/src/Misc.cpp offload/DeviceRTL/src/Parallelism.cpp offload/DeviceRTL/src/Reduction.cpp offload/DeviceRTL/src/State.cpp offload/DeviceRTL/src/Stub.cpp offload/DeviceRTL/src/Synchronization.cpp offload/DeviceRTL/src/Tasking.cpp offload/DeviceRTL/src/Utils.cpp offload/DeviceRTL/src/Workshare.cpp offload/include/DeviceImage.h offload/include/ExclusiveAccess.h offload/include/OffloadEntry.h offload/include/OffloadPolicy.h offload/include/OpenMP/InternalTypes.h offload/include/OpenMP/InteropAPI.h offload/include/OpenMP/Mapping.h offload/include/OpenMP/OMPT/Callback.h offload/include/OpenMP/OMPT/Connector.h offload/include/OpenMP/OMPT/Interface.h offload/include/OpenMP/omp.h offload/include/PluginManager.h offload/include/Shared/APITypes.h offload/include/Shared/Debug.h offload/include/Shared/Environment.h offload/include/Shared/EnvironmentVar.h offload/include/Shared/PluginAPI.h offload/include/Shared/PluginAPI.inc offload/include/Shared/Profile.h offload/include/Shared/Requirements.h offload/include/Shared/SourceInfo.h offload/include/Shared/Utils.h offload/include/Utils/ExponentialBackoff.h offload/include/device.h offload/include/omptarget.h offload/include/rtl.h offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h offload/plugins-nextgen/amdgpu/src/rtl.cpp offload/plugins-nextgen/amdgpu/utils/UtilitiesRTL.h offload/plugins-nextgen/common/OMPT/OmptCallback.cpp offload/plugins-nextgen/common/include/DLWrap.h offload/plugins-nextgen/common/include/GlobalHandler.h offload/plugins-nextgen/common/include/JIT.h offload/plugins-nextgen/common/include/MemoryManager.h offload/plugins-nextgen/common/include/PluginInterface.h offload/plugins-nextgen/common/include/RPC.h offload/plugins-nextgen/common/include/Utils/ELF.h offload/plugins-nextgen/common/src/GlobalHandler.cpp offload/plugins-nextgen/common/src/JIT.cpp offload/plugins-nextgen/common/src/PluginInterface.cpp offload/plugins-nextgen/common/src/RPC.cpp offload/plugins-nextgen/common/src/Utils/ELF.cpp offload/plugins-nextgen/cuda/dynamic_cuda/cuda.cpp offload/plugins-nextgen/cuda/dynamic_cuda/cuda.h offload/plugins-nextgen/cuda/src/rtl.cpp offload/plugins-nextgen/host/dynamic_ffi/ffi.cpp offload/plugins-nextgen/host/dynamic_ffi/ffi.h offload/plugins-nextgen/host/src/rtl.cpp offload/src/DeviceImage.cpp offload/src/LegacyAPI.cpp offload/src/OffloadRTL.cpp offload/src/OpenMP/API.cpp offload/src/OpenMP/InteropAPI.cpp offload/src/OpenMP/Mapping.cpp offload/src/OpenMP/OMPT/Callback.cpp offload/src/PluginManager.cpp offload/src/device.cpp offload/src/interface.cpp offload/src/omptarget.cpp offload/src/private.h offload/test/Inputs/declare_indirect_func.c offload/test/api/assert.c offload/test/api/is_initial_device.c offload/test/api/omp_device_managed_memory.c offload/test/api/omp_device_managed_memory_alloc.c offload/test/api/omp_device_memory.c offload/test/api/omp_dynamic_shared_memory.c offload/test/api/omp_dynamic_shared_memory_amdgpu.c offload/test/api/omp_dynamic_shared_memory_mixed.inc offload/test/api/omp_dynamic_shared_memory_mixed_amdgpu.c offload/test/api/omp_dynamic_shared_memory_mixed_nvptx.c offload/test/api/omp_env_vars.c offload/test/api/omp_get_device_num.c offload/test/api/omp_get_mapped_ptr.c offload/test/api/omp_get_num_devices.c offload/test/api/omp_get_num_devices_with_empty_target.c offload/test/api/omp_get_num_procs.c offload/test/api/omp_host_pinned_memory.c offload/test/api/omp_host_pinned_memory_alloc.c offload/test/api/omp_indirect_call.c offload/test/api/omp_target_memcpy_async1.c offload/test/api/omp_target_memcpy_async2.c offload/test/api/omp_target_memcpy_rect_async1.c offload/test/api/omp_target_memcpy_rect_async2.c offload/test/api/omp_target_memset.c offload/test/api/ompx_3d.c offload/test/api/ompx_3d.cpp offload/test/api/ompx_sync.c offload/test/api/ompx_sync.cpp offload/test/env/base_ptr_ref_count.c offload/test/env/omp_target_debug.c offload/test/jit/empty_kernel.inc offload/test/jit/empty_kernel_lvl1.c offload/test/jit/empty_kernel_lvl2.c offload/test/jit/type_punning.c offload/test/libc/assert.c offload/test/libc/fwrite.c offload/test/libc/global_ctor_dtor.cpp offload/test/libc/host_call.c offload/test/libc/malloc.c offload/test/libc/puts.c offload/test/mapping/alloc_fail.c offload/test/mapping/array_section_implicit_capture.c offload/test/mapping/array_section_use_device_ptr.c offload/test/mapping/auto_zero_copy.cpp offload/test/mapping/auto_zero_copy_apu.cpp offload/test/mapping/auto_zero_copy_globals.cpp offload/test/mapping/data_absent_at_exit.c offload/test/mapping/data_member_ref.cpp offload/test/mapping/declare_mapper_api.cpp offload/test/mapping/declare_mapper_nested_default_mappers.cpp offload/test/mapping/declare_mapper_nested_default_mappers_array.cpp offload/test/mapping/declare_mapper_nested_default_mappers_array_subscript.cpp offload/test/mapping/declare_mapper_nested_default_mappers_complex_structure.cpp offload/test/mapping/declare_mapper_nested_default_mappers_ptr_subscript.cpp offload/test/mapping/declare_mapper_nested_default_mappers_var.cpp offload/test/mapping/declare_mapper_nested_mappers.cpp offload/test/mapping/declare_mapper_target.cpp offload/test/mapping/declare_mapper_target_data.cpp offload/test/mapping/declare_mapper_target_data_enter_exit.cpp offload/test/mapping/declare_mapper_target_update.cpp offload/test/mapping/delete_inf_refcount.c offload/test/mapping/device_ptr_update.c offload/test/mapping/firstprivate_aligned.cpp offload/test/mapping/has_device_addr.cpp offload/test/mapping/implicit_device_ptr.c offload/test/mapping/is_device_ptr.cpp offload/test/mapping/lambda_by_value.cpp offload/test/mapping/lambda_mapping.cpp offload/test/mapping/low_alignment.c offload/test/mapping/map_back_race.cpp offload/test/mapping/ompx_hold/omp_target_disassociate_ptr.c offload/test/mapping/ompx_hold/struct.c offload/test/mapping/ompx_hold/target-data.c offload/test/mapping/ompx_hold/target.c offload/test/mapping/padding_not_mapped.c offload/test/mapping/power_of_two_alignment.c offload/test/mapping/pr38704.c offload/test/mapping/prelock.cpp offload/test/mapping/present/target.c offload/test/mapping/present/target_array_extension.c offload/test/mapping/present/target_data.c offload/test/mapping/present/target_data_array_extension.c offload/test/mapping/present/target_data_at_exit.c offload/test/mapping/present/target_enter_data.c offload/test/mapping/present/target_exit_data_delete.c offload/test/mapping/present/target_exit_data_release.c offload/test/mapping/present/target_update.c offload/test/mapping/present/target_update_array_extension.c offload/test/mapping/present/unified_shared_memory.c offload/test/mapping/present/zero_length_array_section.c offload/test/mapping/present/zero_length_array_section_exit.c offload/test/mapping/private_mapping.c offload/test/mapping/ptr_and_obj_motion.c offload/test/mapping/reduction_implicit_map.cpp offload/test/mapping/target_data_array_extension_at_exit.c offload/test/mapping/target_derefence_array_pointrs.cpp offload/test/mapping/target_has_device_addr.c offload/test/mapping/target_implicit_partial_map.c offload/test/mapping/target_map_for_member_data.cpp offload/test/mapping/target_pointers_members_map.cpp offload/test/mapping/target_update_array_extension.c offload/test/mapping/target_use_device_addr.c offload/test/mapping/target_uses_allocator.c offload/test/mapping/target_wrong_use_device_addr.c offload/test/offloading/assert.cpp offload/test/offloading/atomic-compare-signedness.c offload/test/offloading/back2back_distribute.c offload/test/offloading/barrier_fence.c offload/test/offloading/bug47654.cpp offload/test/offloading/bug49021.cpp offload/test/offloading/bug49334.cpp offload/test/offloading/bug49779.cpp offload/test/offloading/bug50022.cpp offload/test/offloading/bug51781.c offload/test/offloading/bug51982.c offload/test/offloading/bug53727.cpp offload/test/offloading/bug64959.c offload/test/offloading/bug64959_compile_only.c offload/test/offloading/bug74582.c offload/test/offloading/complex_reduction.cpp offload/test/offloading/ctor_dtor.cpp offload/test/offloading/cuda_no_devices.c offload/test/offloading/d2d_memcpy.c offload/test/offloading/d2d_memcpy_sync.c offload/test/offloading/default_thread_limit.c offload/test/offloading/dynamic_module.c offload/test/offloading/dynamic_module_load.c offload/test/offloading/extern.c offload/test/offloading/force-usm.cpp offload/test/offloading/fortran/basic_array.c offload/test/offloading/generic_multiple_parallel_regions.c offload/test/offloading/global_constructor.cpp offload/test/offloading/host_as_target.c offload/test/offloading/indirect_fp_mapping.c offload/test/offloading/info.c offload/test/offloading/interop.c offload/test/offloading/lone_target_exit_data.c offload/test/offloading/looptripcnt.c offload/test/offloading/malloc.c offload/test/offloading/malloc_parallel.c offload/test/offloading/mandatory_but_no_devices.c offload/test/offloading/memory_manager.cpp offload/test/offloading/multiple_reductions_simple.c offload/test/offloading/non_contiguous_update.cpp offload/test/offloading/offloading_success.c offload/test/offloading/offloading_success.cpp offload/test/offloading/ompx_bare.c offload/test/offloading/ompx_coords.c offload/test/offloading/ompx_saxpy_mixed.c offload/test/offloading/parallel_offloading_map.cpp offload/test/offloading/parallel_target_teams_reduction.cpp offload/test/offloading/parallel_target_teams_reduction_max.cpp offload/test/offloading/parallel_target_teams_reduction_min.cpp offload/test/offloading/requires.c offload/test/offloading/runtime_init.c offload/test/offloading/shared_lib_fp_mapping.c offload/test/offloading/small_trip_count.c offload/test/offloading/small_trip_count_thread_limit.cpp offload/test/offloading/spmdization.c offload/test/offloading/static_linking.c offload/test/offloading/std_complex_arithmetic.cpp offload/test/offloading/struct_mapping_with_pointers.cpp offload/test/offloading/target-teams-atomic.c offload/test/offloading/target-tile.c offload/test/offloading/target_constexpr_mapping.cpp offload/test/offloading/target_critical_region.cpp offload/test/offloading/target_depend_nowait.cpp offload/test/offloading/target_map_for_member_data.cpp offload/test/offloading/target_nowait_target.cpp offload/test/offloading/task_in_reduction_target.c offload/test/offloading/taskloop_offload_nowait.cpp offload/test/offloading/test_libc.cpp offload/test/offloading/thread_limit.c offload/test/offloading/thread_state_1.c offload/test/offloading/thread_state_2.c offload/test/offloading/weak.c offload/test/offloading/workshare_chunk.c offload/test/offloading/wtime.c offload/test/ompt/callbacks.h offload/test/ompt/register_both.h offload/test/ompt/register_emi.h offload/test/ompt/register_emi_map.h offload/test/ompt/register_no_device_init.h offload/test/ompt/register_non_emi.h offload/test/ompt/register_non_emi_map.h offload/test/ompt/register_wrong_return.h offload/test/ompt/target_memcpy.c offload/test/ompt/target_memcpy_emi.c offload/test/ompt/veccopy.c offload/test/ompt/veccopy_data.c offload/test/ompt/veccopy_disallow_both.c offload/test/ompt/veccopy_emi.c offload/test/ompt/veccopy_emi_map.c offload/test/ompt/veccopy_map.c offload/test/ompt/veccopy_no_device_init.c offload/test/ompt/veccopy_wrong_return.c offload/test/unified_shared_memory/api.c offload/test/unified_shared_memory/associate_ptr.c offload/test/unified_shared_memory/close_enter_exit.c offload/test/unified_shared_memory/close_manual.c offload/test/unified_shared_memory/close_member.c offload/test/unified_shared_memory/close_modifier.c offload/test/unified_shared_memory/shared_update.c offload/tools/deviceinfo/llvm-omp-device-info.cpp offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp offload/unittests/Plugins/NextgenPluginsTest.cpp

View the diff from clang-format here.

diff --git a/offload/DeviceRTL/include/Types.h b/offload/DeviceRTL/include/Types.h index 2e12d9da03..6d32f100c4 100644 --- a/offload/DeviceRTL/include/Types.h +++ b/offload/DeviceRTL/include/Types.h @@ -115,9 +115,9 @@ enum kmp_sched_t { #define SCHEDULE_WITHOUT_MODIFIERS(s)
(enum kmp_sched_t)(
(s) & ~(kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic)) -#define SCHEDULE_HAS_MONOTONIC(s) (((s)&kmp_sched_modifier_monotonic) != 0) +#define SCHEDULE_HAS_MONOTONIC(s) (((s) & kmp_sched_modifier_monotonic) != 0) #define SCHEDULE_HAS_NONMONOTONIC(s)
- (((s)&kmp_sched_modifier_nonmonotonic) != 0) + (((s) & kmp_sched_modifier_nonmonotonic) != 0) #define SCHEDULE_HAS_NO_MODIFIERS(s)
(((s) & (kmp_sched_modifier_nonmonotonic | kmp_sched_modifier_monotonic)) ==
0) diff --git a/offload/include/OpenMP/OMPT/Interface.h b/offload/include/OpenMP/OMPT/Interface.h index 327fadfcd4..4f624bd00b 100644 --- a/offload/include/OpenMP/OMPT/Interface.h +++ b/offload/include/OpenMP/OMPT/Interface.h @@ -254,8 +254,8 @@ private: // InterfaceRAII's class template argument deduction guide template <typename CallbackPairTy, typename... ArgsTy> -InterfaceRAII(CallbackPairTy Callbacks, ArgsTy... Args) - -> InterfaceRAII<CallbackPairTy, ArgsTy...>; +InterfaceRAII(CallbackPairTy Callbacks, + ArgsTy... Args) -> InterfaceRAII<CallbackPairTy, ArgsTy...>; /// Used to set and reset the thread-local return address. The RAII is expected /// to be created at a runtime entry point when the return address should be diff --git a/offload/include/OpenMP/omp.h b/offload/include/OpenMP/omp.h index b44c6aff1b..201ec5acaf 100644 --- a/offload/include/OpenMP/omp.h +++ b/offload/include/OpenMP/omp.h @@ -101,8 +101,9 @@ int __KAI_KMPC_CONVENTION omp_get_num_interop_properties(const omp_interop_t); * The omp_get_interop_int routine retrieves an integer property from an * omp_interop_t object. / -omp_intptr_t __KAI_KMPC_CONVENTION -omp_get_interop_int(const omp_interop_t, omp_interop_property_t, int ); +omp_intptr_t __KAI_KMPC_CONVENTION omp_get_interop_int(const omp_interop_t, + omp_interop_property_t, + int ); /! * The omp_get_interop_ptr routine retrieves a pointer property from an * omp_interop_t object. @@ -113,8 +114,9 @@ void __KAI_KMPC_CONVENTION omp_get_interop_ptr(const omp_interop_t, * The omp_get_interop_str routine retrieves a string property from an * omp_interop_t object. / -const char __KAI_KMPC_CONVENTION -omp_get_interop_str(const omp_interop_t, omp_interop_property_t, int ); +const char __KAI_KMPC_CONVENTION omp_get_interop_str(const omp_interop_t, + omp_interop_property_t, + int ); /! * The omp_get_interop_name routine retrieves a property name from an * omp_interop_t object. diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h index e8fc27785b..2ebd03af8e 100644 --- a/offload/include/Shared/APITypes.h +++ b/offload/include/Shared/APITypes.h @@ -26,11 +26,11 @@ extern "C" { /// This struct is a record of an entry point or global. For a function /// entry point the size is expected to be zero struct __tgt_offload_entry { - void addr; // Pointer to the offload entry info (function or global) - char name; // Name of the function or global - size_t size; // Size of the entry info (0 if it is a function) - int32_t flags; // Flags associated with the entry, e.g. 'link'. - int32_t data; // Extra data associated with the entry. + void addr; // Pointer to the offload entry info (function or global) + char name; // Name of the function or global + size_t size; // Size of the entry info (0 if it is a function) + int32_t flags; // Flags associated with the entry, e.g. 'link'. + int32_t data; // Extra data associated with the entry. }; /// This struct is a record of the device image information diff --git a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h index 3117763e35..82c607f474 100644 --- a/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h +++ b/offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h @@ -104,11 +104,11 @@ hsa_status_t hsa_amd_agents_allow_access(uint32_t num_agents, const uint32_t flags, const void ptr); -hsa_status_t hsa_amd_memory_lock(void host_ptr, size_t size, - hsa_agent_t agents, int num_agent, - void agent_ptr); +hsa_status_t hsa_amd_memory_lock(void host_ptr, size_t size, + hsa_agent_t agents, int num_agent, + void agent_ptr); -hsa_status_t hsa_amd_memory_unlock(void host_ptr); +hsa_status_t hsa_amd_memory_unlock(void host_ptr); hsa_status_t hsa_amd_memory_fill(void ptr, uint32_t value, size_t count); @@ -156,16 +156,15 @@ typedef enum { typedef struct hsa_amd_pointer_info_s { uint32_t size; hsa_amd_pointer_type_t type; - void agentBaseAddress; - void hostBaseAddress; + void agentBaseAddress; + void hostBaseAddress; size_t sizeInBytes; } hsa_amd_pointer_info_t; -hsa_status_t hsa_amd_pointer_info(const void ptr, - hsa_amd_pointer_info_t info, - void (alloc)(size_t), - uint32_t num_agents_accessible, - hsa_agent_t accessible); +hsa_status_t hsa_amd_pointer_info(const void ptr, hsa_amd_pointer_info_t info, + void (alloc)(size_t), + uint32_t num_agents_accessible, + hsa_agent_t *accessible); #ifdef __cplusplus } diff --git a/offload/plugins-nextgen/common/include/DLWrap.h b/offload/plugins-nextgen/common/include/DLWrap.h index 8934e7e701..00610678f9 100644 --- a/offload/plugins-nextgen/common/include/DLWrap.h +++ b/offload/plugins-nextgen/common/include/DLWrap.h @@ -107,7 +107,9 @@ template struct count { static constexpr size_t N = count<S - 1>::N; }; -template <> struct count<0> { static constexpr size_t N = 0; }; +template <> struct count<0> { + static constexpr size_t N = 0; +}; // Get a constexpr size_t ID, starts at zero #define DLWRAP_ID() (dlwrap::type::count<__LINE__>::N) diff --git a/offload/plugins-nextgen/cuda/dynamic_cuda/cuda.h b/offload/plugins-nextgen/cuda/dynamic_cuda/cuda.h index 32031c28f8..96c22dda1e 100644 --- a/offload/plugins-nextgen/cuda/dynamic_cuda/cuda.h +++ b/offload/plugins-nextgen/cuda/dynamic_cuda/cuda.h @@ -24,12 +24,12 @@ typedef struct CUfunc_st CUfunction; typedef struct CUstream_st CUstream; typedef struct CUevent_st CUevent; -#define CU_DEVICE_INVALID ((CUdevice)-2) +#define CU_DEVICE_INVALID ((CUdevice) - 2) typedef unsigned long long CUmemGenericAllocationHandle_v1; typedef CUmemGenericAllocationHandle_v1 CUmemGenericAllocationHandle; -#define CU_DEVICE_INVALID ((CUdevice)-2) +#define CU_DEVICE_INVALID ((CUdevice) - 2) typedef enum CUmemAllocationGranularity_flags_enum { CU_MEM_ALLOC_GRANULARITY_MINIMUM = 0x0, diff --git a/offload/src/OpenMP/API.cpp b/offload/src/OpenMP/API.cpp index c85f9868e3..77bdabb52f 100644 --- a/offload/src/OpenMP/API.cpp +++ b/offload/src/OpenMP/API.cpp @@ -644,8 +644,8 @@ EXTERN void omp_get_mapped_ptr(const void Ptr, int DeviceNum) { size_t NumDevices = omp_get_initial_device(); if (DeviceNum == NumDevices) { - DP("Device %d is initial device, returning Ptr " DPxMOD ".\n", - DeviceNum, DPxPTR(Ptr)); + DP("Device %d is initial device, returning Ptr " DPxMOD ".\n", DeviceNum, + DPxPTR(Ptr)); return const_cast<void *>(Ptr); } diff --git a/offload/test/api/omp_dynamic_shared_memory_amdgpu.c b/offload/test/api/omp_dynamic_shared_memory_amdgpu.c index 0b4d9d6ea9..c4d7c8da2e 100644 --- a/offload/test/api/omp_dynamic_shared_memory_amdgpu.c +++ b/offload/test/api/omp_dynamic_shared_memory_amdgpu.c @@ -1,4 +1,5 @@ -// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm -openmp-opt-inline-device +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm +// -openmp-opt-inline-device // RUN: env LIBOMPTARGET_SHARED_MEMORY_SIZE=256
// RUN: %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa // REQUIRES: amdgcn-amd-amdhsa diff --git a/offload/test/api/omp_dynamic_shared_memory_mixed.inc b/offload/test/api/omp_dynamic_shared_memory_mixed.inc index 16e0becd09..08e050b935 100644 --- a/offload/test/api/omp_dynamic_shared_memory_mixed.inc +++ b/offload/test/api/omp_dynamic_shared_memory_mixed.inc @@ -9,8 +9,7 @@ int main() { int Result[N], NumThreads; #pragma omp target teams num_teams(1) thread_limit(N)
- ompx_dyn_cgroup_mem(N * sizeof(Result[0]))
- map(from : Result, NumThreads) + ompx_dyn_cgroup_mem(N * sizeof(Result[0])) map(from : Result, NumThreads) { int Buffer[N]; #pragma omp parallel diff --git a/offload/test/api/omp_dynamic_shared_memory_mixed_amdgpu.c b/offload/test/api/omp_dynamic_shared_memory_mixed_amdgpu.c index 656c3a20aa..9de28974d6 100644 --- a/offload/test/api/omp_dynamic_shared_memory_mixed_amdgpu.c +++ b/offload/test/api/omp_dynamic_shared_memory_mixed_amdgpu.c @@ -1,4 +1,5 @@ -// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm -openmp-opt-inline-device -I %S +// RUN: %libomptarget-compile-amdgcn-amd-amdhsa -O1 -mllvm +// -openmp-opt-inline-device -I %S // RUN: env LIBOMPTARGET_NEXTGEN_PLUGINS=1
// RUN: %libomptarget-run-amdgcn-amd-amdhsa | %fcheck-amdgcn-amd-amdhsa // REQUIRES: amdgcn-amd-amdhsa diff --git a/offload/test/api/omp_get_mapped_ptr.c b/offload/test/api/omp_get_mapped_ptr.c index a8e11f912d..6d6fe42bc3 100644 --- a/offload/test/api/omp_get_mapped_ptr.c +++ b/offload/test/api/omp_get_mapped_ptr.c @@ -13,7 +13,7 @@ int main(int argc, char argv[]) { assert(device_ptr == NULL && "the pointer should not be mapped right now"); -#pragma omp target enter data map(to: host_data[:N]) +#pragma omp target enter data map(to : host_data[ : N]) device_ptr = omp_get_mapped_ptr(host_data, 0); @@ -21,7 +21,7 @@ int main(int argc, char argv[]) { void ptr = NULL; -#pragma omp target map(from: ptr) +#pragma omp target map(from : ptr) { ptr = host_data; } assert(ptr == device_ptr && "wrong pointer mapping"); @@ -30,7 +30,7 @@ int main(int argc, char argv[]) { assert(device_ptr && "the pointer with offset should be mapped"); -#pragma omp target map(from: ptr) +#pragma omp target map(from : ptr) { ptr = host_data + OFFSET; } assert(ptr == device_ptr && "wrong pointer mapping"); diff --git a/offload/test/api/ompx_3d.c b/offload/test/api/ompx_3d.c index 2ae17c226f..3045c1d2fc 100644 --- a/offload/test/api/ompx_3d.c +++ b/offload/test/api/ompx_3d.c @@ -6,8 +6,7 @@ void foo(int device) { int tid = 0, bid = 0, bdim = 0; -#pragma omp target teams distribute parallel for map(from
- : tid, bid, bdim)
+#pragma omp target teams distribute parallel for map(from : tid, bid, bdim)
device(device) thread_limit(2) num_teams(5) for (int i = 0; i < 1000; ++i) { if (i == 42) { diff --git a/offload/test/api/ompx_3d.cpp b/offload/test/api/ompx_3d.cpp index 5b5491263a..f25b851532 100644 --- a/offload/test/api/ompx_3d.cpp +++ b/offload/test/api/ompx_3d.cpp @@ -6,8 +6,7 @@ void foo(int device) { int tid = 0, bid = 0, bdim = 0; -#pragma omp target teams distribute parallel for map(from
- : tid, bid, bdim)
+#pragma omp target teams distribute parallel for map(from : tid, bid, bdim)
device(device) thread_limit(2) num_teams(5) for (int i = 0; i < 1000; ++i) { if (i == 42) { diff --git a/offload/test/env/base_ptr_ref_count.c b/offload/test/env/base_ptr_ref_count.c index 9a15568712..876991176d 100644 --- a/offload/test/env/base_ptr_ref_count.c +++ b/offload/test/env/base_ptr_ref_count.c @@ -1,5 +1,5 @@ -// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -// REQUIRES: libomptarget-debug +// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 +// %libomptarget-run-generic 2>&1 | %fcheck-generic REQUIRES: libomptarget-debug #include <stdio.h> #include <stdlib.h> diff --git a/offload/test/env/omp_target_debug.c b/offload/test/env/omp_target_debug.c index ec81873a09..8694448acb 100644 --- a/offload/test/env/omp_target_debug.c +++ b/offload/test/env/omp_target_debug.c @@ -1,6 +1,8 @@ -// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG -// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=0 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=NDEBUG -// REQUIRES: libomptarget-debug +// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 +// %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty +// -check-prefix=DEBUG RUN: %libomptarget-compile-generic && env +// LIBOMPTARGET_DEBUG=0 %libomptarget-run-generic 2>&1 | %fcheck-generic +// -allow-empty -check-prefix=NDEBUG REQUIRES: libomptarget-debug int main(void) { #pragma omp target diff --git a/offload/test/jit/empty_kernel.inc b/offload/test/jit/empty_kernel.inc index 43813891cc..3e60ca994a 100644 --- a/offload/test/jit/empty_kernel.inc +++ b/offload/test/jit/empty_kernel.inc @@ -1,15 +1,14 @@ -int main(int argc, char argv) { - #pragma omp TGT1_DIRECTIVE +int main(int argc, char *argv) { +#pragma omp TGT1_DIRECTIVE { #ifdef LOOP_DIRECTIVE - #pragma omp LOOP_DIRECTIVE +#pragma omp LOOP_DIRECTIVE for (int i = 0; i < argc; ++i) #endif { #ifdef BODY_DIRECTIVE - #pragma omp BODY_DIRECTIVE - { - } +#pragma omp BODY_DIRECTIVE + {} #endif } } @@ -18,14 +17,13 @@ int main(int argc, char** argv) { #pragma omp TGT2_DIRECTIVE { #ifdef LOOP_DIRECTIVE - #pragma omp LOOP_DIRECTIVE +#pragma omp LOOP_DIRECTIVE for (int i = 0; i < argc; ++i) #endif { #ifdef BODY_DIRECTIVE - #pragma omp BODY_DIRECTIVE - { - } +#pragma omp BODY_DIRECTIVE + {} #endif } } diff --git a/offload/test/mapping/alloc_fail.c b/offload/test/mapping/alloc_fail.c index c4ae70fc73..31912d4bdb 100644 --- a/offload/test/mapping/alloc_fail.c +++ b/offload/test/mapping/alloc_fail.c @@ -2,9 +2,11 @@ // RUN: %libomptarget-run-fail-generic 2>&1
// RUN: | %fcheck-generic -// CHECK: omptarget message: explicit extension not allowed: host address specified is 0x{{.}} (8 bytes), but device allocation maps to host at 0x{{.*}} (8 bytes) -// CHECK: omptarget error: Call to getTargetPointer returned null pointer (device failure or illegal mapping). -// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory +// CHECK: omptarget message: explicit extension not allowed: host address +// specified is 0x{{.}} (8 bytes), but device allocation maps to host at +// 0x{{.}} (8 bytes) CHECK: omptarget error: Call to getTargetPointer returned +// null pointer (device failure or illegal mapping). CHECK: omptarget fatal +// error 1: failure of target construct while offloading is mandatory int main() { int arr[4] = {0, 1, 2, 3}; diff --git a/offload/test/mapping/array_section_implicit_capture.c b/offload/test/mapping/array_section_implicit_capture.c index 210b7e51cb..1042bd23bd 100644 --- a/offload/test/mapping/array_section_implicit_capture.c +++ b/offload/test/mapping/array_section_implicit_capture.c @@ -1,4 +1,4 @@ -// RUN: %libomptarget-compile-generic +// RUN: %libomptarget-compile-generic // RUN: %libomptarget-run-generic 2>&1
// RUN: | %fcheck-generic diff --git a/offload/test/mapping/data_absent_at_exit.c b/offload/test/mapping/data_absent_at_exit.c index 5baaf933b2..bbbff1e621 100644 --- a/offload/test/mapping/data_absent_at_exit.c +++ b/offload/test/mapping/data_absent_at_exit.c @@ -16,7 +16,8 @@ int main(void) { #pragma omp target data map(from : f) map(always, from : af) { #pragma omp target exit data map(delete : f, af) - } printf("f = %d, af = %d\n", f, af); + } + printf("f = %d, af = %d\n", f, af); // Check omp target exit data. // CHECK: f = 5, r = 6, d = 7, af = 8 diff --git a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp index c6c5657ae6..6408dae555 100644 --- a/offload/test/mapping/declare_mapper_nested_default_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_default_mappers.cpp @@ -44,8 +44,8 @@ int main() { int spp00fa = -1, spp00fca = -1, spp00fb_r = -1; __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]); -#pragma omp target map(tofrom: spp[0][0]) firstprivate(p)
- map(from: spp00fa, spp00fca, spp00fb_r) +#pragma omp target map(tofrom : spp[0][0]) firstprivate(p)
+ map(from : spp00fa, spp00fca, spp00fb_r) { spp00fa = spp[0][0].f.a; spp00fca = spp[0][0].f.c.a; diff --git a/offload/test/mapping/declare_mapper_nested_mappers.cpp b/offload/test/mapping/declare_mapper_nested_mappers.cpp index a9e3f05e0f..cbcb10bbf1 100644 --- a/offload/test/mapping/declare_mapper_nested_mappers.cpp +++ b/offload/test/mapping/declare_mapper_nested_mappers.cpp @@ -42,8 +42,8 @@ int main() { int spp00fa = -1, spp00fb_r = -1, spp00fg1 = -1, spp00fg_r = -1; __intptr_t p = reinterpret_cast<__intptr_t>(&x[0]), p1 = reinterpret_cast<__intptr_t>(&y[0]); -#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)
- map(from: spp00fa, spp00fb_r, spp00fg1, spp00fg_r) +#pragma omp target map(tofrom : spp[0][0]) firstprivate(p, p1)
+ map(from : spp00fa, spp00fb_r, spp00fg1, spp00fg_r) { spp00fa = spp[0][0].f.a; spp00fb_r = spp[0][0].f.b == reinterpret_cast<void *>(p) ? 1 : 0; @@ -56,7 +56,7 @@ int main() { spp[0][0].f.b[1] = 40; spp[0][0].g[1] = 50; } - printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r); + printf("%d %d %d %d\n", spp00fa, spp00fb_r, spp00fg1, spp00fg_r); // CHECK: 222 0 30 0 printf("%d %d %4.5f %d %d %d\n", spp[0][0].e, spp[0][0].f.a, spp[0][0].f.b[1], spp[0][0].f.b == &x[0] ? 1 : 0, spp[0][0].g[1], diff --git a/offload/test/mapping/device_ptr_update.c b/offload/test/mapping/device_ptr_update.c index c6719d2285..77aab6c579 100644 --- a/offload/test/mapping/device_ptr_update.c +++ b/offload/test/mapping/device_ptr_update.c @@ -15,7 +15,8 @@ int main(void) { s1.p = A; -// DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{[}}[[DEV_OBJ_A:0x[^ ]+]]{{]}} +// DEBUG: Update pointer ([[DEV_PTR:0x[^ ]+]]) -> {{[}}[[DEV_OBJ_A:0x[^ +// ]+]]{{]}} #pragma omp target enter data map(alloc : s1.p[0 : 10]) // DEBUG-NOT: Update pointer ([[DEV_PTR]]) -> {{[}}[[DEV_OBJ_A]]{{]}} diff --git a/offload/test/mapping/padding_not_mapped.c b/offload/test/mapping/padding_not_mapped.c index 3ee70ab640..99f57f4fbb 100644 --- a/offload/test/mapping/padding_not_mapped.c +++ b/offload/test/mapping/padding_not_mapped.c @@ -18,7 +18,11 @@ #include <stdio.h> int main() { - struct S { int x; int y; double z; } s = {1, 2, 3}; + struct S { + int x; + int y; + double z; + } s = {1, 2, 3}; // CHECK: &s.x = 0x[[#%x,HOST_ADDR:]], size = [[#%u,SIZE:]] fprintf(stderr, "&s = %p\n", &s); @@ -26,18 +30,20 @@ int main() { fprintf(stderr, "&s.y = %p\n", &s.y); fprintf(stderr, "&s.z = %p\n", &s.z); - // CHECK: s.x is present: 0 - // CHECK: s.x = 1{{$}} - #pragma omp target enter data map(alloc: s.y, s.z) +// CHECK: s.x is present: 0 +// CHECK: s.x = 1{{$}} +#pragma omp target enter data map(alloc : s.y, s.z) int dev = omp_get_default_device(); fprintf(stderr, "s.x is present: %d\n", omp_target_is_present(&s.x, dev)); - #pragma omp target update from(s.x) // should have no effect +#pragma omp target update from(s.x) // should have no effect fprintf(stderr, "s.x = %d\n", s.x); - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory - #pragma omp target enter data map(present, alloc: s.x) +// CHECK: omptarget message: device mapping required by 'present' map type +// modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] +// bytes) CHECK: omptarget error: Call to getTargetPointer returned null pointer +// ('present' map type modifier). CHECK: omptarget fatal error 1: failure of +// target construct while offloading is mandatory +#pragma omp target enter data map(present, alloc : s.x) return 0; } diff --git a/offload/test/mapping/power_of_two_alignment.c b/offload/test/mapping/power_of_two_alignment.c index faebe4f89f..b31c43ac51 100644 --- a/offload/test/mapping/power_of_two_alignment.c +++ b/offload/test/mapping/power_of_two_alignment.c @@ -42,10 +42,13 @@ #include <stdint.h> #include <stdio.h> -template -void test() { +template void test() { StackPad stackPad; - struct S { char x; char y[7]; char z[8]; }; + struct S { + char x; + char y[7]; + char z[8]; + }; struct S collidePre, s, collidePost; uintptr_t mod16 = (uintptr_t)&s % 16; fprintf(stderr, "&s = %p\n", &s); @@ -56,8 +59,8 @@ void test() { } fprintf(stderr, "&collidePre = %p\n", &collidePre); fprintf(stderr, "&collidePost = %p\n", &collidePost); - #pragma omp target data map(to:s.y, s.z) - #pragma omp target data map(to:collidePre, collidePost) +#pragma omp target data map(to : s.y, s.z) +#pragma omp target data map(to : collidePre, collidePost) ; } diff --git a/offload/test/mapping/present/target.c b/offload/test/mapping/present/target.c index 4344c42c80..e5cab61374 100644 --- a/offload/test/mapping/present/target.c +++ b/offload/test/mapping/present/target.c @@ -18,11 +18,13 @@ int main() { // CHECK: i is present fprintf(stderr, "i is present\n"); - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget error: Call to targetDataBegin failed, abort target. - // CHECK: omptarget error: Failed to process data before launching the kernel. - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] + // bytes) CHECK: omptarget error: Call to getTargetPointer returned null + // pointer ('present' map type modifier). CHECK: omptarget error: Call to + // targetDataBegin failed, abort target. CHECK: omptarget error: Failed to + // process data before launching the kernel. CHECK: omptarget fatal error 1: + // failure of target construct while offloading is mandatory #pragma omp target map(present, alloc : i) ; diff --git a/offload/test/mapping/present/target_array_extension.c b/offload/test/mapping/present/target_array_extension.c index 873b2b3617..132a2430ff 100644 --- a/offload/test/mapping/present/target_array_extension.c +++ b/offload/test/mapping/present/target_array_extension.c @@ -42,12 +42,8 @@ #define SMALL_SIZE (SMALL_END - SMALL_BEG) #define LARGE_SIZE (LARGE_END - LARGE_BEG) -#define SMALL
- SMALL_BEG:
- SMALL_SIZE -#define LARGE
- LARGE_BEG:
- LARGE_SIZE +#define SMALL SMALL_BEG : SMALL_SIZE +#define LARGE LARGE_BEG : LARGE_SIZE int main() { int arr[SIZE]; @@ -70,12 +66,16 @@ int main() { // CHECK: arr is present fprintf(stderr, "arr is present\n"); - // CHECK: omptarget message: explicit extension not allowed: host address specified is 0x{{0}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget error: Call to targetDataBegin failed, abort target. - // CHECK: omptarget error: Failed to process data before launching the kernel. - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: explicit extension not allowed: host address + // specified is 0x{{0}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device + // allocation maps to host at 0x{{0}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0}}[[#LARGE_ADDR]] + // ([[#LARGE_BYTES]] bytes) CHECK: omptarget error: Call to getTargetPointer + // returned null pointer ('present' map type modifier). CHECK: omptarget + // error: Call to targetDataBegin failed, abort target. CHECK: omptarget + // error: Failed to process data before launching the kernel. CHECK: omptarget + // fatal error 1: failure of target construct while offloading is mandatory #pragma omp target data map(alloc : arr[SMALL]) { #pragma omp target map(present, tofrom : arr[LARGE]) diff --git a/offload/test/mapping/present/target_data.c b/offload/test/mapping/present/target_data.c index f894283a0c..ae3b762c34 100644 --- a/offload/test/mapping/present/target_data.c +++ b/offload/test/mapping/present/target_data.c @@ -18,8 +18,10 @@ int main() { // CHECK: i is present fprintf(stderr, "i is present\n"); - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] + // bytes) CHECK: omptarget fatal error 1: failure of target construct while + // offloading is mandatory #pragma omp target data map(present, alloc : i) ; diff --git a/offload/test/mapping/present/target_data_array_extension.c b/offload/test/mapping/present/target_data_array_extension.c index 794543a246..7124664756 100644 --- a/offload/test/mapping/present/target_data_array_extension.c +++ b/offload/test/mapping/present/target_data_array_extension.c @@ -42,12 +42,8 @@ #define SMALL_SIZE (SMALL_END - SMALL_BEG) #define LARGE_SIZE (LARGE_END - LARGE_BEG) -#define SMALL
- SMALL_BEG:
- SMALL_SIZE -#define LARGE
- LARGE_BEG:
- LARGE_SIZE +#define SMALL SMALL_BEG : SMALL_SIZE +#define LARGE LARGE_BEG : LARGE_SIZE int main() { int arr[SIZE]; @@ -70,10 +66,14 @@ int main() { // CHECK: arr is present fprintf(stderr, "arr is present\n"); - // CHECK: omptarget message: explicit extension not allowed: host address specified is 0x{{0}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device allocation maps to host at 0x{{0}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: explicit extension not allowed: host address + // specified is 0x{{0}}[[#LARGE_ADDR]] ([[#LARGE_BYTES]] bytes), but device + // allocation maps to host at 0x{{0}}[[#SMALL_ADDR]] ([[#SMALL_BYTES]] bytes) + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0}}[[#LARGE_ADDR]] + // ([[#LARGE_BYTES]] bytes) CHECK: omptarget error: Call to getTargetPointer + // returned null pointer ('present' map type modifier). CHECK: omptarget fatal + // error 1: failure of target construct while offloading is mandatory #pragma omp target data map(alloc : arr[SMALL]) { #pragma omp target data map(present, tofrom : arr[LARGE]) diff --git a/offload/test/mapping/present/target_enter_data.c b/offload/test/mapping/present/target_enter_data.c index 871a05290e..633277ddf0 100644 --- a/offload/test/mapping/present/target_enter_data.c +++ b/offload/test/mapping/present/target_enter_data.c @@ -18,9 +18,11 @@ int main() { // CHECK: i is present fprintf(stderr, "i is present\n"); - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] + // bytes) CHECK: omptarget error: Call to getTargetPointer returned null + // pointer ('present' map type modifier). CHECK: omptarget fatal error 1: + // failure of target construct while offloading is mandatory #pragma omp target enter data map(present, alloc : i) // CHECK-NOT: i is present diff --git a/offload/test/mapping/present/target_exit_data_delete.c b/offload/test/mapping/present/target_exit_data_delete.c index 0fb812b299..d791166147 100644 --- a/offload/test/mapping/present/target_exit_data_delete.c +++ b/offload/test/mapping/present/target_exit_data_delete.c @@ -17,8 +17,10 @@ int main() { // CHECK: i was present fprintf(stderr, "i was present\n"); -// CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0}}[[#HOST_ADDR]] ([[#SIZE]] bytes) -// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory +// CHECK: omptarget message: device mapping required by 'present' map type +// modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] +// bytes) CHECK: omptarget fatal error 1: failure of target construct while +// offloading is mandatory #pragma omp target exit data map(present, delete : i) // CHECK-NOT: i was present diff --git a/offload/test/mapping/present/target_exit_data_release.c b/offload/test/mapping/present/target_exit_data_release.c index 14be22faba..083f6c080d 100644 --- a/offload/test/mapping/present/target_exit_data_release.c +++ b/offload/test/mapping/present/target_exit_data_release.c @@ -17,8 +17,10 @@ int main() { // CHECK: i was present fprintf(stderr, "i was present\n"); -// CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) -// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory +// CHECK: omptarget message: device mapping required by 'present' map type +// modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] +// bytes) CHECK: omptarget fatal error 1: failure of target construct while +// offloading is mandatory #pragma omp target exit data map(present, release : i) // CHECK-NOT: i was present diff --git a/offload/test/mapping/present/target_update.c b/offload/test/mapping/present/target_update.c index 9f6783b6ef..5df4eee6fc 100644 --- a/offload/test/mapping/present/target_update.c +++ b/offload/test/mapping/present/target_update.c @@ -32,8 +32,10 @@ int main() { // CHECK: i is present fprintf(stderr, "i is present\n"); - // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' motion + // modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] + // bytes) CHECK: omptarget fatal error 1: failure of target construct while + // offloading is mandatory #pragma omp target update CLAUSE(present : i) // CHECK-NOT: i is present diff --git a/offload/test/mapping/present/target_update_array_extension.c b/offload/test/mapping/present/target_update_array_extension.c index 11ad4a8d49..f8f3c94d21 100644 --- a/offload/test/mapping/present/target_update_array_extension.c +++ b/offload/test/mapping/present/target_update_array_extension.c @@ -66,8 +66,10 @@ int main() { // CHECK: arr is present fprintf(stderr, "arr is present\n"); - // CHECK: omptarget message: device mapping required by 'present' motion modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] bytes) - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' motion + // modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] ([[#SIZE]] + // bytes) CHECK: omptarget fatal error 1: failure of target construct while + // offloading is mandatory #pragma omp target data map(alloc : arr[SMALL]) { #pragma omp target update CLAUSE(present : arr[LARGE]) diff --git a/offload/test/mapping/present/zero_length_array_section.c b/offload/test/mapping/present/zero_length_array_section.c index e903a268ef..4759585845 100644 --- a/offload/test/mapping/present/zero_length_array_section.c +++ b/offload/test/mapping/present/zero_length_array_section.c @@ -20,11 +20,13 @@ int main() { // arr[0:0] doesn't create an actual mapping in the first directive. // - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) - // CHECK: omptarget error: Call to getTargetPointer returned null pointer ('present' map type modifier). - // CHECK: omptarget error: Call to targetDataBegin failed, abort target. - // CHECK: omptarget error: Failed to process data before launching the kernel. - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) + // CHECK: omptarget error: Call to getTargetPointer returned null pointer + // ('present' map type modifier). CHECK: omptarget error: Call to + // targetDataBegin failed, abort target. CHECK: omptarget error: Failed to + // process data before launching the kernel. CHECK: omptarget fatal error 1: + // failure of target construct while offloading is mandatory #pragma omp target data map(alloc : arr[0 : 0]) #pragma omp target map(present, alloc : arr[0 : 0]) ; diff --git a/offload/test/mapping/present/zero_length_array_section_exit.c b/offload/test/mapping/present/zero_length_array_section_exit.c index 5a7360542e..c88c47f115 100644 --- a/offload/test/mapping/present/zero_length_array_section_exit.c +++ b/offload/test/mapping/present/zero_length_array_section_exit.c @@ -19,8 +19,10 @@ int main() { // arr[0:0] doesn't create an actual mapping in the first directive. // - // CHECK: omptarget message: device mapping required by 'present' map type modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) - // CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory + // CHECK: omptarget message: device mapping required by 'present' map type + // modifier does not exist for host address 0x{{0*}}[[#HOST_ADDR]] (0 bytes) + // CHECK: omptarget fatal error 1: failure of target construct while + // offloading is mandatory #pragma omp target enter data map(alloc : arr[0 : 0]) #pragma omp target exit data map(present, release : arr[0 : 0]) diff --git a/offload/test/mapping/private_mapping.c b/offload/test/mapping/private_mapping.c index 1329a66a5d..157cc4c511 100644 --- a/offload/test/mapping/private_mapping.c +++ b/offload/test/mapping/private_mapping.c @@ -10,9 +10,10 @@ int main() { int data3[3] = {100, 200, 500}; int sum[16] = {0}; - for (int i=0; i<16; i++) sum[i] = 10000; + for (int i = 0; i < 16; i++) + sum[i] = 10000; -#pragma omp target teams distribute parallel for map(tofrom : sum[:16])
+#pragma omp target teams distribute parallel for map(tofrom : sum[ : 16])
firstprivate(data1, data2, data3) for (int i = 0; i < 16; ++i) { for (int j = 0; j < 3; ++j) { diff --git a/offload/test/mapping/target_data_array_extension_at_exit.c b/offload/test/mapping/target_data_array_extension_at_exit.c index e300c800b2..154dc66080 100644 --- a/offload/test/mapping/target_data_array_extension_at_exit.c +++ b/offload/test/mapping/target_data_array_extension_at_exit.c @@ -39,12 +39,8 @@ #error EXTENDS undefined #endif -#define SMALL
- SMALL_BEG:
- (SMALL_END - SMALL_BEG) -#define LARGE
- LARGE_BEG:
- (LARGE_END - LARGE_BEG) +#define SMALL SMALL_BEG : (SMALL_END - SMALL_BEG) +#define LARGE LARGE_BEG : (LARGE_END - LARGE_BEG) void check_not_present() { int arr[SIZE]; diff --git a/offload/test/mapping/target_implicit_partial_map.c b/offload/test/mapping/target_implicit_partial_map.c index a8b2cc893d..fe4b9df72c 100644 --- a/offload/test/mapping/target_implicit_partial_map.c +++ b/offload/test/mapping/target_implicit_partial_map.c @@ -29,7 +29,7 @@ int main() { arr[50] = 5; arr[51] = 6; } // must treat as present (dec ref count) even though full size not present - } // wouldn't delete if previous ref count dec didn't happen + } // wouldn't delete if previous ref count dec didn't happen // CHECK: arr[50] still present: 0 fprintf(stderr, "arr[50] still present: %d\n", diff --git a/offload/test/mapping/target_wrong_use_device_addr.c b/offload/test/mapping/target_wrong_use_device_addr.c index 7a5babd692..28ec6857fa 100644 --- a/offload/test/mapping/target_wrong_use_device_addr.c +++ b/offload/test/mapping/target_wrong_use_device_addr.c @@ -14,7 +14,7 @@ int main() { // CHECK: host addr=0x[[#%x,HOST_ADDR:]] fprintf(stderr, "host addr=%p\n", x); -#pragma omp target data map(to : x [0:10]) +#pragma omp target data map(to : x[0 : 10]) { // CHECK: omptarget device 0 info: variable x does not have a valid device // counterpart @@ -27,4 +27,3 @@ int main() { return 0; }

diff --git a/offload/test/offloading/back2back_distribute.c b/offload/test/offloading/back2back_distribute.c index 63cabd0c66..36d2512deb 100644 --- a/offload/test/offloading/back2back_distribute.c +++ b/offload/test/offloading/back2back_distribute.c @@ -1,4 +1,5 @@ -// RUN: %libomptarget-compile-generic -O3 && %libomptarget-run-generic | %fcheck-generic +// RUN: %libomptarget-compile-generic -O3 && %libomptarget-run-generic | +// %fcheck-generic #include <omp.h> #include <stdio.h> @@ -7,10 +8,10 @@ #define MAX_N 25000 void reset_input(double *a, double *a_h, double *b, double *c) { - for(int i = 0 ; i < MAX_N ; i++) { + for (int i = 0; i < MAX_N; i++) { a[i] = a_h[i] = i; - b[i] = i*2; - c[i] = i-3; + b[i] = i * 2; + c[i] = i - 3; } } @@ -22,15 +23,16 @@ int main(int argc, char *argv[]) { double *b = (double *)calloc(MAX_N, sizeof(double)); double *c = (double *)calloc(MAX_N, sizeof(double)); -#pragma omp target enter data map(to:a[:MAX_N],b[:MAX_N],c[:MAX_N],d[:MAX_N]) +#pragma omp target enter data map(to : a[ : MAX_N], b[ : MAX_N], c[ : MAX_N],
+ d[ : MAX_N]) - for (int n = 32 ; n < MAX_N ; n+=5000) { + for (int n = 32; n < MAX_N; n += 5000) { reset_input(a, a_h, b, c); -#pragma omp target update to(a[:n],b[:n],c[:n],d[:n]) +#pragma omp target update to(a[ : n], b[ : n], c[ : n], d[ : n]) int t = 0; - for (int tms = 1 ; tms <= 256 ; tms *= 2) { // 8 times - for (int ths = 32 ; ths <= 1024 ; ths *= 2) { // 6 times + for (int tms = 1; tms <= 256; tms *= 2) { // 8 times + for (int ths = 32; ths <= 1024; ths *= 2) { // 6 times t++; #pragma omp target #pragma omp teams num_teams(tms) thread_limit(ths) @@ -41,29 +43,31 @@ int main(int argc, char *argv[]) { } #pragma omp distribute parallel for for (int i = 0; i < n; ++i) { - d[i] -= b[i] + c[i]; + d[i] -= b[i] + c[i]; } } } // loop over 'ths' } // loop over 'tms' // check results for each 'n' - for (int times = 0 ; times < t ; times++) { + for (int times = 0; times < t; times++) { for (int i = 0; i < n; ++i) { a_h[i] += b[i] + c[i]; } for (int i = 0; i < n; ++i) d_h[i] -= b[i] + c[i]; } -#pragma omp target update from(a[:n],d[:n]) +#pragma omp target update from(a[ : n], d[ : n]) for (int i = 0; i < n; ++i) { if (a_h[i] != a[i]) { - printf("A Error at n = %d, i = %d: host = %f, device = %f\n", n, i, a_h[i], a[i]); + printf("A Error at n = %d, i = %d: host = %f, device = %f\n", n, i, + a_h[i], a[i]); return 1; } if (d_h[i] != d[i]) { - printf("D Error at n = %d, i = %d: host = %lf, device = %lf\n", n, i, d_h[i], d[i]); + printf("D Error at n = %d, i = %d: host = %lf, device = %lf\n", n, i, + d_h[i], d[i]); return 1; } } diff --git a/offload/test/offloading/bug49021.cpp b/offload/test/offloading/bug49021.cpp index 37da8ce5d8..8b03c4580e 100644 --- a/offload/test/offloading/bug49021.cpp +++ b/offload/test/offloading/bug49021.cpp @@ -1,7 +1,8 @@ // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic -// RUN: %libomptarget-compilexx-generic -O3 -ffast-math && %libomptarget-run-generic -// RUN: %libomptarget-compileoptxx-generic -O3 && %libomptarget-run-generic -// RUN: %libomptarget-compileoptxx-generic -O3 -ffast-math && %libomptarget-run-generic +// RUN: %libomptarget-compilexx-generic -O3 -ffast-math && +// %libomptarget-run-generic RUN: %libomptarget-compileoptxx-generic -O3 && +// %libomptarget-run-generic RUN: %libomptarget-compileoptxx-generic -O3 +// -ffast-math && %libomptarget-run-generic #include diff --git a/offload/test/offloading/bug49334.cpp b/offload/test/offloading/bug49334.cpp index 1f19dab378..a8f6e814c5 100644 --- a/offload/test/offloading/bug49334.cpp +++ b/offload/test/offloading/bug49334.cpp @@ -93,7 +93,7 @@ int BlockMatMul_TargetNowait(BlockMatrix &A, BlockMatrix &B, BlockMatrix &C) { for (int k = 0; k < N / BS; ++k) { float *BlockA = A.GetBlock(i, k); float *BlockB = B.GetBlock(k, j); -// clang-format off + // clang-format off #pragma omp target depend(in: BlockA[0], BlockB[0]) depend(inout: BlockC[0])
map(to: BlockA[:BS * BS], BlockB[:BS * BS])
map(tofrom: BlockC[:BS * BS]) nowait diff --git a/offload/test/offloading/complex_reduction.cpp b/offload/test/offloading/complex_reduction.cpp index 3239f48743..32f06ae4cb 100644 --- a/offload/test/offloading/complex_reduction.cpp +++ b/offload/test/offloading/complex_reduction.cpp @@ -1,5 +1,6 @@ // RUN: %libomptarget-compilexx-generic -O3 && %libomptarget-run-generic -// RUN: %libomptarget-compilexx-generic -O3 -ffast-math && %libomptarget-run-generic +// RUN: %libomptarget-compilexx-generic -O3 -ffast-math && +// %libomptarget-run-generic #include #include @@ -20,8 +21,8 @@ template void test_map() { } #if !defined(__NO_UDR) -#pragma omp declare reduction(+ : std::complex : omp_out += omp_in) -#pragma omp declare reduction(+ : std::complex : omp_out += omp_in) +#pragma omp declare reduction(+ : std::complex : omp_out += omp_in) +#pragma omp declare reduction(+ : std::complex : omp_out += omp_in) #endif template class initiator { @@ -43,7 +44,8 @@ template void test_reduction() { sum_host += array[i]; } -#pragma omp target teams distribute parallel for map(to : array[:size]) reduction(+ : sum) +#pragma omp target teams distribute parallel for map(to : array[ : size])
+ reduction(+ : sum) for (int i = 0; i < size; i++) sum += array[i]; @@ -55,10 +57,8 @@ template void test_reduction() { const int nblock(10), block_size(10); T block_sum[nblock]; -#pragma omp target teams distribute map(to
- : array[:size])
- map(from
- : block_sum[:nblock]) +#pragma omp target teams distribute map(to : array[ : size])
+ map(from : block_sum[ : nblock]) for (int ib = 0; ib < nblock; ib++) { T partial_sum(0); const int istart = ib * block_size; diff --git a/offload/test/offloading/dynamic_module_load.c b/offload/test/offloading/dynamic_module_load.c index 5393f33e84..6e39669021 100644 --- a/offload/test/offloading/dynamic_module_load.c +++ b/offload/test/offloading/dynamic_module_load.c @@ -1,4 +1,6 @@ -// RUN: %libomptarget-compile-generic -DSHARED -fPIC -shared -o %t.so && %clang %flags %s -o %t -ldl && %libomptarget-run-generic %t.so 2>&1 | %fcheck-generic +// RUN: %libomptarget-compile-generic -DSHARED -fPIC -shared -o %t.so && %clang +// %flags %s -o %t -ldl && %libomptarget-run-generic %t.so 2>&1 | +// %fcheck-generic #ifdef SHARED #include <stdio.h> diff --git a/offload/test/offloading/extern.c b/offload/test/offloading/extern.c index 35cf905ee0..df9285ce61 100644 --- a/offload/test/offloading/extern.c +++ b/offload/test/offloading/extern.c @@ -1,11 +1,12 @@ // RUN: %libomptarget-compile-generic -DVAR -c -o %t.o -// RUN: %libomptarget-compile-generic %t.o && %libomptarget-run-generic | %fcheck-generic +// RUN: %libomptarget-compile-generic %t.o && %libomptarget-run-generic | +// %fcheck-generic #ifdef VAR int x = 1; #else -#include <stdio.h> #include <assert.h> +#include <stdio.h> extern int x; int main() { @@ -22,6 +23,6 @@ int main() { assert(value == 999); // CHECK: PASS - printf ("PASS\n"); + printf("PASS\n"); } #endif diff --git a/offload/test/offloading/fortran/basic_array.c b/offload/test/offloading/fortran/basic_array.c index 4daf1d4b0a..0928814bdd 100644 --- a/offload/test/offloading/fortran/basic_array.c +++ b/offload/test/offloading/fortran/basic_array.c @@ -14,7 +14,7 @@ void increment_at(int i, int *array); #pragma omp end declare target void increment_array(int b, int n) { -#pragma omp target map(tofrom : b [0:n]) +#pragma omp target map(tofrom : b[0 : n]) for (int i = 0; i < n; i++) { increment_at(i, b); } diff --git a/offload/test/offloading/global_constructor.cpp b/offload/test/offloading/global_constructor.cpp index eb68c5f783..6260d8994c 100644 --- a/offload/test/offloading/global_constructor.cpp +++ b/offload/test/offloading/global_constructor.cpp @@ -1,4 +1,5 @@ -// RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | %fcheck-generic +// RUN: %libomptarget-compilexx-generic && %libomptarget-run-generic | +// %fcheck-generic #include diff --git a/offload/test/offloading/info.c b/offload/test/offloading/info.c index da8e4c44c5..5e63a252e5 100644 --- a/offload/test/offloading/info.c +++ b/offload/test/offloading/info.c @@ -25,7 +25,7 @@ int main() { int C[N]; int val = 1; -// clang-format off + // clang-format off // INFO: info: Entering OpenMP data region with being_mapper at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments: // INFO: info: alloc(A[0:64])[256] // INFO: info: tofrom(B[0:64])[256] @@ -60,7 +60,7 @@ int main() { // INFO: info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]: // INFO: info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration // INFO: info: [[#%#x,]] [[#%#x,]] 4 INF 0 global at unknown:0:0 -// clang-format on + // clang-format on #pragma omp target data map(alloc : A[0 : N])
map(ompx_hold, tofrom : B[0 : N]) map(to : C[0 : N]) #pragma omp target firstprivate(val) diff --git a/offload/test/offloading/looptripcnt.c b/offload/test/offloading/looptripcnt.c index d1a095038b..5547bbac16 100644 --- a/offload/test/offloading/looptripcnt.c +++ b/offload/test/offloading/looptripcnt.c @@ -1,5 +1,6 @@ -// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG -// REQUIRES: libomptarget-debug +// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 +// %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty +// -check-prefix=DEBUG REQUIRES: libomptarget-debug / Test for looptripcount being popped from runtime stack. diff --git a/offload/test/offloading/mandatory_but_no_devices.c b/offload/test/offloading/mandatory_but_no_devices.c index ecdee72aca..c44b3f6ba8 100644 --- a/offload/test/offloading/mandatory_but_no_devices.c +++ b/offload/test/offloading/mandatory_but_no_devices.c @@ -47,7 +47,8 @@ #include <omp.h> #include <stdio.h> -// CHECK: omptarget fatal error 1: failure of target construct while offloading is mandatory +// CHECK: omptarget fatal error 1: failure of target construct while offloading +// is mandatory int main(void) { int X; #pragma omp DIR device(omp_get_initial_device()) diff --git a/offload/test/offloading/non_contiguous_update.cpp b/offload/test/offloading/non_contiguous_update.cpp index 609f0f967f..3d9ce82361 100644 --- a/offload/test/offloading/non_contiguous_update.cpp +++ b/offload/test/offloading/non_contiguous_update.cpp @@ -1,5 +1,6 @@ -// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty -check-prefix=DEBUG -// REQUIRES: libomptarget-debug +// RUN: %libomptarget-compile-generic && env LIBOMPTARGET_DEBUG=1 +// %libomptarget-run-generic 2>&1 | %fcheck-generic -allow-empty +// -check-prefix=DEBUG REQUIRES: libomptarget-debug #include #include diff --git a/offload/test/offloading/ompx_bare.c b/offload/test/offloading/ompx_bare.c index 3dabdcd15e..991672a21a 100644 --- a/offload/test/offloading/ompx_bare.c +++ b/offload/test/offloading/ompx_bare.c @@ -20,9 +20,11 @@ int main(int argc, char *argv[]) { const int N = num_blocks * block_size; int data = (int )malloc(N * sizeof(int)); - // CHECK: "PluginInterface" device 0 info: Launching kernel _omp_offloading{{.}} with 64 blocks and 64 threads in SPMD mode + // CHECK: "PluginInterface" device 0 info: Launching kernel + // _omp_offloading{{.}} with 64 blocks and 64 threads in SPMD mode -#pragma omp target teams ompx_bare num_teams(num_blocks) thread_limit(block_size) map(from: data[0:N]) +#pragma omp target teams ompx_bare num_teams(num_blocks)
+ thread_limit(block_size) map(from : data[0 : N]) { int bid = ompx_block_id_x(); int bdim = ompx_block_dim_x(); diff --git a/offload/test/offloading/ompx_coords.c b/offload/test/offloading/ompx_coords.c index 5e4e14b4c6..7a6eb27351 100644 --- a/offload/test/offloading/ompx_coords.c +++ b/offload/test/offloading/ompx_coords.c @@ -28,7 +28,7 @@ int main(int argc, char **argv) { int TL = 256; int NT = (N + TL - 1) / TL; -#pragma omp target data map(tofrom : X [0:N]) +#pragma omp target data map(tofrom : X[0 : N]) #pragma omp target teams num_teams(NT) thread_limit(TL) { #pragma omp parallel diff --git a/offload/test/offloading/ompx_saxpy_mixed.c b/offload/test/offloading/ompx_saxpy_mixed.c index f479be8a48..b55c7e7341 100644 --- a/offload/test/offloading/ompx_saxpy_mixed.c +++ b/offload/test/offloading/ompx_saxpy_mixed.c @@ -30,7 +30,7 @@ int main(int argc, char **argv) { int TL = 256; int NT = (N + TL - 1) / TL; -#pragma omp target data map(to : X [0:N]) map(Y [0:N]) +#pragma omp target data map(to : X[0 : N]) map(Y[0 : N]) #pragma omp target teams num_teams(NT) thread_limit(TL) { #pragma omp parallel diff --git a/offload/test/offloading/spmdization.c b/offload/test/offloading/spmdization.c index 77913bec83..26bd7263c0 100644 --- a/offload/test/offloading/spmdization.c +++ b/offload/test/offloading/spmdization.c @@ -31,14 +31,16 @@ int main(void) { nested = omp_get_nested(); tid = omp_get_thread_num(); maxt = omp_get_max_threads(); - #pragma omp parallel - noop();
+#pragma omp parallel + noop(); } - printf("NumThreads: %i, InParallel: %i, Level: %i, ActiveLevel: %i, Nested: %i, " - "ThreadNum: %i, MaxThreads: %i\n", - nthreads, ip, lvl, alvl, nested, tid, maxt); + printf( + "NumThreads: %i, InParallel: %i, Level: %i, ActiveLevel: %i, Nested: %i, " + "ThreadNum: %i, MaxThreads: %i\n", + nthreads, ip, lvl, alvl, nested, tid, maxt); // GENERIC: Generic mode // SPMD: Generic-SPMD mode - // CHECK: NumThreads: 1, InParallel: 0, Level: 0, ActiveLevel: 0, Nested: 0, ThreadNum: 0, MaxThreads: + // CHECK: NumThreads: 1, InParallel: 0, Level: 0, ActiveLevel: 0, Nested: 0, + // ThreadNum: 0, MaxThreads: return 0; } diff --git a/offload/test/offloading/static_linking.c b/offload/test/offloading/static_linking.c index 7be95a10ff..864ec4f294 100644 --- a/offload/test/offloading/static_linking.c +++ b/offload/test/offloading/static_linking.c @@ -1,6 +1,7 @@ // RUN: %libomptarget-compile-generic -DLIBRARY -c -o %t.o // RUN: ar rcs %t.a %t.o -// RUN: %libomptarget-compile-generic %t.a && %libomptarget-run-generic 2>&1 | %fcheck-generic +// RUN: %libomptarget-compile-generic %t.a && %libomptarget-run-generic 2>&1 | +// %fcheck-generic #ifdef LIBRARY int x = 42; diff --git a/offload/test/offloading/target-tile.c b/offload/test/offloading/target-tile.c index 8460b43b6f..2e67ed8ede 100644 --- a/offload/test/offloading/target-tile.c +++ b/offload/test/offloading/target-tile.c @@ -14,10 +14,10 @@ int main() { int order[I_NTILES][J_NTILES][I_NELEMS][J_NELEMS]; int i, j; - #pragma omp target map(tofrom: i, j) +#pragma omp target map(tofrom : i, j) { int next = 0; - #pragma omp tile sizes(I_NELEMS, J_NELEMS) +#pragma omp tile sizes(I_NELEMS, J_NELEMS) for (i = 0; i < I_NTILES * I_NELEMS; ++i) { for (j = 0; j < J_NTILES * J_NELEMS; ++j) { int iTile = i / I_NELEMS; @@ -35,8 +35,8 @@ int main() { for (int jElem = 0; jElem < J_NELEMS; ++jElem) { int actual = order[iTile][jTile][iElem][jElem]; if (expected != actual) { - printf("error: order[%d][%d][%d][%d] = %d, expected %d\n", - iTile, jTile, iElem, jElem, actual, expected); + printf("error: order[%d][%d][%d][%d] = %d, expected %d\n", iTile, + jTile, iElem, jElem, actual, expected); return 1; } ++expected; diff --git a/offload/test/offloading/target_constexpr_mapping.cpp b/offload/test/offloading/target_constexpr_mapping.cpp index 14cf92a7cc..3148b64fb0 100644 --- a/offload/test/offloading/target_constexpr_mapping.cpp +++ b/offload/test/offloading/target_constexpr_mapping.cpp @@ -18,7 +18,7 @@ constexpr static double anotherPi = 3.14; int main() { double a[2]; -#pragma omp target map(tofrom : a[:2]) +#pragma omp target map(tofrom : a[ : 2]) { a[0] = A::pi; a[1] = anotherPi; diff --git a/offload/test/offloading/target_critical_region.cpp b/offload/test/offloading/target_critical_region.cpp index 495632bf76..4f513f7686 100644 --- a/offload/test/offloading/target_critical_region.cpp +++ b/offload/test/offloading/target_critical_region.cpp @@ -24,9 +24,7 @@ int main() { sum[0] = 0; #pragma omp target teams distribute parallel for num_teams(256)
- schedule(static, 1) map(to
- : A[:N]) map(tofrom
- : sum[:1]) + schedule(static, 1) map(to : A[ : N]) map(tofrom : sum[ : 1]) { for (int i = 0; i < N; i++) { #pragma omp critical diff --git a/offload/test/offloading/target_depend_nowait.cpp b/offload/test/offloading/target_depend_nowait.cpp index b0d5408770..f728ec5c79 100644 --- a/offload/test/offloading/target_depend_nowait.cpp +++ b/offload/test/offloading/target_depend_nowait.cpp @@ -43,7 +43,7 @@ int main() { #pragma omp target exit data map(release : A) map(from : B) depend(out : A[0])
nowait } // if - } // parallel + } // parallel int Sum = 0; for (int i = 0; i < N; i++) diff --git a/offload/test/ompt/veccopy_disallow_both.c b/offload/test/ompt/veccopy_disallow_both.c index a011c21955..0cf46c123a 100644 --- a/offload/test/ompt/veccopy_disallow_both.c +++ b/offload/test/ompt/veccopy_disallow_both.c @@ -72,35 +72,27 @@ int main() { /// CHECK-NOT: dest=(nil) /// 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-NOT: dest=(nil) -/// 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-NOT: dest=(nil) 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-NOT: dest=(nil) /// 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: +/// 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/offload/test/ompt/veccopy_map.c b/offload/test/ompt/veccopy_map.c index 56a0dd48f5..7eb76651a9 100644 --- a/offload/test/ompt/veccopy_map.c +++ b/offload/test/ompt/veccopy_map.c @@ -56,31 +56,39 @@ int main() { 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 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: +/// 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/offload/test/ompt/veccopy_no_device_init.c b/offload/test/ompt/veccopy_no_device_init.c index c1a03191c3..f8872e29b3 100644 --- a/offload/test/ompt/veccopy_no_device_init.c +++ b/offload/test/ompt/veccopy_no_device_init.c @@ -54,26 +54,35 @@ int main() { /// 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 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: +/// 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/offload/test/ompt/veccopy_wrong_return.c b/offload/test/ompt/veccopy_wrong_return.c index 2d07b4e1bf..d88ba65e10 100644 --- a/offload/test/ompt/veccopy_wrong_return.c +++ b/offload/test/ompt/veccopy_wrong_return.c @@ -54,26 +54,35 @@ int main() { /// 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 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 +/// 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