1

I'm building a project with CUDA code, using a recent CMake which has intrinsic CUDA support (version >= 3.8 or later if necessary).

How do I tell CMake to (also) generate PTX files for my various kernels?

Something I've tried which doesn't (?) work:

Beginning with CMake 3.9, we can define an object library to have PTXes rather than kinds of objects, using the CUDA_PTX_COMPILATION property:

add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

However, this isn't a proper solution to the problem - something is missing. Suppose we have in a.cu:

__global__ void foo() { return; }

and in b.cu:

__global__ void bar(int* a) { *a = 5; }

and we run cmake and make with the following CMakeLists.txt:

cmake_minimum_required(VERSION 3.9)
add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

no PTX files get generated and nvcc isn't run. Not sure why.

einpoklum
  • 118,144
  • 57
  • 340
  • 684
  • 1
    @Tsyvarev: Moved the half-answer into the question. – einpoklum May 05 '19 at 19:14
  • What do you want the PTX for? Appropriately compiled binary payloads contain both PTX and SASS. Isn't that sufficient? – talonmies May 06 '19 at 07:27
  • @talonmies: I don't mind whether the PTX is extracted from the binary, aposteriori. I just want to inspect the PTX code as well as build the binaries. – einpoklum May 06 '19 at 08:21

2 Answers2

2

Try turning on CUDA language.

cmake_minimum_required(VERSION 3.9)
enable_language(CUDA)
add_library(myptx OBJECT a.cu b.cu)
set_property(TARGET myptx PROPERTY CUDA_PTX_COMPILATION ON)

PTXs are located in ${CMAKE_BINARY_DIR}/CMakeFiles/myptx.dir

zingdle
  • 491
  • 5
  • 17
1

You can pass the --keep command-line option to nvcc, which will make it keep all intermediate files in the same directory it's using for other outputs. With newer CMake versions, you would write:

target_compile_options(
    some_target
    another_target
    yet_another_target_etc
    PRIVATE 
    "--keep"
)

and that should do it. It seems the .ptx ends up where the binary is. But - note you'll also get a gazillion other files: .cpp4.ii, .cudafe1.c, .cudafe1.cpp, .fatbin and so on. So while this works, it's not such a great solution.

einpoklum
  • 118,144
  • 57
  • 340
  • 684