[PATCH] D84260: [OpenMP][AMDGCN] Support OpenMP offloading for AMDGCN architecture - Part 3

[Alexey Bataev via Phabricator via cfe-commits]

ABataev added inline comments.


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntimeGPU.h:201-205
   /// Declare generalized virtual functions which need to be defined
-  /// by all specializations of OpenMPGPURuntime Targets.
+  /// by all specializations of OpenMPGPURuntime Targets like AMDGCN
+  /// and NVPTX.
+
+  /// Get the GPU warp size.
----------------
Add these notes to the specialized functions too. It is required for better doxygen docs


================
Comment at: clang/test/OpenMP/amdgcn_target_codegen.cpp:17-19
+// CHECK-DAG: [[VAR:%.+]] = trunc i64 [[NUM_THREADS]] to i32
+// CHECK-DAG: sub nuw i32 [[VAR]], 64
+// CHECK-DAG: call i32 @llvm.amdgcn.workitem.id.x()
----------------
Same, the order of these checks should be strict


================
Comment at: clang/test/OpenMP/amdgcn_target_codegen.cpp:34-35
+// CHECK: [[NUM_THREADS:%.+]] = call i64 @__ockl_get_local_size(i32 0)
+// CHECK-DAG: [[VAR:%.+]] = trunc i64 [[NUM_THREADS]] to i32
+// CHECK-DAG: call void @__kmpc_spmd_kernel_init(i32 [[VAR]], i16 0, i16 0)
+#pragma omp target simd
----------------
It should not be `CHECK-DAG`, the order of these 2 instructions is defined and the second one should definitely follow the first one


================
Comment at: clang/test/OpenMP/amdgcn_target_init_temp_alloca.cpp:11
+  // CHECK-LABEL: test_amdgcn_target_temp_alloca
+  // CHECK-LABEL: entry:
+
----------------
`entry:` again


================
Comment at: clang/test/OpenMP/amdgcn_target_init_temp_alloca.cpp:15-17
+  // CHECK:      [[VAR:%.+]].addr = alloca [100 x i32]*, align 8, addrspace(5)
+  // CHECK-NEXT: [[VAR]].addr.ascast = addrspacecast [100 x i32]* addrspace(5)* [[VAR]].addr to [100 x i32]**
+  // CHECK-DAG:  store [100 x i32]* [[VAR]], [100 x i32]** [[VAR]].addr.ascast, align 8
----------------
Do not rely on the names again. Even as part of regexps


================
Comment at: clang/test/OpenMP/amdgcn_target_init_temp_alloca.cpp:17
+  // CHECK-NEXT: [[VAR]].addr.ascast = addrspacecast [100 x i32]* addrspace(5)* [[VAR]].addr to [100 x i32]**
+  // CHECK-DAG:  store [100 x i32]* [[VAR]], [100 x i32]** [[VAR]].addr.ascast, align 8
+
----------------
Just `CHECK`


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D84260