[libc] Add Kernel Resource Usage to nvptx-loader (#97503)

jameshu15869 · jhuber6 · web-flow · commit 1ecffdaf27cb · 2024-07-17T16:07:12.000-05:00
This PR allows `nvptx-loader` to read the resource usage of `_start`,
`_begin`, and `_end` when executing CUDA binaries.

Example output:
```
$ nvptx-loader --print-resource-usage libc/benchmarks/gpu/src/ctype/libc.benchmarks.gpu.src.ctype.isalnum_benchmark.__build__
[ RUN      ] LlvmLibcIsAlNumGpuBenchmark.IsAlnumWrapper
[       OK ] LlvmLibcIsAlNumGpuBenchmark.IsAlnumWrapper: 93 cycles, 76 min, 470 max, 23 iterations, 78000 ns, 80 stddev
_begin registers: 25
_start registers: 80
_end registers: 62
  ```

---------

Co-authored-by: Joseph Huber &lt;huberjn@outlook.com&gt;
diff --git a/libc/benchmarks/gpu/CMakeLists.txt b/libc/benchmarks/gpu/CMakeLists.txt
@@ -15,13 +15,15 @@ function(add_benchmark benchmark_name)
   endif()
   add_libc_hermetic(
     ${benchmark_name}
-    IS_BENCHMARK
+    IS_GPU_BENCHMARK
     LINK_LIBRARIES
       LibcGpuBenchmark.hermetic
       ${BENCHMARK_LINK_LIBRARIES}
     ${BENCHMARK_UNPARSED_ARGUMENTS}
   )
   get_fq_target_name(${benchmark_name} fq_target_name)
+  set(fq_build_target_name ${fq_target_name}.__build__)
+
   add_dependencies(gpu-benchmark ${fq_target_name})
 endfunction(add_benchmark)
 
diff --git a/libc/cmake/modules/LLVMLibCTestRules.cmake b/libc/cmake/modules/LLVMLibCTestRules.cmake
@@ -553,7 +553,7 @@ function(add_libc_hermetic test_name)
   endif()
   cmake_parse_arguments(
     "HERMETIC_TEST"
-    "IS_BENCHMARK" # Optional arguments
+    "IS_GPU_BENCHMARK" # Optional arguments
     "SUITE" # Single value arguments
     "SRCS;HDRS;DEPENDS;ARGS;ENV;COMPILE_OPTIONS;LINK_LIBRARIES;LOADER_ARGS" # Multi-value arguments
     ${ARGN}
@@ -709,14 +709,24 @@ function(add_libc_hermetic test_name)
       $<TARGET_FILE:${fq_build_target_name}> ${HERMETIC_TEST_ARGS})
   add_custom_target(
     ${fq_target_name}
+    DEPENDS ${fq_target_name}-cmd
+  )
+
+  add_custom_command(
+    OUTPUT ${fq_target_name}-cmd
     COMMAND ${test_cmd}
     COMMAND_EXPAND_LISTS
     COMMENT "Running hermetic test ${fq_target_name}"
     ${LIBC_HERMETIC_TEST_JOB_POOL}
   )
 
+  set_source_files_properties(${fq_target_name}-cmd
+    PROPERTIES
+      SYMBOLIC "TRUE"
+  )
+
   add_dependencies(${HERMETIC_TEST_SUITE} ${fq_target_name})
-  if(NOT ${HERMETIC_TEST_IS_BENCHMARK})
+  if(NOT ${HERMETIC_TEST_IS_GPU_BENCHMARK})
     # If it is a benchmark, it will already have been added to the
     # gpu-benchmark target
     add_dependencies(libc-hermetic-tests ${fq_target_name})
diff --git a/libc/utils/gpu/loader/Loader.h b/libc/utils/gpu/loader/Loader.h
@@ -54,7 +54,7 @@ struct end_args_t {
 /// kernel on the target device. Copies \p argc and \p argv to the device.
 /// Returns the final value of the `main` function on the device.
 int load(int argc, char **argv, char **evnp, void *image, size_t size,
-         const LaunchParameters &params);
+         const LaunchParameters &params, bool print_resource_usage);
 
 /// Return \p V aligned "upwards" according to \p Align.
 template <typename V, typename A> inline V align_up(V val, A align) {
diff --git a/libc/utils/gpu/loader/Main.cpp b/libc/utils/gpu/loader/Main.cpp
@@ -20,7 +20,8 @@
 
 int main(int argc, char **argv, char **envp) {
   if (argc < 2) {
-    printf("USAGE: ./loader [--threads <n>, --blocks <n>] <device_image> "
+    printf("USAGE: ./loader [--threads <n>, --blocks <n>, "
+           "--print-resource-usage] <device_image> "
            "<args>, ...\n");
     return EXIT_SUCCESS;
   }
@@ -29,6 +30,7 @@ int main(int argc, char **argv, char **envp) {
   FILE *file = nullptr;
   char *ptr;
   LaunchParameters params = {1, 1, 1, 1, 1, 1};
+  bool print_resource_usage = false;
   while (!file && ++offset < argc) {
     if (argv[offset] == std::string("--threads") ||
         argv[offset] == std::string("--threads-x")) {
@@ -62,6 +64,9 @@ int main(int argc, char **argv, char **envp) {
           offset + 1 < argc ? strtoul(argv[offset + 1], &ptr, 10) : 1;
       offset++;
       continue;
+    } else if (argv[offset] == std::string("--print-resource-usage")) {
+      print_resource_usage = true;
+      continue;
     } else {
       file = fopen(argv[offset], "r");
       if (!file) {
@@ -87,7 +92,8 @@ int main(int argc, char **argv, char **envp) {
   fclose(file);
 
   // Drop the loader from the program arguments.
-  int ret = load(argc - offset, &argv[offset], envp, image, size, params);
+  int ret = load(argc - offset, &argv[offset], envp, image, size, params,
+                 print_resource_usage);
 
   free(image);
   return ret;
diff --git a/libc/utils/gpu/loader/amdgpu/Loader.cpp b/libc/utils/gpu/loader/amdgpu/Loader.cpp
@@ -125,6 +125,10 @@ hsa_status_t get_agent(hsa_agent_t *output_agent) {
   return iterate_agents(cb);
 }
 
+void print_kernel_resources(char *kernel_name) {
+  fprintf("Kernel resources on AMDGPU is not supported yet.\n");
+}
+
 /// Retrieve a global memory pool with a \p flag from the agent.
 template <hsa_amd_memory_pool_global_flag_t flag>
 hsa_status_t get_agent_memory_pool(hsa_agent_t agent,
@@ -156,8 +160,9 @@ hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
                            hsa_amd_memory_pool_t coarsegrained_pool,
                            hsa_queue_t *queue, rpc_device_t device,
                            const LaunchParameters &params,
-                           const char *kernel_name, args_t kernel_args) {
-  // Look up the '_start' kernel in the loaded executable.
+                           const char *kernel_name, args_t kernel_args,
+                           bool print_resource_usage) {
+  // Look up the kernel in the loaded executable.
   hsa_executable_symbol_t symbol;
   if (hsa_status_t err = hsa_executable_get_symbol_by_name(
           executable, kernel_name, &dev_agent, &symbol))
@@ -220,7 +225,7 @@ hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
     handle_error(err);
   hsa_amd_agents_allow_access(1, &dev_agent, nullptr, args);
 
-  // Initialie all the arguments (explicit and implicit) to zero, then set the
+  // Initialize all the arguments (explicit and implicit) to zero, then set the
   // explicit arguments to the values created above.
   std::memset(args, 0, args_size);
   std::memcpy(args, &kernel_args, sizeof(args_t));
@@ -270,6 +275,9 @@ hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
           hsa_signal_create(1, 0, nullptr, &packet->completion_signal))
     handle_error(err);
 
+  if (print_resource_usage)
+    print_kernel_resources(kernel_name);
+
   // Initialize the packet header and set the doorbell signal to begin execution
   // by the HSA runtime.
   uint16_t header =
@@ -327,7 +335,7 @@ static hsa_status_t hsa_memcpy(void *dst, hsa_agent_t dst_agent,
 }
 
 int load(int argc, char **argv, char **envp, void *image, size_t size,
-         const LaunchParameters &params) {
+         const LaunchParameters &params, bool print_resource_usage) {
   // Initialize the HSA runtime used to communicate with the device.
   if (hsa_status_t err = hsa_init())
     handle_error(err);
@@ -545,15 +553,16 @@ int load(int argc, char **argv, char **envp, void *image, size_t size,
 
   LaunchParameters single_threaded_params = {1, 1, 1, 1, 1, 1};
   begin_args_t init_args = {argc, dev_argv, dev_envp};
-  if (hsa_status_t err = launch_kernel(
-          dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
-          device, single_threaded_params, "_begin.kd", init_args))
+  if (hsa_status_t err = launch_kernel(dev_agent, executable, kernargs_pool,
+                                       coarsegrained_pool, queue, device,
+                                       single_threaded_params, "_begin.kd",
+                                       init_args, print_resource_usage))
     handle_error(err);
 
   start_args_t args = {argc, dev_argv, dev_envp, dev_ret};
-  if (hsa_status_t err = launch_kernel(dev_agent, executable, kernargs_pool,
-                                       coarsegrained_pool, queue, device,
-                                       params, "_start.kd", args))
+  if (hsa_status_t err = launch_kernel(
+          dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
+          device, params, "_start.kd", args, print_resource_usage))
     handle_error(err);
 
   void *host_ret;
@@ -571,9 +580,10 @@ int load(int argc, char **argv, char **envp, void *image, size_t size,
   int ret = *static_cast<int *>(host_ret);
 
   end_args_t fini_args = {ret};
-  if (hsa_status_t err = launch_kernel(
-          dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
-          device, single_threaded_params, "_end.kd", fini_args))
+  if (hsa_status_t err = launch_kernel(dev_agent, executable, kernargs_pool,
+                                       coarsegrained_pool, queue, device,
+                                       single_threaded_params, "_end.kd",
+                                       fini_args, print_resource_usage))
     handle_error(err);
 
   if (rpc_status_t err = rpc_server_shutdown(
diff --git a/libc/utils/gpu/loader/nvptx/Loader.cpp b/libc/utils/gpu/loader/nvptx/Loader.cpp
@@ -152,10 +152,23 @@ Expected<void *> get_ctor_dtor_array(const void *image, const size_t size,
   return dev_memory;
 }
 
+void print_kernel_resources(CUmodule binary, const char *kernel_name) {
+  CUfunction function;
+  if (CUresult err = cuModuleGetFunction(&function, binary, kernel_name))
+    handle_error(err);
+  int num_regs;
+  if (CUresult err =
+          cuFuncGetAttribute(&num_regs, CU_FUNC_ATTRIBUTE_NUM_REGS, function))
+    handle_error(err);
+  printf("Executing kernel %s:\n", kernel_name);
+  printf("%6s registers: %d\n", kernel_name, num_regs);
+}
+
 template <typename args_t>
 CUresult launch_kernel(CUmodule binary, CUstream stream,
                        rpc_device_t rpc_device, const LaunchParameters &params,
-                       const char *kernel_name, args_t kernel_args) {
+                       const char *kernel_name, args_t kernel_args,
+                       bool print_resource_usage) {
   // look up the '_start' kernel in the loaded module.
   CUfunction function;
   if (CUresult err = cuModuleGetFunction(&function, binary, kernel_name))
@@ -208,6 +221,9 @@ CUresult launch_kernel(CUmodule binary, CUstream stream,
       },
       &memory_stream);
 
+  if (print_resource_usage)
+    print_kernel_resources(binary, kernel_name);
+
   // Call the kernel with the given arguments.
   if (CUresult err = cuLaunchKernel(
           function, params.num_blocks_x, params.num_blocks_y,
@@ -230,7 +246,7 @@ CUresult launch_kernel(CUmodule binary, CUstream stream,
 }
 
 int load(int argc, char **argv, char **envp, void *image, size_t size,
-         const LaunchParameters &params) {
+         const LaunchParameters &params, bool print_resource_usage) {
   if (CUresult err = cuInit(0))
     handle_error(err);
   // Obtain the first device found on the system.
@@ -323,14 +339,15 @@ int load(int argc, char **argv, char **envp, void *image, size_t size,
 
   LaunchParameters single_threaded_params = {1, 1, 1, 1, 1, 1};
   begin_args_t init_args = {argc, dev_argv, dev_envp};
-  if (CUresult err = launch_kernel(binary, stream, rpc_device,
-                                   single_threaded_params, "_begin", init_args))
+  if (CUresult err =
+          launch_kernel(binary, stream, rpc_device, single_threaded_params,
+                        "_begin", init_args, print_resource_usage))
     handle_error(err);
 
   start_args_t args = {argc, dev_argv, dev_envp,
                        reinterpret_cast<void *>(dev_ret)};
-  if (CUresult err =
-          launch_kernel(binary, stream, rpc_device, params, "_start", args))
+  if (CUresult err = launch_kernel(binary, stream, rpc_device, params, "_start",
+                                   args, print_resource_usage))
     handle_error(err);
 
   // Copy the return value back from the kernel and wait.
@@ -342,8 +359,9 @@ int load(int argc, char **argv, char **envp, void *image, size_t size,
     handle_error(err);
 
   end_args_t fini_args = {host_ret};
-  if (CUresult err = launch_kernel(binary, stream, rpc_device,
-                                   single_threaded_params, "_end", fini_args))
+  if (CUresult err =
+          launch_kernel(binary, stream, rpc_device, single_threaded_params,
+                        "_end", fini_args, print_resource_usage))
     handle_error(err);
 
   // Free the memory allocated for the device.
