Differential D80649 Diff 268589 openmp/libomptarget/test/offloading/d2d_memcpy.c

Changeset View

Standalone View

openmp/libomptarget/test/offloading/d2d_memcpy.c

This file was added.

				// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu \| %fcheck-aarch64-unknown-linux-gnu
				// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu \| %fcheck-powerpc64-ibm-linux-gnu
				// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu \| %fcheck-powerpc64le-ibm-linux-gnu
				// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu \| %fcheck-x86_64-pc-linux-gnu -allow-empty

				#include <assert.h>
				#include <omp.h>
				#include <stdio.h>
				#include <stdlib.h>

				const int magic_num = 7;

				int main(int argc, char *argv[]) {
				const int N = 128;
				const int num_devices = omp_get_num_devices();

				// No target device, just return
				if (num_devices == 0) {
				printf("PASS\n");
				return 0;
				}

				const int src_device = 0;
				int dst_device = 1;
				if (dst_device >= num_devices)
				dst_device = num_devices - 1;
				protze.joachimUnsubmitted Not Done Reply Inline Actions What is this supposed to do? I guess you intended to use a different condition in the if, right? Why not unconditionally set the dst_device to num_devices-1? protze.joachim: What is this supposed to do? I guess you intended to use a different condition in the if…
				tianshilei1992AuthorUnsubmitted Done Reply Inline Actions Yeah, that would be better. tianshilei1992: Yeah, that would be better.

				int length = N * sizeof(int);
				int *src_ptr = omp_target_alloc(length, src_device);
				int *dst_ptr = omp_target_alloc(length, dst_device);

				assert(src_ptr && "src_ptr is NULL");
				assert(dst_ptr && "dst_ptr is NULL");

				#pragma omp target teams distribute parallel for device(src_device) \
				is_device_ptr(src_ptr)
				for (int i = 0; i < N; ++i) {
				src_ptr[i] = magic_num;
				}

				int rc =
				omp_target_memcpy(dst_ptr, src_ptr, length, 0, 0, dst_device, src_device);

				assert(rc == 0 && "error in omp_target_memcpy");

				int *buffer = malloc(length);

				assert(buffer && "failed to allocate host buffer");

				#pragma omp target teams distribute parallel for device(dst_device) \
				map(from: buffer[0:N]) is_device_ptr(dst_ptr)
				for (int i = 0; i < N; ++i) {
				buffer[i] = dst_ptr[i] + magic_num;
				}

				for (int i = 0; i < N; ++i)
				assert(buffer[i] == 2 * magic_num);

				printf("PASS\n");

				// Free host and device memory
				free(buffer);
				omp_target_free(src_ptr, src_device);
				omp_target_free(dst_ptr, dst_device);

				return 0;
				}

				// CHECK: PASS