[Openmp-commits] [PATCH] D134186: [Clang][OpenMP] Codegen generation for has_device_addr claues.

[Jennifer Yu via Phabricator via Openmp-commits]

jyu2 added inline comments.


================
Comment at: clang/test/OpenMP/target_has_device_addr_codegen.cpp:317
+// CHECK-NEXT:    store [10 x i32]* [[AA]], [10 x i32]** [[RAA]], align 8
+// CHECK-NEXT:    [[TMP2:%.*]] = load i32*, i32** [[K]], align 8
+// CHECK-NEXT:    [[TMP3:%.*]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0
----------------
abhinavgaba wrote:
> This looks the same as `is_device_ptr(k)`. I don't think that's the correct thing to do for `has_device_addr(k)`. For `has_device_addr(k)`, the value being stored to `%offload.baseptrs` should be `i32** [[K]]`, and not a load from it.
> 
> Similarly, the kernel argument should be `[[K]]` and not `[[TMP2]]` (which is a load from `[[K]]`).
Good point.  I mad change by passing &var to BP instead pointer.


================
Comment at: openmp/libomptarget/test/mapping/has_device_addr.cpp:19
+    data_device = (int *)omp_target_alloc(bytes, device_id);
+#pragma omp target teams distribute parallel for has_device_addr(data_device)
+    for (int i = 0; i < size; ++i)
----------------
abhinavgaba wrote:
> The clause used here should either be `has_device_addr(data_device[0])` or `is_device_ptr(data_device)`. `data_device` itself is an `i32**` which is allocated on the host (`%data_device = alloca i32*`). Its pointee has a device address (`%data_device_0 = load i32*, i32** %data_device`).
Changed.  Thanks.  


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D134186