[libc-commits] [PATCH] D148486: [libc] Test the RPC interface with multiple blocks

[Joseph Huber via Phabricator via libc-commits]

jhuber6 created this revision.
jhuber6 added reviewers: jdoerfert, tianshilei1992, JonChesterfield, lntue, sivachandra, tra.
Herald added subscribers: libc-commits, ecnelises, tschuett.
Herald added projects: libc-project, All.
jhuber6 requested review of this revision.

The RPC interface can support multiple independent clients. This support
currently only supports many single-thread warps / workgroups
coordinating over a single lock. This patch uses the support added in
the previous patch to test the RPC interface with multiple blocks.

Note that this does not work with multiple threads currently because of
the effect of warps / workgroups executing in lockstep incorrectly. This
will be added later.

Depends on D148485 <https://reviews.llvm.org/D148485>


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D148486

Files:
  libc/cmake/modules/LLVMLibCTestRules.cmake
  libc/test/integration/startup/gpu/CMakeLists.txt
  libc/test/integration/startup/gpu/rpc_test.cpp


Index: libc/test/integration/startup/gpu/rpc_test.cpp
===================================================================
--- libc/test/integration/startup/gpu/rpc_test.cpp
+++ libc/test/integration/startup/gpu/rpc_test.cpp
@@ -11,10 +11,20 @@
 
 using namespace __llvm_libc;
 
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+uint32_t __nvvm_read_ptx_sreg_ctaid_x();
+uint32_t get_block_id() { return __nvvm_read_ptx_sreg_ctaid_x(); }
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+uint32_t __builtin_amdgcn_workitem_id_x();
+uint32_t get_block_id() { return __builtin_amdgcn_workgroup_id_x(); }
+#else
+uint32_t get_block_id() { return 0; }
+#endif
+
 static void test_add_simple() {
-  constexpr int num_additions = 10000;
+  uint32_t num_additions = 1000 + 10 * get_block_id();
   uint64_t cnt = 0;
-  for (int i = 0; i < num_additions; ++i) {
+  for (uint32_t i = 0; i < num_additions; ++i) {
     rpc::Port port = rpc::client.open(rpc::TEST_INCREMENT);
     port.send_and_recv([=](rpc::Buffer *buffer) { buffer->data[0] = cnt; },
                        [&](rpc::Buffer *buffer) { cnt = buffer->data[0]; });
Index: libc/test/integration/startup/gpu/CMakeLists.txt
===================================================================
--- libc/test/integration/startup/gpu/CMakeLists.txt
+++ libc/test/integration/startup/gpu/CMakeLists.txt
@@ -20,4 +20,6 @@
     rpc_test.cpp
   DEPENDS
     libc.src.__support.RPC.rpc_client
+  LOADER_ARGS
+    --blocks 16
 )
Index: libc/cmake/modules/LLVMLibCTestRules.cmake
===================================================================
--- libc/cmake/modules/LLVMLibCTestRules.cmake
+++ libc/cmake/modules/LLVMLibCTestRules.cmake
@@ -421,7 +421,7 @@
     "INTEGRATION_TEST"
     "" # No optional arguments
     "SUITE" # Single value arguments
-    "SRCS;HDRS;DEPENDS;ARGS;ENV;COMPILE_OPTIONS" # Multi-value arguments
+    "SRCS;HDRS;DEPENDS;ARGS;ENV;COMPILE_OPTIONS;LOADER_ARGS" # Multi-value arguments
     ${ARGN}
   )
 
@@ -532,6 +532,7 @@
     ${fq_target_name}
     COMMAND ${INTEGRATION_TEST_ENV}
             $<$<BOOL:${LIBC_TARGET_ARCHITECTURE_IS_GPU}>:${gpu_loader_exe}>
+            ${INTEGRATION_TEST_LOADER_ARGS}
             $<TARGET_FILE:${fq_build_target_name}> ${INTEGRATION_TEST_ARGS}
     COMMENT "Running integration test ${fq_target_name}"
   )


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D148486.514073.patch
Type: text/x-patch
Size: 2302 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/libc-commits/attachments/20230417/37cb3854/attachment-0001.bin>