[libc] Support fopen / fclose on the GPU

[jhuber6]

Summary

This patch adds the necessary support for the fopen and fclose functions
to work on the GPU via RPC. I added a new test that enables testing this
with the minimal features we have on the GPU. I will update it once we
have fread and fwrite to actually check the outputted strings. For
now I just relied on checking manually via the outpuot temp file.

[sivachandra]

OK for the non-GPU parts.

[JonChesterfield]

There's a design question in here I think. Specifically whether FILE* is some handle to a host only thing or something that lives on the GPU. I can see arguments both ways, this LGTM too.

[jhuber6]

In D154519#4475537, @JonChesterfield wrote:

There's a design question in here I think. Specifically whether FILE* is some handle to a host only thing or something that lives on the GPU. I can see arguments both ways, this LGTM too.

So that's something I was wondering about, because as it stands the following will not work if we're doing an offloading language

int main() {
  FILE *fp = fopen("foo.txt", "w");
#pragma omp target
  { fputs("foo", fp); }
}

This is because these both represent fundamentally different structs. We could potentially make the interface just a uintptr_t to a host FILE *, but this would necessitate bypassing completely the file abstractions in libc and implementing each entrypoint separate on the GPU, versus what we do now where we more or less treat RPC as the syscall. We could potentially detect if a FILE * came from us or the host via pointer tagging in the LSB, that would allow us to wrap it in a new file if needed, but going from the GPU to the CPU then still wouldn't work. Alternatively we can just be like "Don't mix the two they're not the same libc" and call it a day.

libc/config/gpu/entrypoints.txt

