CUDA: Allow sources to be compiled to .ptx files

When the target property `CUDA_PTX_COMPILATION` is enabled CUDA OBJECT
libraries will generate ptx files instead of object files.

6 files changed, 145 insertions, 0 deletions

diff --git a/Tests/CudaOnly/CMakeLists.txt b/Tests/CudaOnly/CMakeLists.txt
index 85a2051..a3bd707 100644
--- a/Tests/CudaOnly/CMakeLists.txt
+++ b/Tests/CudaOnly/CMakeLists.txt
@@ -1,4 +1,5 @@
 
 ADD_TEST_MACRO(CudaOnly.EnableStandard CudaOnlyEnableStandard)
+ADD_TEST_MACRO(CudaOnly.ExportPTX CudaOnlyExportPTX)
 ADD_TEST_MACRO(CudaOnly.SeparateCompilation CudaOnlySeparateCompilation)
 ADD_TEST_MACRO(CudaOnly.WithDefs CudaOnlyWithDefs)
diff --git a/Tests/CudaOnly/ExportPTX/CMakeLists.txt b/Tests/CudaOnly/ExportPTX/CMakeLists.txt
new file mode 100644
index 0000000..10249c6
--- /dev/null
+++ b/Tests/CudaOnly/ExportPTX/CMakeLists.txt
@@ -0,0 +1,82 @@
+cmake_minimum_required(VERSION 3.8)
+project (CudaOnlyExportPTX CUDA)
+
+#Goal for this example:
+# How to generate PTX files instead of OBJECT files
+# How to reference PTX files for custom commands
+# How to install PTX files
+
+add_library(CudaPTX OBJECT kernelA.cu kernelB.cu)
+set_property(TARGET CudaPTX PROPERTY CUDA_PTX_COMPILATION ON)
+
+#Test ObjectFiles with file(GENERATE)
+file(GENERATE
+     OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/gen_$<LOWER_CASE:$<CONFIG>/>path_to_objs.h
+     CONTENT [[
+
+#include <vector>
+#include <string>
+
+#ifndef path_to_objs
+#define path_to_objs
+
+static std::string ptx_paths = "$<TARGET_OBJECTS:CudaPTX>";
+
+#endif
+
+]]
+)
+#We are going to need a wrapper around bin2c for multiple reasons
+# 1. bin2c only converts a single file at a time
+# 2. bin2c has only standard out support, so we have to manually
+# redirect to a cmake buffer
+# 3. We want to pack everything into a single output file, so we
+# need to also pass the --name option
+set(output_file ${CMAKE_CURRENT_BINARY_DIR}/embedded_objs.h)
+
+get_filename_component(cuda_compiler_bin "${CMAKE_CUDA_COMPILER}" DIRECTORY)
+find_program(bin_to_c
+  NAMES bin2c
+  PATHS ${cuda_compiler_bin}
+  )
+if(NOT bin_to_c)
+  message(FATAL_ERROR
+    "bin2c not found:\n"
+    "  CMAKE_CUDA_COMPILER='${CMAKE_CUDA_COMPILER}'\n"
+    "  cuda_compiler_bin='${cuda_compiler_bin}'\n"
+    )
+endif()
+
+add_custom_command(
+  OUTPUT "${output_file}"
+  COMMAND ${CMAKE_COMMAND}
+    "-DBIN_TO_C_COMMAND=${bin_to_c}"
+    "-DOBJECTS=$<TARGET_OBJECTS:CudaPTX>"
+    "-DOUTPUT=${output_file}"
+    -P ${CMAKE_CURRENT_SOURCE_DIR}/bin2c_wrapper.cmake
+  VERBATIM
+  DEPENDS $<TARGET_OBJECTS:CudaPTX>
+  COMMENT "Converting Object files to a C header"
+  )
+
+add_executable(CudaOnlyExportPTX main.cu ${output_file})
+add_dependencies(CudaOnlyExportPTX CudaPTX)
+target_include_directories(CudaOnlyExportPTX PRIVATE
+                           ${CMAKE_CURRENT_BINARY_DIR} )
+target_compile_definitions(CudaOnlyExportPTX PRIVATE
+                           "CONFIG_TYPE=gen_$<LOWER_CASE:$<CONFIG>>")
+
+if(APPLE)
+  # We need to add the default path to the driver (libcuda.dylib) as an rpath, so that
+  # the static cuda runtime can find it at runtime.
+  target_link_libraries(CudaOnlyExportPTX PRIVATE -Wl,-rpath,/usr/local/cuda/lib)
+endif()
+
+#Verify that we can install object targets properly
+install(TARGETS CudaPTX CudaOnlyExportPTX
+   EXPORT cudaPTX
+   RUNTIME DESTINATION bin
+   LIBRARY DESTINATION lib
+   OBJECTS DESTINATION objs
+   )
+install(EXPORT cudaPTX DESTINATION lib/cudaPTX)
diff --git a/Tests/CudaOnly/ExportPTX/bin2c_wrapper.cmake b/Tests/CudaOnly/ExportPTX/bin2c_wrapper.cmake
new file mode 100644
index 0000000..0baf934
--- /dev/null
+++ b/Tests/CudaOnly/ExportPTX/bin2c_wrapper.cmake
@@ -0,0 +1,19 @@
+
+set(file_contents)
+foreach(obj ${OBJECTS})
+  get_filename_component(obj_ext ${obj} EXT)
+  get_filename_component(obj_name ${obj} NAME_WE)
+  get_filename_component(obj_dir ${obj} DIRECTORY)
+
+  if(obj_ext MATCHES ".ptx")
+    set(args --name ${obj_name} ${obj})
+    execute_process(COMMAND "${BIN_TO_C_COMMAND}" ${args}
+                    WORKING_DIRECTORY ${obj_dir}
+                    RESULT_VARIABLE result
+                    OUTPUT_VARIABLE output
+                    ERROR_VARIABLE error_var
+                    )
+    set(file_contents "${file_contents} \n${output}")
+  endif()
+endforeach()
+file(WRITE "${OUTPUT}" "${file_contents}")
diff --git a/Tests/CudaOnly/ExportPTX/kernelA.cu b/Tests/CudaOnly/ExportPTX/kernelA.cu
new file mode 100644
index 0000000..fbe0d26
--- /dev/null
+++ b/Tests/CudaOnly/ExportPTX/kernelA.cu
@@ -0,0 +1,7 @@
+
+__global__ void kernelA(float* r, float* x, float* y, float* z, int size)
+{
+  for (int i = threadIdx.x; i < size; i += blockDim.x) {
+    r[i] = x[i] * y[i] + z[i];
+  }
+}
diff --git a/Tests/CudaOnly/ExportPTX/kernelB.cu b/Tests/CudaOnly/ExportPTX/kernelB.cu
new file mode 100644
index 0000000..11872e4
--- /dev/null
+++ b/Tests/CudaOnly/ExportPTX/kernelB.cu
@@ -0,0 +1,8 @@
+
+
+__global__ void kernelB(float* r, float* x, float* y, float* z, int size)
+{
+  for (int i = threadIdx.x; i < size; i += blockDim.x) {
+    r[i] = x[i] * y[i] + z[i];
+  }
+}
diff --git a/Tests/CudaOnly/ExportPTX/main.cu b/Tests/CudaOnly/ExportPTX/main.cu
new file mode 100644
index 0000000..132377c
--- /dev/null
+++ b/Tests/CudaOnly/ExportPTX/main.cu
@@ -0,0 +1,28 @@
+
+#include <iostream>
+
+/*
+  Define GENERATED_HEADER macro to allow c++ files to include headers
+  generated based on different configuration types.
+*/
+
+/* clang-format off */
+#define GENERATED_HEADER(x) GENERATED_HEADER0(CONFIG_TYPE/x)
+/* clang-format on */
+#define GENERATED_HEADER0(x) GENERATED_HEADER1(x)
+#define GENERATED_HEADER1(x) <x>
+
+#include GENERATED_HEADER(path_to_objs.h)
+
+#include "embedded_objs.h"
+
+int main(int argc, char** argv)
+{
+  (void)argc;
+  (void)argv;
+
+  unsigned char* ka = kernelA;
+  unsigned char* kb = kernelB;
+
+  return (ka != NULL && kb != NULL) ? 0 : 1;
+}