{"id":390,"date":"2022-05-04T03:47:43","date_gmt":"2022-05-04T01:47:43","guid":{"rendered":"https:\/\/threedots.ovh\/blog\/?p=390"},"modified":"2022-05-04T13:25:42","modified_gmt":"2022-05-04T11:25:42","slug":"gccs-nvptx-target","status":"publish","type":"post","link":"https:\/\/threedots.ovh\/blog\/2022\/05\/gccs-nvptx-target\/","title":{"rendered":"GCC’s NVPTX target"},"content":{"rendered":"\n

The GNU Compiler Collection<\/a> has an NVIDIA PTX target, supporting both offload modes for OpenMP<\/a> and<\/em> OpenACC<\/a> on NVIDIA GPUs.<\/p>\n\n\n\n

It also supports a quite special standalone<\/em> mode, for use cases that want to run a single-threaded<\/em> load on GPUs – this is of course not remotely performant by orders of magnitude – but is very useful for testing the compiler itself.<\/p>\n\n\n\n

Building it<\/h2>\n\n\n\n

Common<\/h3>\n\n\n\n

nvptx-tools<\/a> are the binutils equivalent for this GCC target.<\/p>\n\n\n\n

git clone https:\/\/github.com\/MentorEmbedded\/nvptx-tools\ncd nvptx-tools\nPATH=\/usr\/local\/cuda\/bin .\/configure --with-cuda-driver=\/usr\/local\/cuda\nmake\nsudo make install<\/code><\/pre>\n\n\n\n
Offload mode<\/h3>\n\n\n\n
This is the mode that you’ll most likely want to use. Offload mode is currently not functional on 64-bit Arm due to GCC bug 96265<\/a>.<\/p>\n\n\n\n
git clone git:\/\/gcc.gnu.org\/git\/gcc.git\ngit clone https:\/\/sourceware.org\/git\/newlib-cygwin.git\ncd gcc\nln -s ..\/newlib-cygwin\/newlib newlib\nmkdir -p build\/nvptx-offload\nmkdir -p build\/host\ncd build\/nvptx-offload\n..\/..\/configure --target=nvptx-none --enable-as-accelerator-for=$HOST_TRIPLET --disable-sjlj-exceptions --enable-newlib-io-long-long --with-build-time-tools=\/usr\/local\/nvptx-none\/bin --with-as=\/usr\/local\/bin\/nvptx-none-as\nmake -j160 && sudo make install DESTDIR=$PREFIX\ncd ..\/host\n..\/..\/configure --build=$HOST_TRIPLET --host=$HOST_TRIPLET --target=$HOST_TRIPLET --enable-languages=c,c++,fortran,lto --enable-offload-targets=nvptx-none=$PREFIX\/usr\/local\/nvptx-none --with-cuda-driver=\/usr\/local\/cuda\nmake -j160 && sudo make install DESTDIR=$PREFIX<\/code><\/pre>\n\n\n\n
Standalone mode<\/h3>\n\n\n\n
git clone git:\/\/gcc.gnu.org\/git\/gcc.git\ngit clone https:\/\/sourceware.org\/git\/newlib-cygwin.git\ncd gcc\nln -s ..\/newlib-cygwin\/newlib newlib\nmkdir -p build\/nvptx\ncd build\/nvptx\n..\/..\/configure --target=nvptx-none --disable-sjlj-exceptions --enable-newlib-io-long-long --with-build-time-tools=\/usr\/local\/nvptx-none\/bin  --with-as=\/usr\/local\/bin\/nvptx-none-as\nmake -j160 && sudo make install DESTDIR=$PREFIX<\/code><\/pre>\n\n\n\n
Fun with standalone mode<\/h2>\n\n\n\n
Standalone mode allows to have regular<\/em> programs run, very slowly (GPUs aren’t exactly known for good single SIMT lane performance) and with the limitations applicable to GCC’s PTX backend. That allows to have the GCC test suite run for example – but we can also have some fun with it.<\/p>\n\n\n\n
Let’s have a (totally useless by the way) test program to see how long a printf takes on hardware:<\/p>\n\n\n\n
#include <stdio.h>\n#include <time.h>\n#include <stdint.h>\n\nextern uint64_t clock64();\n\nint main (int argc, char** argv) {\n\tuint64_t b = clock64();\n\tputs(\"hello world from a GPU thread!\");\n\tuint64_t e = clock64();\n\tprintf(\"time elapsed: %i\\n\", e - b);\n}<\/code><\/pre>\n\n\n\n
and<\/p>\n\n\n\n
.version 6.0\n.target sm_30\n.address_size 64\n.visible .func (.param .u64 %value_out) clock64;\n.visible .func (.param .u64 %value_out) clock64\n{\n.reg .u64 %value;\n.reg .u64 %r23;\nmov.u64 %r23,%clock64;\nmov.u64 %value,%r23;\nst.param.u64 [%value_out],%value;\nret;\n}\n<\/code><\/pre>\n\n\n\n
We can compile the program and then run it:<\/p>\n\n\n\n
$ nvptx-none-gcc -mmainkernel -misa=sm_70 clock.s -O3 hello.c -o hello\n$ nvptx-none-run hello\nnvptx-none-run hello\nhello world from a GPU thread!\ntime elapsed: 48359<\/code><\/pre>\n\n\n\n
What does -mmainkernel do?<\/h3>\n\n\n\n
It adds a startup stub analogous to the one below to allow regular unmodified test cases to run.<\/p>\n\n\n\n
int main(int argc, char** argv);\n\n__attribute__((kernel))\nvoid __main(int *exitval_ptr, int argc, char *argv[])\n{\n  *exitval_ptr = main(argc, argv);\n}\n<\/code><\/pre>\n","protected":false},"excerpt":{"rendered":"
The GNU Compiler Collection has an NVIDIA PTX target, supporting both offload modes for OpenMP and OpenACC on NVIDIA GPUs. It also supports a quite special standalone mode, for use cases that want to run a single-threaded load on GPUs – this is of course not remotely performant by orders of magnitude – but is… Read More »GCC’s NVPTX target<\/span><\/a><\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"neve_meta_sidebar":"","neve_meta_container":"","neve_meta_enable_content_width":"","neve_meta_content_width":0,"neve_meta_title_alignment":"","neve_meta_author_avatar":"","neve_post_elements_order":"","neve_meta_disable_header":"","neve_meta_disable_footer":"","neve_meta_disable_title":"","footnotes":""},"categories":[1],"tags":[],"class_list":["post-390","post","type-post","status-publish","format-standard","hentry","category-uncategorized"],"_links":{"self":[{"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/posts\/390","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/comments?post=390"}],"version-history":[{"count":6,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/posts\/390\/revisions"}],"predecessor-version":[{"id":397,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/posts\/390\/revisions\/397"}],"wp:attachment":[{"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/media?parent=390"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/categories?post=390"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/threedots.ovh\/blog\/wp-json\/wp\/v2\/tags?post=390"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}