[PATCH] D67867: [libc] Add few docs and implementation of strcpy and strcat.

[Siva Chandra via Phabricator via llvm-commits]

sivachandra added a comment.

In D67867#1686056 <https://reviews.llvm.org/D67867#1686056>, @hfinkel wrote:

> Maybe everything is fine, but given this setup, does anyone see any potential problems with compiling these functions for nvptx? I'd like to eventually see a mode where we compile an appropriate subset of these functions for GPU targets -- either in bitcode form for linking with our device-side OpenMP runtime library or as a native object -- to provide a more feature-complete offloading environment.
>
> The one thing that caught by eye was the use of the section attribute, and I was curious what nvptx does with that. As far as I can tell, perhaps the answer is nothing. I took this:
>
>   cat /tmp/f.c 
>   int __attribute__((section(".llvm.libc.entrypoint.foo"))) foo ()  {
>     return 0;
>   }
>  
>   clang -target nvptx64 -O3 -S -o - /tmp/f.c
>   //
>   // Generated by LLVM NVPTX Back-End
>   //
>    
>   .version 3.2
>   .target sm_20
>   .address_size 64
>    
>     // .globl foo
>    
>   .visible .func  (.param .b32 func_retval0) foo()
>   {
>     .reg .b32   %r<2>;
>    
>     mov.u32   %r1, 0;
>     st.param.b32  [func_retval0+0], %r1;
>     ret;
>    
>   }
>   
>
> so maybe that part will just work?


The way it is setup, it is very ELF-ish (or should I be saying elvish :)

Very likely, non-ELF systems like Windows and nvptx will require some other mechanism. I working on the windows story in parallel, but I am totally new to nvptx. If you have any suggestions on how to make it work there, we can surely incorporate it.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D67867/new/

https://reviews.llvm.org/D67867