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

[Hal Finkel via Phabricator via llvm-commits]

hfinkel added a comment.

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?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D67867