TIFitis wrote:
As @agozillon requested, here's a sample of how Clang lowers declare mapper.
**C Code:**
```
typedef struct {
int *ptr;
int buf_size;
} T;
#pragma omp declare mapper(deep_copy : T abc) map(abc, abc.ptr[ : abc.buf_size])
int main() {
T xyz;
#pragma omp target data map(mapper(deep_copy), tofrom : xyz)
{
}
return 0;
}
```
**LLVM IR:**
```
@.offload_sizes = private unnamed_addr constant [1 x i64] [i64 16]
@.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 3]
; Function Attrs: noinline nounwind optnone uwtable
define dso_local i32 @main() #0 {
entry:
%retval = alloca i32, align 4
%xyz = alloca %struct.T, align 8
%.offload_baseptrs = alloca [1 x ptr], align 8
%.offload_ptrs = alloca [1 x ptr], align 8
%.offload_mappers = alloca [1 x ptr], align 8
store i32 0, ptr %retval, align 4
%0 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
store ptr %xyz, ptr %0, align 8
%1 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
store ptr %xyz, ptr %1, align 8
%2 = getelementptr inbounds [1 x ptr], ptr %.offload_mappers, i64 0, i64 0
store ptr @.omp_mapper._ZTS1T.deep_copy, ptr %2, align 8
%3 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
%4 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
call void @__tgt_target_data_begin_mapper(ptr @1, i64 -1, i32 1, ptr %3, ptr
%4, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr
%.offload_mappers)
%5 = getelementptr inbounds [1 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
%6 = getelementptr inbounds [1 x ptr], ptr %.offload_ptrs, i32 0, i32 0
call void @__tgt_target_data_end_mapper(ptr @1, i64 -1, i32 1, ptr %5, ptr
%6, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr
%.offload_mappers)
ret i32 0
}
; Function Attrs: noinline nounwind uwtable
define internal void @.omp_mapper._ZTS1T.deep_copy(ptr noundef %0, ptr noundef
%1, ptr noundef %2, i64 noundef %3, i64 noundef %4, ptr noundef %5) #1 {
entry:
%.addr = alloca ptr, align 8
%.addr1 = alloca ptr, align 8
%.addr2 = alloca ptr, align 8
%.addr3 = alloca i64, align 8
%.addr4 = alloca i64, align 8
%.addr5 = alloca ptr, align 8
store ptr %0, ptr %.addr, align 8
store ptr %1, ptr %.addr1, align 8
store ptr %2, ptr %.addr2, align 8
store i64 %3, ptr %.addr3, align 8
store i64 %4, ptr %.addr4, align 8
store ptr %5, ptr %.addr5, align 8
%6 = load i64, ptr %.addr3, align 8
%7 = load ptr, ptr %.addr, align 8
%8 = load ptr, ptr %.addr1, align 8
%9 = load ptr, ptr %.addr2, align 8
%10 = udiv exact i64 %6, 16
%11 = getelementptr %struct.T, ptr %9, i64 %10
%12 = load i64, ptr %.addr4, align 8
%13 = load ptr, ptr %.addr5, align 8
%omp.arrayinit.isarray = icmp sgt i64 %10, 1
%14 = and i64 %12, 8
%15 = icmp ne ptr %8, %9
%16 = and i64 %12, 16
%17 = icmp ne i64 %16, 0
%18 = and i1 %15, %17
%19 = or i1 %omp.arrayinit.isarray, %18
%.omp.array..init..delete = icmp eq i64 %14, 0
%20 = and i1 %19, %.omp.array..init..delete
br i1 %20, label %.omp.array..init, label %omp.arraymap.head
.omp.array..init: ; preds = %entry
%21 = mul nuw i64 %10, 16
%22 = and i64 %12, -4
%23 = or i64 %22, 512
call void @__tgt_push_mapper_component(ptr %7, ptr %8, ptr %9, i64 %21, i64
%23, ptr %13)
br label %omp.arraymap.head
omp.arraymap.head: ; preds = %.omp.array..init,
%entry
%omp.arraymap.isempty = icmp eq ptr %9, %11
br i1 %omp.arraymap.isempty, label %omp.done, label %omp.arraymap.body
omp.arraymap.body: ; preds = %omp.type.end19,
%omp.arraymap.head
%omp.arraymap.ptrcurrent = phi ptr [ %9, %omp.arraymap.head ], [
%omp.arraymap.next, %omp.type.end19 ]
%ptr = getelementptr inbounds nuw %struct.T, ptr %omp.arraymap.ptrcurrent,
i32 0, i32 0
%ptr6 = getelementptr inbounds nuw %struct.T, ptr %omp.arraymap.ptrcurrent,
i32 0, i32 0
%24 = load ptr, ptr %ptr6, align 8
%arrayidx = getelementptr inbounds nuw i32, ptr %24, i64 0
%buf_size = getelementptr inbounds nuw %struct.T, ptr
%omp.arraymap.ptrcurrent, i32 0, i32 1
%25 = load i32, ptr %buf_size, align 8
%conv = sext i32 %25 to i64
%26 = mul nuw i64 %conv, 4
%27 = getelementptr %struct.T, ptr %omp.arraymap.ptrcurrent, i32 1
%28 = ptrtoint ptr %27 to i64
%29 = ptrtoint ptr %omp.arraymap.ptrcurrent to i64
%30 = sub i64 %28, %29
%31 = sdiv exact i64 %30, ptrtoint (ptr getelementptr (i8, ptr null, i32 1)
to i64)
%32 = call i64 @__tgt_mapper_num_components(ptr %7)
%33 = shl i64 %32, 48
%34 = add nuw i64 0, %33
%35 = and i64 %12, 3
%36 = icmp eq i64 %35, 0
br i1 %36, label %omp.type.alloc, label %omp.type.alloc.else
omp.type.alloc: ; preds = %omp.arraymap.body
%37 = and i64 %34, -4
br label %omp.type.end
omp.type.alloc.else: ; preds = %omp.arraymap.body
%38 = icmp eq i64 %35, 1
br i1 %38, label %omp.type.to, label %omp.type.to.else
omp.type.to: ; preds = %omp.type.alloc.else
%39 = and i64 %34, -3
br label %omp.type.end
omp.type.to.else: ; preds = %omp.type.alloc.else
%40 = icmp eq i64 %35, 2
br i1 %40, label %omp.type.from, label %omp.type.end
omp.type.from: ; preds = %omp.type.to.else
%41 = and i64 %34, -2
br label %omp.type.end
omp.type.end: ; preds = %omp.type.from,
%omp.type.to.else, %omp.type.to, %omp.type.alloc
%omp.maptype = phi i64 [ %37, %omp.type.alloc ], [ %39, %omp.type.to ], [
%41, %omp.type.from ], [ %34, %omp.type.to.else ]
call void @__tgt_push_mapper_component(ptr %7, ptr %omp.arraymap.ptrcurrent,
ptr %omp.arraymap.ptrcurrent, i64 %31, i64 %omp.maptype, ptr null)
%42 = add nuw i64 281474976710659, %33
%43 = and i64 %12, 3
%44 = icmp eq i64 %43, 0
br i1 %44, label %omp.type.alloc7, label %omp.type.alloc.else8
omp.type.alloc7: ; preds = %omp.type.end
%45 = and i64 %42, -4
br label %omp.type.end12
omp.type.alloc.else8: ; preds = %omp.type.end
%46 = icmp eq i64 %43, 1
br i1 %46, label %omp.type.to9, label %omp.type.to.else10
omp.type.to9: ; preds =
%omp.type.alloc.else8
%47 = and i64 %42, -3
br label %omp.type.end12
omp.type.to.else10: ; preds =
%omp.type.alloc.else8
%48 = icmp eq i64 %43, 2
br i1 %48, label %omp.type.from11, label %omp.type.end12
omp.type.from11: ; preds = %omp.type.to.else10
%49 = and i64 %42, -2
br label %omp.type.end12
omp.type.end12: ; preds = %omp.type.from11,
%omp.type.to.else10, %omp.type.to9, %omp.type.alloc7
%omp.maptype13 = phi i64 [ %45, %omp.type.alloc7 ], [ %47, %omp.type.to9 ], [
%49, %omp.type.from11 ], [ %42, %omp.type.to.else10 ]
call void @__tgt_push_mapper_component(ptr %7, ptr %omp.arraymap.ptrcurrent,
ptr %omp.arraymap.ptrcurrent, i64 16, i64 %omp.maptype13, ptr null)
%50 = add nuw i64 281474976710675, %33
%51 = and i64 %12, 3
%52 = icmp eq i64 %51, 0
br i1 %52, label %omp.type.alloc14, label %omp.type.alloc.else15
omp.type.alloc14: ; preds = %omp.type.end12
%53 = and i64 %50, -4
br label %omp.type.end19
omp.type.alloc.else15: ; preds = %omp.type.end12
%54 = icmp eq i64 %51, 1
br i1 %54, label %omp.type.to16, label %omp.type.to.else17
omp.type.to16: ; preds =
%omp.type.alloc.else15
%55 = and i64 %50, -3
br label %omp.type.end19
omp.type.to.else17: ; preds =
%omp.type.alloc.else15
%56 = icmp eq i64 %51, 2
br i1 %56, label %omp.type.from18, label %omp.type.end19
omp.type.from18: ; preds = %omp.type.to.else17
%57 = and i64 %50, -2
br label %omp.type.end19
omp.type.end19: ; preds = %omp.type.from18,
%omp.type.to.else17, %omp.type.to16, %omp.type.alloc14
%omp.maptype20 = phi i64 [ %53, %omp.type.alloc14 ], [ %55, %omp.type.to16 ],
[ %57, %omp.type.from18 ], [ %50, %omp.type.to.else17 ]
call void @__tgt_push_mapper_component(ptr %7, ptr %ptr, ptr %arrayidx, i64
%26, i64 %omp.maptype20, ptr null)
%omp.arraymap.next = getelementptr %struct.T, ptr %omp.arraymap.ptrcurrent,
i32 1
%omp.arraymap.isdone = icmp eq ptr %omp.arraymap.next, %11
br i1 %omp.arraymap.isdone, label %omp.arraymap.exit, label %omp.arraymap.body
omp.arraymap.exit: ; preds = %omp.type.end19
%omp.arrayinit.isarray21 = icmp sgt i64 %10, 1
%58 = and i64 %12, 8
%.omp.array..del..delete = icmp ne i64 %58, 0
%59 = and i1 %omp.arrayinit.isarray21, %.omp.array..del..delete
br i1 %59, label %.omp.array..del, label %omp.done
.omp.array..del: ; preds = %omp.arraymap.exit
%60 = mul nuw i64 %10, 16
%61 = and i64 %12, -4
%62 = or i64 %61, 512
call void @__tgt_push_mapper_component(ptr %7, ptr %8, ptr %9, i64 %60, i64
%62, ptr %13)
br label %omp.done
omp.done: ; preds = %.omp.array..del,
%omp.arraymap.exit, %omp.arraymap.head
ret void
}
```
https://github.com/llvm/llvm-project/pull/117046
_______________________________________________
llvm-branch-commits mailing list
llvm-branch-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