set(TARGET_LIBC_ENTRYPOINTS
     libc.src.stdlib.free
 
     # errno.h entrypoints
     libc.src.errno.errno
 
     # stdio.h entrypoints
     libc.src.stdio.puts
     libc.src.stdio.fputs
+    libc.src.stdio.fclose
+    libc.src.stdio.fopen
     libc.src.stdio.stdin
     libc.src.stdio.stdout
     libc.src.stdio.stderr
 
     # inttypes.h entrypoints
     libc.src.inttypes.imaxabs
     libc.src.inttypes.imaxdiv
     libc.src.inttypes.strtoimax

libc/docs/gpu/support.rst

 stdio.h
 --------
 
 =============  =========  ============
 Function Name  Available  RPC Required
 =============  =========  ============
 puts           |check|    |check|
 fputs          |check|    |check|
+fclose         |check|    |check|
+fopen          |check|    |check|
 =============  =========  ============

libc/include/llvm-libc-types/rpc_opcodes_t.h

 #define __LLVM_LIBC_TYPES_RPC_OPCODE_H__
 
 typedef enum : unsigned short {
   RPC_NOOP = 0,
   RPC_EXIT = 1,
   RPC_WRITE_TO_STDOUT = 2,
   RPC_WRITE_TO_STDERR = 3,
   RPC_WRITE_TO_STREAM = 4,
-  RPC_MALLOC = 5,
-  RPC_FREE = 6,
-  RPC_TEST_INCREMENT = 7,
-  RPC_TEST_INTERFACE = 8,
-  RPC_TEST_STREAM = 9,
+  RPC_OPEN_FILE = 5,
+  RPC_CLOSE_FILE = 6,
+  RPC_MALLOC = 7,
+  RPC_FREE = 8,
+  // TODO: Move these out of here and handle then with custom handlers in the
+  // loader.
+  RPC_TEST_INCREMENT = 1000,
+  RPC_TEST_INTERFACE = 1001,
+  RPC_TEST_STREAM = 1002,
 } rpc_opcode_t;
 
 #endif // __LLVM_LIBC_TYPES_RPC_OPCODE_H__

libc/src/__support/File/gpu/file.cpp

 //===--- GPU specialization of the File data structure --------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #include "src/__support/File/file.h"
 
+#include "llvm-libc-types/rpc_opcodes_t.h"
 #include "src/__support/RPC/rpc_client.h"
 #include "src/errno/libc_errno.h" // For error macros
+#include "src/string/string_utils.h"
 
 #include <stdio.h>
 
 namespace __llvm_libc {
 
 namespace {
 
 FileIOResult write_func(File *, const void *, size_t);
+int close_func(File *);
 
 } // namespace
 
 class GPUFile : public File {
   uintptr_t file;
 
 public:
   constexpr GPUFile(uintptr_t file, File::ModeFlags modeflags)
-      : File(&write_func, nullptr, nullptr, nullptr, nullptr, 0, _IONBF, false,
-             modeflags),
+      : File(&write_func, nullptr, nullptr, &close_func, nullptr, 0, _IONBF,
+             false, modeflags),
         file(file) {}
 
   uintptr_t get_file() const { return file; }
 };
 
 namespace {
 
 int write_to_stdout(const void *data, size_t size) {
else if (gpu_file == stderr)
     ret = write_to_stderr(data, size);
   else
     ret = write_to_stream(gpu_file->get_file(), data, size);
   if (ret < 0)
     return {0, -ret};
   return ret;
 }
 
+int close_func(File *file) {
+  int ret = 0;
+  GPUFile *gpu_file = reinterpret_cast<GPUFile *>(file);
+  rpc::Client::Port port = rpc::client.open<RPC_CLOSE_FILE>();
+  port.send_and_recv(
+      [=](rpc::Buffer *buffer) { buffer->data[0] = gpu_file->get_file(); },
+      [&](rpc::Buffer *buffer) { ret = buffer->data[0]; });
+  port.close();
+
+  return ret;
+}
+
 } // namespace
 
+void *ptr;
+
+ErrorOr<File *> openfile(const char *path, const char *mode) {
+  auto modeflags = File::mode_flags(mode);
+  if (modeflags == 0)
+    return Error(EINVAL);
+
+  uintptr_t file;
+  rpc::Client::Port port = rpc::client.open<RPC_OPEN_FILE>();
+  port.send_n(path, internal::string_length(path) + 1);
+  port.send_and_recv(
+      [=](rpc::Buffer *buffer) {
+        inline_memcpy(buffer->data, mode, internal::string_length(mode) + 1);
+      },
+      [&](rpc::Buffer *buffer) { file = buffer->data[0]; });
+  port.close();
+
+  static GPUFile gpu_file(0, 0);
+  gpu_file = GPUFile(file, modeflags);
+  return &gpu_file;
+}
+
 static GPUFile StdIn(0UL, File::ModeFlags(File::OpenMode::READ));
 File *stdin = &StdIn;
 
 static GPUFile StdOut(0UL, File::ModeFlags(File::OpenMode::APPEND));
 File *stdout = &StdOut;
 
 static GPUFile StdErr(0UL, File::ModeFlags(File::OpenMode::APPEND));
 File *stderr = &StdErr;

libc/test/src/stdio/CMakeLists.txt

add_libc_test(
   SUITE
     libc_stdio_unittests
   SRCS
     fputs_test.cpp
   DEPENDS
     libc.src.stdio.fputs
 )
 
+add_libc_test(
+  fopen_test
+  SUITE
+    libc_stdio_unittests
+  SRCS
+    fopen_test.cpp
+  DEPENDS
+    libc.src.stdio.fputs
+    libc.src.stdio.fclose
+    libc.src.stdio.fopen
+)
+
 add_libc_unittest(
   putc_test
   SUITE
     libc_stdio_unittests
   SRCS
     putc_test.cpp
   DEPENDS
     libc.src.stdio.putc
DEPENDS
     libc.src.stdio.fopen
     libc.src.stdio.fread
     libc.src.stdio.fseek
     libc.src.stdio.ftell
     libc.src.stdio.fwrite
     libc.src.stdio.setvbuf
 )
 
+# Create an output directory for any temporary test files.
+file(MAKE_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/testdata)
+
 if(LIBC_TARGET_ARCHITECTURE_IS_GPU)
   return()
 endif()
 
 add_subdirectory(printf_core)
 add_subdirectory(scanf_core)
 add_subdirectory(testdata)

libc/test/src/stdio/fopen_test.cpp

+//===-- Unittests for fopen / fclose --------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/__support/File/file.h"
+#include "src/stdio/fclose.h"
+#include "src/stdio/fopen.h"
+#include "src/stdio/fputs.h"
+
+#include "test/UnitTest/Test.h"
+
+TEST(LlvmLibcFOpenTest, PrintToFile) {
+  int result;
+
+  FILE *file = __llvm_libc::fopen("./testdata/test_data.txt", "w");
+  ASSERT_FALSE(file == nullptr);
+
+  constexpr char another[] = "A simple string written to a file\n";
+  result = __llvm_libc::fputs(another, file);
+  EXPECT_GE(result, 0);
+
+  ASSERT_EQ(0, __llvm_libc::fclose(file));
+}

libc/utils/gpu/server/rpc_server.cpp

case RPC_WRITE_TO_STDOUT: {
         std::memcpy(buffer->data, &ret, sizeof(int));
       });
       for (uint64_t i = 0; i < rpc::MAX_LANE_SIZE; ++i) {
         if (strs[i])
           delete[] reinterpret_cast<uint8_t *>(strs[i]);
       }
       break;
     }
+    case RPC_OPEN_FILE: {
+      uint64_t sizes[rpc::MAX_LANE_SIZE] = {0};
+      void *paths[rpc::MAX_LANE_SIZE] = {nullptr};
+      port->recv_n(paths, sizes, [&](uint64_t size) { return new char[size]; });
+      port->recv_and_send([&](rpc::Buffer *buffer, uint32_t id) {
+        FILE *file = fopen(reinterpret_cast<char *>(paths[id]),
+                           reinterpret_cast<char *>(buffer->data));
+        buffer->data[0] = reinterpret_cast<uintptr_t>(file);
+      });
+      break;
+    }
+    case RPC_CLOSE_FILE: {
+      port->recv_and_send([&](rpc::Buffer *buffer, uint32_t id) {
+        FILE *file = reinterpret_cast<FILE *>(buffer->data[0]);
+        buffer->data[0] = fclose(file);
+      });
+      break;
+    }
     case RPC_EXIT: {
       // Send a response to the client to signal that we are ready to exit.
       port->recv_and_send([](rpc::Buffer *) {});
       port->recv([](rpc::Buffer *buffer) {
         int status = 0;
         std::memcpy(&status, buffer->data, sizeof(int));
         exit(status);
       });