[Clang][OpenMP] Support expression semantics in target update fields with non-contiguous array sections (#176708)
### Issue: Variable stride not recognized as non-contiguous `CGOpenMPRuntime.cpp` failed to detect `DeclRefExpr`, `MemberExpr`, `ArraySubscriptExpr` as non-contiguous. **Fixes**: `clang/lib/CodeGen/CGOpenMPRuntime.cpp` - Variable stride detection + dimension count logic Detect variable stride expressions (`DeclRefExpr/MemberExpr/ArraySubscriptExpr`) as non-contiguous Added testcases to cover stack arrays, heap pointers, struct members, etc., for expression semantics in non-contiguous update.
This commit is contained in:
Amit Tiwari
GitHub
parent
3548ec9517
commit
14de1bb711
@ -8054,11 +8054,17 @@ private:
|
||||
if (!StrideExpr)
|
||||
return false;
|
||||
|
||||
assert(StrideExpr->getType()->isIntegerType() &&
|
||||
"Stride expression must be of integer type");
|
||||
|
||||
// If stride is not evaluatable as a constant, treat as
|
||||
// non-contiguous.
|
||||
const auto Constant =
|
||||
StrideExpr->getIntegerConstantExpr(CGF.getContext());
|
||||
if (!Constant)
|
||||
return false;
|
||||
return true;
|
||||
|
||||
// Treat non-unitary strides as non-contiguous.
|
||||
return !Constant->isOne();
|
||||
});
|
||||
|
||||
|
||||
135
clang/test/OpenMP/target_update_variable_stride_codegen.c
Normal file
135
clang/test/OpenMP/target_update_variable_stride_codegen.c
Normal file
@ -0,0 +1,135 @@
|
||||
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ --global-value-regex "\.offload_.*" --global-hex-value-regex ".offload_maptypes.*" --version 3
|
||||
|
||||
// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -o - \
|
||||
// RUN: | FileCheck %s
|
||||
|
||||
// Check same results after serialization round-trip
|
||||
// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-pch -o %t %s
|
||||
// RUN: %clang_cc1 -fopenmp -x c -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -include-pch %t -emit-llvm %s -o - \
|
||||
// RUN: | FileCheck %s
|
||||
|
||||
// expected-no-diagnostics
|
||||
#ifndef HEADER
|
||||
#define HEADER
|
||||
|
||||
// Test that variable stride expressions in target update directives correctly
|
||||
// set the OMP_MAP_NON_CONTIG flag. For NON_CONTIG entries, offload_sizes contains
|
||||
// dimension count (2 for 1D array section). For non-NON_CONTIG entries,
|
||||
// offload_sizes contains byte size (5 elements * 4 bytes = 20).
|
||||
int data[10];
|
||||
int stride;
|
||||
|
||||
void test_variable_stride_to() {
|
||||
#pragma omp target update to(data[0:5:stride])
|
||||
}
|
||||
|
||||
void test_variable_stride_from() {
|
||||
#pragma omp target update from(data[0:5:stride])
|
||||
}
|
||||
|
||||
void test_constant_stride_one() {
|
||||
#pragma omp target update to(data[0:5:1])
|
||||
}
|
||||
|
||||
#endif // HEADER
|
||||
//.
|
||||
// CHECK: @.offload_sizes = private unnamed_addr constant [1 x i64] [i64 2]
|
||||
// CHECK: @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 [[#0x100000000001]]]
|
||||
// CHECK: @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 2]
|
||||
// CHECK: @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 [[#0x100000000002]]]
|
||||
// CHECK: @.offload_sizes.3 = private unnamed_addr constant [1 x i64] [i64 20]
|
||||
// CHECK: @.offload_maptypes.4 = private unnamed_addr constant [1 x i64] [i64 [[#0x1]]]
|
||||
//.
|
||||
// CHECK-LABEL: define dso_local void @test_variable_stride_to(
|
||||
// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DIMS:%.*]] = alloca [2 x [[STRUCT_DESCRIPTOR_DIM:%.*]]], align 8
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr @stride, align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = mul nuw i64 4, [[TMP1]]
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP3]], align 8
|
||||
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP4]], align 8
|
||||
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
||||
// CHECK-NEXT: store ptr null, ptr [[TMP5]], align 8
|
||||
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM]]], ptr [[DIMS]], i64 0, i64 0
|
||||
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 0
|
||||
// CHECK-NEXT: store i64 0, ptr [[TMP7]], align 8
|
||||
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 1
|
||||
// CHECK-NEXT: store i64 5, ptr [[TMP8]], align 8
|
||||
// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP6]], i32 0, i32 2
|
||||
// CHECK-NEXT: store i64 [[TMP2]], ptr [[TMP9]], align 8
|
||||
// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM]]], ptr [[DIMS]], i64 0, i64 1
|
||||
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 0
|
||||
// CHECK-NEXT: store i64 0, ptr [[TMP11]], align 8
|
||||
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 1
|
||||
// CHECK-NEXT: store i64 1, ptr [[TMP12]], align 8
|
||||
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM]], ptr [[TMP10]], i32 0, i32 2
|
||||
// CHECK-NEXT: store i64 4, ptr [[TMP13]], align 8
|
||||
// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr [[DIMS]], ptr [[TMP14]], align 8
|
||||
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: call void @__tgt_target_data_update_mapper(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 1, ptr [[TMP15]], ptr [[TMP16]], ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
//
|
||||
// CHECK-LABEL: define dso_local void @test_variable_stride_from(
|
||||
// CHECK-SAME: ) #[[ATTR0]] {
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DIMS:%.*]] = alloca [2 x [[STRUCT_DESCRIPTOR_DIM_0:%.*]]], align 8
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr @stride, align 4
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = mul nuw i64 4, [[TMP1]]
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP3]], align 8
|
||||
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP4]], align 8
|
||||
// CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
||||
// CHECK-NEXT: store ptr null, ptr [[TMP5]], align 8
|
||||
// CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM_0]]], ptr [[DIMS]], i64 0, i64 0
|
||||
// CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 0
|
||||
// CHECK-NEXT: store i64 0, ptr [[TMP7]], align 8
|
||||
// CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 1
|
||||
// CHECK-NEXT: store i64 5, ptr [[TMP8]], align 8
|
||||
// CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP6]], i32 0, i32 2
|
||||
// CHECK-NEXT: store i64 [[TMP2]], ptr [[TMP9]], align 8
|
||||
// CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds [2 x [[STRUCT_DESCRIPTOR_DIM_0]]], ptr [[DIMS]], i64 0, i64 1
|
||||
// CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 0
|
||||
// CHECK-NEXT: store i64 0, ptr [[TMP11]], align 8
|
||||
// CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 1
|
||||
// CHECK-NEXT: store i64 1, ptr [[TMP12]], align 8
|
||||
// CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT_DESCRIPTOR_DIM_0]], ptr [[TMP10]], i32 0, i32 2
|
||||
// CHECK-NEXT: store i64 4, ptr [[TMP13]], align 8
|
||||
// CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr [[DIMS]], ptr [[TMP14]], align 8
|
||||
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: call void @__tgt_target_data_update_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP15]], ptr [[TMP16]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr null, ptr null)
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
//
|
||||
// CHECK-LABEL: define dso_local void @test_constant_stride_one(
|
||||
// CHECK-SAME: ) #[[ATTR0]] {
|
||||
// CHECK-NEXT: entry:
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8
|
||||
// CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP0]], align 8
|
||||
// CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: store ptr @data, ptr [[TMP1]], align 8
|
||||
// CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0
|
||||
// CHECK-NEXT: store ptr null, ptr [[TMP2]], align 8
|
||||
// CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0
|
||||
// CHECK-NEXT: call void @__tgt_target_data_update_mapper(ptr @[[GLOB1]], i64 -1, i32 1, ptr [[TMP3]], ptr [[TMP4]], ptr @.offload_sizes.3, ptr @.offload_maptypes.4, ptr null, ptr null)
|
||||
// CHECK-NEXT: ret void
|
||||
//
|
||||
@ -0,0 +1,145 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests multiple arrays with different variable strides in single update
|
||||
// clause.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
void test_1_update_from_multiple() {
|
||||
int stride1 = 2;
|
||||
int stride2 = 2;
|
||||
double data1[10], data2[10];
|
||||
|
||||
// Initialize data on host
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] = i;
|
||||
data2[i] = i * 10;
|
||||
}
|
||||
|
||||
printf("Test 1: Update FROM - Multiple arrays\n");
|
||||
|
||||
#pragma omp target data map(to : stride1, stride2, data1[0 : 10], data2[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] += i;
|
||||
data2[i] += 100;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
|
||||
}
|
||||
|
||||
printf("from target data1:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data1[i]);
|
||||
|
||||
printf("\nfrom target data2:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data2[i]);
|
||||
}
|
||||
|
||||
void test_2_update_to_multiple() {
|
||||
int stride1 = 2;
|
||||
int stride2 = 2;
|
||||
double data1[10], data2[10];
|
||||
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] = i;
|
||||
data2[i] = i * 10;
|
||||
}
|
||||
|
||||
printf("\nTest 2: Update TO - Multiple arrays\n");
|
||||
|
||||
#pragma omp target data map(tofrom : stride1, stride2, data1[0 : 10], \
|
||||
data2[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] = 100.0;
|
||||
data2[i] = 20.0;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < 10; i += 2) {
|
||||
data1[i] = 10.0;
|
||||
data2[i] = 5.0;
|
||||
}
|
||||
|
||||
#pragma omp target update to(data1[0 : 5 : stride1], data2[0 : 5 : stride2])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] += 2.0;
|
||||
data2[i] += 2.0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
printf("device data1 after update to:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data1[i]);
|
||||
|
||||
printf("\ndevice data2 after update to:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data2[i]);
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_1_update_from_multiple();
|
||||
test_2_update_to_multiple();
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: Test 1: Update FROM - Multiple arrays
|
||||
// CHECK: from target data1:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 16.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: from target data2:
|
||||
// CHECK-NEXT: 100.000000
|
||||
// CHECK-NEXT: 10.000000
|
||||
// CHECK-NEXT: 120.000000
|
||||
// CHECK-NEXT: 30.000000
|
||||
// CHECK-NEXT: 140.000000
|
||||
// CHECK-NEXT: 50.000000
|
||||
// CHECK-NEXT: 160.000000
|
||||
// CHECK-NEXT: 70.000000
|
||||
// CHECK-NEXT: 180.000000
|
||||
// CHECK-NEXT: 90.000000
|
||||
|
||||
// CHECK: Test 2: Update TO - Multiple arrays
|
||||
// CHECK: device data1 after update to:
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
|
||||
// CHECK: device data2 after update to:
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
@ -0,0 +1,136 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests combining variable count expression AND variable stride in array
|
||||
// sections.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
void test_1_update_from() {
|
||||
int len = 10;
|
||||
int stride = 2;
|
||||
double data[len];
|
||||
|
||||
// Initialize data on host
|
||||
for (int i = 0; i < len; i++) {
|
||||
data[i] = i;
|
||||
}
|
||||
|
||||
printf("Test 1: Update FROM - Variable count and stride\n");
|
||||
printf("original values:\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", data[i]);
|
||||
|
||||
#pragma omp target data map(to : len, stride, data[0 : len])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
data[i] += i;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(data[0 : len / 2 : stride])
|
||||
}
|
||||
|
||||
printf("from target results:\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", data[i]);
|
||||
}
|
||||
|
||||
void test_2_update_to() {
|
||||
int len = 10;
|
||||
int stride = 2;
|
||||
double data[len];
|
||||
|
||||
for (int i = 0; i < len; i++) {
|
||||
data[i] = i;
|
||||
}
|
||||
|
||||
printf("\nTest 2: Update TO - Variable count and stride\n");
|
||||
printf("original values:\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", data[i]);
|
||||
|
||||
#pragma omp target data map(tofrom : len, stride, data[0 : len])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
data[i] = 50.0;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < len / 2; i++) {
|
||||
data[i * stride] = 10.0;
|
||||
}
|
||||
|
||||
#pragma omp target update to(data[0 : len / 2 : stride])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
data[i] += 5.0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
printf("device values after update to:\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", data[i]);
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_1_update_from();
|
||||
test_2_update_to();
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: Test 1: Update FROM - Variable count and stride
|
||||
// CHECK: original values:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: from target results:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 16.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: Test 2: Update TO - Variable count and stride
|
||||
// CHECK: original values:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: device values after update to:
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
134
offload/test/offloading/strided_update_variable_stride.c
Normal file
134
offload/test/offloading/strided_update_variable_stride.c
Normal file
@ -0,0 +1,134 @@
|
||||
// This test checks "update from" and "update to" with variable stride.
|
||||
// Tests data[0:5:stride] where stride is a variable, making it non-contiguous.
|
||||
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
void test_1_update_from() {
|
||||
int stride = 2;
|
||||
double data[10];
|
||||
|
||||
// Initialize data on host
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data[i] = i;
|
||||
}
|
||||
|
||||
printf("Test 1: Update FROM device\n");
|
||||
printf("original values:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data[i]);
|
||||
|
||||
#pragma omp target data map(to : stride, data[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data[i] += i;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(data[0 : 5 : stride])
|
||||
}
|
||||
|
||||
printf("from target results:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data[i]);
|
||||
}
|
||||
|
||||
void test_2_update_to() {
|
||||
int stride = 2;
|
||||
double data[10];
|
||||
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data[i] = i;
|
||||
}
|
||||
|
||||
printf("\nTest 2: Update TO device\n");
|
||||
printf("original values:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data[i]);
|
||||
|
||||
#pragma omp target data map(tofrom : stride, data[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data[i] = 50.0;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < 10; i += 2) {
|
||||
data[i] = 10.0;
|
||||
}
|
||||
|
||||
#pragma omp target update to(data[0 : 5 : stride])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data[i] += 5.0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
printf("device values after update to:\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data[i]);
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_1_update_from();
|
||||
test_2_update_to();
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: Test 1: Update FROM device
|
||||
// CHECK: original values:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: from target results:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 12.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 16.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: Test 2: Update TO device
|
||||
// CHECK: original values:
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: device values after update to:
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 15.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
@ -0,0 +1,88 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Miscellaneous variable stride tests: stride=1, stride=array_size, stride from
|
||||
// array subscript.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
void test_1_variable_stride_one() {
|
||||
int stride_one = 1;
|
||||
double data1[10];
|
||||
|
||||
// Initialize data on host
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] = i;
|
||||
}
|
||||
|
||||
#pragma omp target data map(to : stride_one, data1[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data1[i] += i;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(data1[0 : 10 : stride_one])
|
||||
}
|
||||
|
||||
printf("Test 1: Variable stride = 1\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data1[i]);
|
||||
}
|
||||
|
||||
void test_2_variable_stride_large() {
|
||||
int stride_large = 5;
|
||||
double data2[10];
|
||||
|
||||
// Initialize data on host
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data2[i] = i;
|
||||
}
|
||||
|
||||
#pragma omp target data map(to : stride_large, data2[0 : 10])
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < 10; i++) {
|
||||
data2[i] += i;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(data2[0 : 2 : stride_large])
|
||||
}
|
||||
|
||||
printf("\nTest 2: Variable stride = 5\n");
|
||||
for (int i = 0; i < 10; i++)
|
||||
printf("%f\n", data2[i]);
|
||||
}
|
||||
|
||||
int main() {
|
||||
test_1_variable_stride_one();
|
||||
test_2_variable_stride_large();
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: Test 1: Variable stride = 1
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 5.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
|
||||
// CHECK: Test 2: Variable stride = 5
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 1.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 3.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 10.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 7.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
@ -0,0 +1,94 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests heap-allocated pointers with both variable count expression and
|
||||
// variable stride.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int main() {
|
||||
int len = 10;
|
||||
int stride = 2;
|
||||
double *result = (double *)malloc(len * sizeof(double));
|
||||
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = 0;
|
||||
}
|
||||
|
||||
#pragma omp target enter data map(to : len, stride, result[0 : len])
|
||||
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = i;
|
||||
}
|
||||
}
|
||||
|
||||
// Test FROM: Variable count and stride
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] += i * 10;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(result[0 : len / 2 : stride])
|
||||
|
||||
printf("heap ptr variable count and stride (from):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", result[i]);
|
||||
|
||||
// Test TO: Reset, modify host, update to device
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = i * 2;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < len / 2; i++) {
|
||||
result[i * stride] = i + 100;
|
||||
}
|
||||
|
||||
#pragma omp target update to(result[0 : len / 2 : stride])
|
||||
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] += 100;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(result[0 : len])
|
||||
|
||||
printf("heap ptr variable count and stride (to):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", result[i]);
|
||||
|
||||
#pragma omp target exit data map(delete : len, stride, result[0 : len])
|
||||
free(result);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: heap ptr variable count and stride (from):
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 44.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 66.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 88.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK: heap ptr variable count and stride (to):
|
||||
// CHECK-NEXT: 200.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 201.000000
|
||||
// CHECK-NEXT: 106.000000
|
||||
// CHECK-NEXT: 202.000000
|
||||
// CHECK-NEXT: 110.000000
|
||||
// CHECK-NEXT: 203.000000
|
||||
// CHECK-NEXT: 114.000000
|
||||
// CHECK-NEXT: 204.000000
|
||||
// CHECK-NEXT: 118.000000
|
||||
95
offload/test/offloading/target_update_ptr_variable_stride.c
Normal file
95
offload/test/offloading/target_update_ptr_variable_stride.c
Normal file
@ -0,0 +1,95 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests non-contiguous array sections with variable stride on heap-allocated
|
||||
// pointers.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
int main() {
|
||||
int stride = 2;
|
||||
int len = 10;
|
||||
double *result = (double *)malloc(len * sizeof(double));
|
||||
|
||||
// Initialize
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = 0;
|
||||
}
|
||||
|
||||
#pragma omp target enter data map(to : stride, len, result[0 : len])
|
||||
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = i;
|
||||
}
|
||||
}
|
||||
|
||||
// Test FROM
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] += i * 10;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(result[0 : 5 : stride])
|
||||
|
||||
printf("heap ptr variable stride (from):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", result[i]);
|
||||
|
||||
// Test TO: Reset, modify host, update to device
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] = i * 2;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < 5; i++) {
|
||||
result[i * stride] = i + 100;
|
||||
}
|
||||
|
||||
#pragma omp target update to(result[0 : 5 : stride])
|
||||
|
||||
#pragma omp target map(alloc : result[0 : len])
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
result[i] += 100;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(result[0 : len])
|
||||
|
||||
printf("heap ptr variable stride (to):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", result[i]);
|
||||
|
||||
#pragma omp target exit data map(delete : stride, len, result[0 : len])
|
||||
free(result);
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: heap ptr variable stride (from):
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 22.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 44.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 66.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 88.000000
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK: heap ptr variable stride (to):
|
||||
// CHECK-NEXT: 200.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 201.000000
|
||||
// CHECK-NEXT: 106.000000
|
||||
// CHECK-NEXT: 202.000000
|
||||
// CHECK-NEXT: 110.000000
|
||||
// CHECK-NEXT: 203.000000
|
||||
// CHECK-NEXT: 114.000000
|
||||
// CHECK-NEXT: 204.000000
|
||||
// CHECK-NEXT: 118.000000
|
||||
@ -0,0 +1,96 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests struct member arrays with both variable count expression and variable
|
||||
// stride.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
struct S {
|
||||
int len;
|
||||
int stride;
|
||||
double data[20];
|
||||
};
|
||||
|
||||
int main() {
|
||||
struct S s;
|
||||
s.len = 10;
|
||||
s.stride = 2;
|
||||
|
||||
// Initialize
|
||||
#pragma omp target map(tofrom : s)
|
||||
{
|
||||
for (int i = 0; i < s.len; i++) {
|
||||
s.data[i] = i;
|
||||
}
|
||||
}
|
||||
|
||||
// Test FROM: Variable count and stride
|
||||
#pragma omp target data map(to : s)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < s.len; i++) {
|
||||
s.data[i] += i * 10;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(s.data[0 : s.len / 2 : s.stride])
|
||||
}
|
||||
|
||||
printf("struct variable count and stride (from):\n");
|
||||
for (int i = 0; i < s.len; i++)
|
||||
printf("%f\n", s.data[i]);
|
||||
|
||||
// Test TO: Reset, modify host, update to device
|
||||
#pragma omp target map(tofrom : s)
|
||||
{
|
||||
for (int i = 0; i < s.len; i++) {
|
||||
s.data[i] = i * 2;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < s.len / 2; i++) {
|
||||
s.data[i * s.stride] = i + 100;
|
||||
}
|
||||
|
||||
#pragma omp target data map(to : s)
|
||||
{
|
||||
#pragma omp target update to(s.data[0 : s.len / 2 : s.stride])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < s.len; i++) {
|
||||
s.data[i] += 100;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
printf("struct variable count and stride (to):\n");
|
||||
for (int i = 0; i < s.len; i++)
|
||||
printf("%f\n", s.data[i]);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: struct variable count and stride (from):
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 11.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 33.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 77.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
// CHECK: struct variable count and stride (to):
|
||||
// CHECK-NEXT: 100.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 101.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 10.000000
|
||||
// CHECK-NEXT: 103.000000
|
||||
// CHECK-NEXT: 14.000000
|
||||
// CHECK-NEXT: 104.000000
|
||||
// CHECK-NEXT: 18.000000
|
||||
@ -0,0 +1,95 @@
|
||||
// RUN: %libomptarget-compile-run-and-check-generic
|
||||
// Tests non-contiguous array sections with variable stride on struct member
|
||||
// arrays.
|
||||
|
||||
#include <omp.h>
|
||||
#include <stdio.h>
|
||||
|
||||
struct S {
|
||||
int stride;
|
||||
double data[20];
|
||||
};
|
||||
|
||||
int main() {
|
||||
struct S s;
|
||||
s.stride = 2;
|
||||
int len = 10;
|
||||
|
||||
// Initialize
|
||||
#pragma omp target map(tofrom : s, len)
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
s.data[i] = i;
|
||||
}
|
||||
}
|
||||
|
||||
// Test FROM
|
||||
#pragma omp target data map(to : s, len)
|
||||
{
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
s.data[i] += i * 10;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma omp target update from(s.data[0 : 5 : s.stride])
|
||||
}
|
||||
|
||||
printf("struct variable stride (from):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", s.data[i]);
|
||||
|
||||
// Test TO: Reset, modify host, update to device
|
||||
#pragma omp target map(tofrom : s)
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
s.data[i] = i * 2;
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < 5; i++) {
|
||||
s.data[i * s.stride] = i + 100;
|
||||
}
|
||||
|
||||
#pragma omp target data map(to : s)
|
||||
{
|
||||
#pragma omp target update to(s.data[0 : 5 : s.stride])
|
||||
|
||||
#pragma omp target
|
||||
{
|
||||
for (int i = 0; i < len; i++) {
|
||||
s.data[i] += 100;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
printf("struct variable stride (to):\n");
|
||||
for (int i = 0; i < len; i++)
|
||||
printf("%f\n", s.data[i]);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
// CHECK: struct variable stride (from):
|
||||
// CHECK-NEXT: 0.000000
|
||||
// CHECK-NEXT: 11.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 33.000000
|
||||
// CHECK-NEXT: 4.000000
|
||||
// CHECK-NEXT: 55.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 77.000000
|
||||
// CHECK-NEXT: 8.000000
|
||||
// CHECK-NEXT: 9.000000
|
||||
// CHECK: struct variable stride (to):
|
||||
// CHECK-NEXT: 100.000000
|
||||
// CHECK-NEXT: 2.000000
|
||||
// CHECK-NEXT: 101.000000
|
||||
// CHECK-NEXT: 6.000000
|
||||
// CHECK-NEXT: 102.000000
|
||||
// CHECK-NEXT: 10.000000
|
||||
// CHECK-NEXT: 103.000000
|
||||
// CHECK-NEXT: 14.000000
|
||||
// CHECK-NEXT: 104.000000
|
||||
// CHECK-NEXT: 18.000000
|
||||
Loading…
x
Reference in New Issue
Block a user