[PATCH] D146648: [MLIR][OpenMP] Added MLIR translation support for use_device clauses

[Akash Banerjee via Phabricator via llvm-commits]

TIFitis added a comment.

Here is an example of what clang generates:

C code:

  void foo(int *a, int *b, int *c){
  #pragma omp target data map(from:b, c) use_device_ptr(a) use_device_addr(b){
      *a = 10;
      *b = 20;
      *c = 30;
    }
  }

LLVM IR:

  @.offload_sizes = private unnamed_addr constant [3 x i64] [i64 8, i64 8, i64 0]
  @.offload_maptypes = private unnamed_addr constant [3 x i64] [i64 66, i64 2, i64 64]
  
  define dso_local void @foo(ptr noundef %a, ptr noundef %b, ptr noundef %c) #0 {
  entry:
    %a.addr = alloca ptr, align 8
    %b.addr = alloca ptr, align 8
    %c.addr = alloca ptr, align 8
    %.offload_baseptrs = alloca [3 x ptr], align 8
    %.offload_ptrs = alloca [3 x ptr], align 8
    %.offload_mappers = alloca [3 x ptr], align 8
    %0 = alloca ptr, align 8
    store ptr %a, ptr %a.addr, align 8
    store ptr %b, ptr %b.addr, align 8
    store ptr %c, ptr %c.addr, align 8
    %1 = load ptr, ptr %a.addr, align 8
    %2 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
    store ptr %b.addr, ptr %2, align 8
    %3 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
    store ptr %b.addr, ptr %3, align 8
    %4 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 0
    store ptr null, ptr %4, align 8
    %5 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 1
    store ptr %c.addr, ptr %5, align 8
    %6 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 1
    store ptr %c.addr, ptr %6, align 8
    %7 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 1
    store ptr null, ptr %7, align 8
    %8 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 2
    store ptr %1, ptr %8, align 8
    %9 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 2
    store ptr %1, ptr %9, align 8
    %10 = getelementptr inbounds [3 x ptr], ptr %.offload_mappers, i64 0, i64 2
    store ptr null, ptr %10, align 8
    %11 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
    %12 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
    call void @__tgt_target_data_begin_mapper(ptr @1, i64 -1, i32 3, ptr %11, ptr %12, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
    %13 = load ptr, ptr %8, align 8
    store ptr %13, ptr %0, align 8
    %14 = load ptr, ptr %2, align 8
    %15 = load ptr, ptr %0, align 8
    store i32 10, ptr %15, align 4
    %16 = load ptr, ptr %14, align 8
    store i32 20, ptr %16, align 4
    %17 = load ptr, ptr %c.addr, align 8
    store i32 30, ptr %17, align 4
    %18 = getelementptr inbounds [3 x ptr], ptr %.offload_baseptrs, i32 0, i32 0
    %19 = getelementptr inbounds [3 x ptr], ptr %.offload_ptrs, i32 0, i32 0
    call void @__tgt_target_data_end_mapper(ptr @1, i64 -1, i32 3, ptr %18, ptr %19, ptr @.offload_sizes, ptr @.offload_maptypes, ptr null, ptr null)
    ret void
  }


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D146648/new/

https://reviews.llvm.org/D146648