[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/CGOpenMPRuntimeAMDGCN.cpp:51
+
+/// Get the maximum number of threads in a block of the GPU.
+llvm::Value *CGOpenMPRuntimeAMDGCN::getGPUNumThreads(CodeGenFunction &CGF) {
----------------
Move these comments to the header instead and remove them from the .cpp module.


================
Comment at: clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp:43
+
+/// Get the id of the current thread on the GPU.
+llvm::Value *CGOpenMPRuntimeNVPTX::getGPUThreadID(CodeGenFunction &CGF) {
----------------
Same, move these comments to the header.


================
Comment at: clang/test/OpenMP/amdgcn_target_codegen.cpp:13
+  // CHECK-LABEL: test_amdgcn_target_tid_threads
+  // CHECK-LABEL: entry:
+
----------------
`entry:` not always exist in the LLVM IR functions, better not to check for it.


================
Comment at: clang/test/OpenMP/amdgcn_target_codegen.cpp:19-20
+// CHECK-DAG: [[VAR1:%[0-9]+]] = trunc i64 %nvptx_num_threads to i32
+// CHECK-DAG: %thread_limit = sub nuw i32 [[VAR1]], 64
+// CHECK-DAG: %nvptx_tid{{[0-9]*}} = call i32 @llvm.amdgcn.workitem.id.x()
+#pragma omp target
----------------
Do not rely on names for the locals, some release builds do not generate them. Use regexps instead.


================
Comment at: clang/test/OpenMP/amdgcn_target_init_temp_alloca.cpp:15-16
+
+  // CHECK:      %arr.addr = alloca [100 x i32]*, align 8, addrspace(5)
+  // CHECK-NEXT: %arr.addr.ascast = addrspacecast [100 x i32]* addrspace(5)* %arr.addr to [100 x i32]**
+  // CHECK-DAG:  store [100 x i32]* %arr, [100 x i32]** %arr.addr.ascast, align 8
----------------
same, use regexps instead of %-like names


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D84260