replica
/
hanchenye-llvm-project
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[OPENMP]Use the real pointer value as base, not indexed value. 

After fix for PR48174 the base pointer for pointer-based
array-sections/array-subscripts will be emitted as `&ptr[idx]`, but
actually it should be just `ptr`, i.e. the address stored in the ponter
to point correctly to the beginning of the array. Currently it may lead
to a crash in the runtime.

Differential Revision: https://reviews.llvm.org/D91805
This commit is contained in:
Alexey Bataev 2020-11-19 09:16:25 -08:00
parent 77e25b5bc8
commit c964f30814
2 changed files with 14 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
clang/lib/CodeGen/CGOpenMPRuntime.cpp
View File

@ -7904,8 +7904,11 @@ private:
IsCaptureFirstInfo = false;
FirstPointerInComplexData = false;
} else if (FirstPointerInComplexData) {
BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
.getAddress(CGF);
QualType Ty = Components.rbegin()
->getAssociatedDeclaration()
->getType()
.getNonReferenceType();
BP = CGF.EmitLoadOfPointer(BP, Ty->castAs<PointerType>());
FirstPointerInComplexData = false;
}
}

11
clang/test/OpenMP/target_data_map_pointer_array_subscript_codegen.cpp
View File

@ -38,10 +38,17 @@ MyObject *objects;
// CHECK-DAG: [[MAPS1:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710673]
// CHECK: @main
int main(void) {
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %{{.+}}, i32 0, i32 0
// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject***
// CHECK: store %struct.MyObject** @objects, %struct.MyObject*** [[BPTRC0]],
// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[SIZES0]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPS0]], i32 0, i32 0), i8** null, i8** null)
#pragma omp target enter data map(to : objects [0:1])
#pragma omp target enter data map(to : objects [1:1])
// CHECK: [[OBJ:%.+]] = load %struct.MyObject*, %struct.MyObject** @objects,
// CHECK: [[BPTR0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
// CHECK: [[BPTRC0:%.+]] = bitcast i8** [[BPTR0]] to %struct.MyObject**
// CHECK: store %struct.MyObject* [[OBJ]], %struct.MyObject** [[BPTRC0]],
// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 2, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPS1]], i32 0, i32 0), i8** null, i8** null)
#pragma omp target enter data map(to : objects[0].arr [0:1])
#pragma omp target enter data map(to : objects[1].arr [0:1])
return 0;
}