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.
Details
Diff Detail
- Repository
- rG LLVM Github Monorepo
Event Timeline
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.