diff --git a/libc/cmake/modules/LLVMLibCTestRules.cmake b/libc/cmake/modules/LLVMLibCTestRules.cmake --- a/libc/cmake/modules/LLVMLibCTestRules.cmake +++ b/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} $<$:${gpu_loader_exe}> + ${INTEGRATION_TEST_LOADER_ARGS} $ ${INTEGRATION_TEST_ARGS} COMMENT "Running integration test ${fq_target_name}" ) diff --git a/libc/test/integration/startup/gpu/CMakeLists.txt b/libc/test/integration/startup/gpu/CMakeLists.txt --- a/libc/test/integration/startup/gpu/CMakeLists.txt +++ b/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 ) diff --git a/libc/test/integration/startup/gpu/rpc_test.cpp b/libc/test/integration/startup/gpu/rpc_test.cpp --- a/libc/test/integration/startup/gpu/rpc_test.cpp +++ b/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]; });