Handle optimization case w.r.t. load on GEP

Author: ahueck

lib/type/DITypeExtractor.cpp

@@ -83,7 +83,7 @@ bool load_to(const llvm::LoadInst* load) {
   return detail::get_operand_to<T>(load).has_value();
 }
 
-std::optional<llvm::DIType*> reset_load_related_basic(const dataflow::ValuePath&, llvm::DIType* type_to_reset,
+std::optional<llvm::DIType*> reset_load_related_basic(const dataflow::ValuePath& path, llvm::DIType* type_to_reset,
                                                       const llvm::LoadInst* load) {
   auto* type = type_to_reset;
 
@@ -101,6 +101,25 @@ std::optional<llvm::DIType*> reset_load_related_basic(const dataflow::ValuePath&
     return type;
   }
 
+  // a (last?) load to a GEP of a composite likely loads the first member in an optimized context:
+  const bool last_load = path.start_value().value_or(nullptr) == load;
+  if (last_load && load_to<llvm::GetElementPtrInst>(load)) {
+    if (auto* may_be_member_type = llvm::dyn_cast<llvm::DIDerivedType>(type)) {
+      LOG_DEBUG("Load on GEP, return basetype " << log::ditype_str(may_be_member_type->getBaseType()))
+      if (di::util::is_member(*may_be_member_type)) {
+        auto type_of_member = may_be_member_type->getBaseType();
+        if (auto member = llvm::dyn_cast<llvm::DIDerivedType>(type_of_member)) {
+          may_be_member_type = member;
+          if (auto* member_composite_type = llvm::dyn_cast<llvm::DICompositeType>(member->getBaseType())) {
+            auto members_of_composite_type = di::util::get_composite_members(*member_composite_type);
+            assert(members_of_composite_type.size() > 0 && "Load to composite expects at least one member");
+            return (*members_of_composite_type.begin());
+          }
+        }
+      }
+    }
+  }
+
   if (auto* maybe_ptr_to_type = llvm::dyn_cast<llvm::DIDerivedType>(type)) {
     if (di::util::is_pointer(*maybe_ptr_to_type)) {
       LOG_DEBUG("Load of pointer-like " << log::ditype_str(maybe_ptr_to_type))

lib/type/Dimeta.cpp

@@ -253,6 +253,7 @@ std::optional<QualifiedType> type_for(const llvm::Value* value) {
 
   return {};
 }
+
 }  // namespace experimental
 
 }  // namespace dimeta

lib/type/Dimeta.h

@@ -27,12 +27,12 @@ class Module;
 
 namespace dimeta {
 
-using DimetaDIVar = std::variant<llvm::DILocalVariable*, llvm::DIGlobalVariable*>;
+using DimetaDIVariable = std::variant<llvm::DILocalVariable*, llvm::DIGlobalVariable*>;
 
 struct DimetaData {
   enum MemLoc { kStack = 0, kHeap, kGlobal };
   MemLoc memory_location{kStack};
-  std::optional<DimetaDIVar> di_variable{};        // if existing the named variable w.r.t. allocation
+  std::optional<DimetaDIVariable> di_variable{};   // if existing the named variable w.r.t. allocation
   std::optional<llvm::DIType*> entry_type{};       // determined to be the allocation including "pointer" DITypes
   std::optional<llvm::DIType*> base_type{};        // The base type (int, struct X...) of the allocated memory
   std::optional<llvm::DILocation*> di_location{};  // Loc of call (malloc etc.)/alloca. Not set for global
@@ -59,7 +59,8 @@ std::optional<CompileUnitTypeList> compile_unit_types(const llvm::Module*);
 
 namespace experimental {
 std::optional<QualifiedType> type_for(const llvm::Value*);
-}
+
+}  // namespace experimental
 
 }  // namespace dimeta
 

test/pass/cuda/heap_cudamalloc_member_nested.c

@@ -0,0 +1,19 @@
+// RUN: %c-to-llvm %s | %apply-verifier 2>&1 | %filecheck %s
+// RUN: %c-to-llvm %s | %opt -O2 -S | %apply-verifier 2>&1 | %filecheck %s
+
+extern int cudaMalloc(void** devPtr, int size);
+
+struct B {
+  double* x;
+};
+
+struct A {
+  float b;
+  float c;
+  struct B* a;
+};
+
+void foo(struct A* ar) {
+  // CHECK: Final Type: {{.*}} = !DIBasicType(name: "double", size: 64, encoding: DW_ATE_float)
+  cudaMalloc((void**)&ar->a->x, sizeof(double));
+}

test/pass/gep/array.c

@@ -1,4 +1,5 @@
 // RUN: %c-to-llvm %s | %apply-verifier 2>&1 | %filecheck %s
+// RUN: %c-to-llvm %s | %opt -O2 | %apply-verifier 2>&1 | %filecheck %s
 
 #include <stdlib.h>