Add support for unity builds with CUDA sources

Fixes #22750

Help/prop_tgt/UNITY_BUILD.rst

@@ -38,6 +38,9 @@ Unity builds are supported for the following languages:
 ``CXX``
   .. versionadded:: 3.16
 
+``CUDA``
+  .. versionadded:: 3.31
+
 ``OBJC``
   .. versionadded:: 3.29
 

Help/release/dev/unity-build-supports-cuda.rst

@@ -0,0 +1,5 @@
+unity-build-supports-cuda
+-------------------------
+
+* The :prop_tgt:`UNITY_BUILD` target property now supports the
+  CUDA (``CUDA``) language.

Source/cmLocalGenerator.cxx

@@ -3162,6 +3162,8 @@ std::string unity_file_extension(std::string const& lang)
     extension = "_c.c";
   } else if (lang == "CXX") {
     extension = "_cxx.cxx";
+  } else if (lang == "CUDA") {
+    extension = "_cu.cu";
   } else if (lang == "OBJC") {
     extension = "_m.m";
   } else if (lang == "OBJCXX") {
@@ -3291,7 +3293,7 @@ void cmLocalGenerator::AddUnityBuild(cmGeneratorTarget* target)
   cmValue afterInclude = target->GetProperty("UNITY_BUILD_CODE_AFTER_INCLUDE");
   cmValue unityMode = target->GetProperty("UNITY_BUILD_MODE");
 
-  for (std::string lang : { "C", "CXX", "OBJC", "OBJCXX" }) {
+  for (std::string lang : { "C", "CXX", "OBJC", "OBJCXX", "CUDA" }) {
     std::vector<UnityBatchedSource> filtered_sources;
     std::copy_if(unitySources.begin(), unitySources.end(),
                  std::back_inserter(filtered_sources),

Source/cmVisualStudio10TargetGenerator.cxx

@@ -2686,6 +2686,12 @@ void cmVisualStudio10TargetGenerator::WriteAllSources(Elem& e0)
           }
         }
       }
+      if (haveUnityBuild && strcmp(tool, "CudaCompile") == 0 &&
+          si.Source->GetProperty("UNITY_SOURCE_FILE")) {
+        if (!si.Source->GetPropertyAsBool("SKIP_UNITY_BUILD_INCLUSION")) {
+          exclude_configs = all_configs;
+        }
+      }
 
       if (si.Kind == cmGeneratorTarget::SourceKindObjectSource ||
           si.Kind == cmGeneratorTarget::SourceKindUnityBatched) {

Tests/CudaOnly/CMakeLists.txt

@@ -19,6 +19,7 @@ add_cuda_test_macro(CudaOnly.Toolkit CudaOnlyToolkit)
 add_cuda_test_macro(CudaOnly.ToolkitBeforeLang CudaOnlyToolkitBeforeLang)
 add_cuda_test_macro(CudaOnly.ToolkitMultipleDirs CudaOnlyToolkitMultipleDirs)
 add_cuda_test_macro(CudaOnly.TryCompileTargetStatic CudaOnlyTryCompileTargetStatic)
+add_cuda_test_macro(CudaOnly.Unity CudaOnlyUnity)
 add_cuda_test_macro(CudaOnly.WithDefs CudaOnlyWithDefs)
 add_cuda_test_macro(CudaOnly.CircularLinkLine CudaOnlyCircularLinkLine)
 add_cuda_test_macro(CudaOnly.ResolveDeviceSymbols CudaOnlyResolveDeviceSymbols)

Tests/CudaOnly/Unity/CMakeLists.txt

@@ -0,0 +1,14 @@
+cmake_minimum_required(VERSION 3.30)
+project(Unity CUDA)
+
+set(CMAKE_UNITY_BUILD 1)
+
+add_library(UnityObjects STATIC a.cu b.cu)
+
+add_executable(CudaOnlyUnity main.cu)
+target_link_libraries(CudaOnlyUnity PRIVATE UnityObjects)
+
+if(APPLE)
+  # Help the static cuda runtime find the driver (libcuda.dyllib) at runtime.
+  set_property(TARGET CudaOnlyUnity PROPERTY BUILD_RPATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES})
+endif()

Tests/CudaOnly/Unity/a.cu

@@ -0,0 +1,12 @@
+
+__device__ int function(int a, int b);
+
+__global__ void kernel()
+{
+  function(2, 3);
+}
+
+void test_unity_functions()
+{
+  kernel<<<1, 1, 1>>>();
+}

Tests/CudaOnly/Unity/b.cu

@@ -0,0 +1,5 @@
+
+__device__ int function(int a, int b)
+{
+  return b * b * a + 1;
+}

Tests/CudaOnly/Unity/main.cu

@@ -0,0 +1,8 @@
+
+void test_unity_functions();
+
+int main(int argc, char** argv)
+{
+  test_unity_functions();
+  return 0;
+}

Tests/RunCMake/CMakeLists.txt

@@ -1201,7 +1201,10 @@ add_RunCMake_test(PrecompileHeaders -DCMAKE_C_COMPILER_ID=${CMAKE_C_COMPILER_ID}
   -DCMAKE_C_SIMULATE_ID=${CMAKE_C_SIMULATE_ID}
   -DCMAKE_C_COMPILER_VERSION=${CMAKE_C_COMPILER_VERSION})
 
-add_RunCMake_test(UnityBuild -DCMake_TEST_OBJC=${CMake_TEST_OBJC})
+add_RunCMake_test(UnityBuild -DCMake_TEST_OBJC=${CMake_TEST_OBJC} -DCMake_TEST_CUDA=${CMake_TEST_CUDA})
+set_property(TEST RunCMake.UnityBuild APPEND
+  PROPERTY LABELS "CUDA")
+
 add_RunCMake_test(CMakePresets
   -DPython_EXECUTABLE=${Python_EXECUTABLE}
   -DCMake_TEST_JSON_SCHEMA=${CMake_TEST_JSON_SCHEMA}

Tests/RunCMake/UnityBuild/RunCMakeTest.cmake

@@ -21,6 +21,11 @@ if(CMake_TEST_OBJC)
   run_cmake(unitybuild_objcxx_group)
   run_cmake(unitybuild_c_and_cxx_and_objc_and_objcxx)
 endif()
+if(CMake_TEST_CUDA)
+  run_cmake(unitybuild_cuda)
+  run_cmake(unitybuild_cuda_group)
+  run_cmake(unitybuild_cxx_and_cuda)
+endif()
 run_cmake(unitybuild_batchsize)
 run_cmake(unitybuild_default_batchsize)
 run_cmake(unitybuild_skip)

Tests/RunCMake/UnityBuild/f.cu

@@ -0,0 +1,5 @@
+int f(int x)
+{
+  (void)x;
+  return 0;
+}

Tests/RunCMake/UnityBuild/unitybuild_cuda-check.cmake

@@ -0,0 +1,5 @@
+set(unitybuild_cu "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_0_cu.cu")
+if(NOT EXISTS "${unitybuild_cu}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_cu} does not exist.")
+  return()
+endif()

Tests/RunCMake/UnityBuild/unitybuild_cuda.cmake

@@ -0,0 +1,14 @@
+
+set(CMAKE_CUDA_ARCHITECTURES all-major)
+project(unitybuild_cuda CUDA)
+
+set(srcs "")
+foreach(s RANGE 1 8)
+  set(src "${CMAKE_CURRENT_BINARY_DIR}/s${s}.cu")
+  file(WRITE "${src}" "int s${s}(void) { return 0; }\n")
+  list(APPEND srcs "${src}")
+endforeach()
+
+add_library(tgt SHARED ${srcs})
+
+set_target_properties(tgt PROPERTIES UNITY_BUILD ON)

Tests/RunCMake/UnityBuild/unitybuild_cuda_group-check.cmake

@@ -0,0 +1,27 @@
+set(unitybuild_a_cu "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_a_cu.cu")
+if(NOT EXISTS "${unitybuild_a_cu}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_a_cu} does not exist.")
+  return()
+else()
+  #verify that odr2 is not part of this source set
+  file(STRINGS ${unitybuild_a_cu} unitybuild_a_cu_strings)
+  string(REGEX MATCH ".*#include.*odr2.cu" matched_code ${unitybuild_a_cu_strings})
+  if(matched_code)
+    set(RunCMake_TEST_FAILED "Generated unity file includes un-expected ord2.cu source file")
+    return()
+  endif()
+endif()
+
+set(unitybuild_b_cu "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_b_cu.cu")
+if(NOT EXISTS "${unitybuild_b_cu}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_b_cu} does not exist.")
+  return()
+else()
+  #verify that odr1 is not part of this source set
+  file(STRINGS ${unitybuild_b_cu} unitybuild_b_cu_strings)
+  string(REGEX MATCH ".*#include.*odr1.cu" matched_code ${unitybuild_b_cu_strings})
+  if(matched_code)
+    set(RunCMake_TEST_FAILED "Generated unity file includes un-expected ord1.cu source file")
+    return()
+  endif()
+endif()

Tests/RunCMake/UnityBuild/unitybuild_cuda_group.cmake

@@ -0,0 +1,28 @@
+set(CMAKE_CUDA_ARCHITECTURES all-major)
+project(unitybuild_cu CUDA)
+
+set(srcs "")
+foreach(s RANGE 1 4)
+  set(src "${CMAKE_CURRENT_BINARY_DIR}/s${s}.cu")
+  file(WRITE "${src}" "int s${s}(void) { return 0; }\n")
+  list(APPEND srcs "${src}")
+endforeach()
+
+foreach(s RANGE 1 2)
+  set(src "${CMAKE_CURRENT_BINARY_DIR}/odr${s}.cu")
+  file(WRITE "${src}" "namespace odr { int s${s}(void) { return 0; } }\n")
+  list(APPEND srcs "${src}")
+endforeach()
+
+add_library(tgt SHARED ${srcs})
+
+set_target_properties(tgt PROPERTIES UNITY_BUILD ON
+                          UNITY_BUILD_MODE GROUP
+                          )
+
+set_source_files_properties(s1.cu s2.cu odr1.cu
+                            PROPERTIES UNITY_GROUP "a"
+                            )
+set_source_files_properties(s3.cu s4.cu odr2.cu
+                            PROPERTIES UNITY_GROUP "b"
+                            )

Tests/RunCMake/UnityBuild/unitybuild_cxx_and_cuda-check.cmake

@@ -0,0 +1,42 @@
+set(unitybuild_a_cu "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_a_cu.cu")
+if(NOT EXISTS "${unitybuild_a_cu}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_a_cu} does not exist.")
+  return()
+else()
+  #verify that the 4 c file is part of this grouping and therefore UNITY_BUILD_BATCH_SIZE
+  #was ignored
+  file(STRINGS ${unitybuild_a_cu} unitybuild_a_c_strings)
+  string(REGEX MATCH ".*#include.*s1.cu.*#include.*s2.cu.*#include.*s3.cu.*#include.*s4.cu.*" matched_code ${unitybuild_a_c_strings})
+  if(NOT matched_code)
+    set(RunCMake_TEST_FAILED "Generated unity file doesn't include expected source files")
+    return()
+  endif()
+endif()
+
+set(unitybuild_b_cu "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_b_cu.cu")
+if(NOT EXISTS "${unitybuild_b_cu}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_b_cu} does not exist.")
+  return()
+endif()
+
+
+set(unitybuild_a_cxx "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_a_cxx.cxx")
+if(NOT EXISTS "${unitybuild_a_cxx}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_a_cxx} does not exist.")
+  return()
+else()
+  #verify that the 4 cxx file are part of this grouping and therefore UNITY_BUILD_BATCH_SIZE
+  #was ignored
+  file(STRINGS ${unitybuild_a_cxx} unitybuild_a_cxx_strings)
+  string(REGEX MATCH ".*#include.*s1.cxx.*#include.*s2.cxx.*#include.*s3.cxx.*#include.*s4.cxx.*" matched_code ${unitybuild_a_cxx_strings})
+  if(NOT matched_code)
+    set(RunCMake_TEST_FAILED "Generated unity file doesn't include expected source files")
+    return()
+  endif()
+endif()
+
+set(unitybuild_b_cxx "${RunCMake_TEST_BINARY_DIR}/CMakeFiles/tgt.dir/Unity/unity_b_cxx.cxx")
+if(NOT EXISTS "${unitybuild_b_cxx}")
+  set(RunCMake_TEST_FAILED "Generated unity source files ${unitybuild_b_cxx} does not exist.")
+  return()
+endif()

Tests/RunCMake/UnityBuild/unitybuild_cxx_and_cuda.cmake

@@ -0,0 +1,40 @@
+set(CMAKE_CUDA_ARCHITECTURES all-major)
+project(unitybuild_c_and_cxx CUDA CXX)
+
+set(srcs f.cu)
+foreach(s RANGE 1 8)
+  set(src_cu "${CMAKE_CURRENT_BINARY_DIR}/s${s}.cu")
+  file(WRITE "${src_cu}" "
+int f(int);\n
+int s${s}(void) { return f(${s}); }\n"
+  )
+
+  set(src_cxx "${CMAKE_CURRENT_BINARY_DIR}/s${s}.cxx")
+  file(WRITE "${src_cxx}" "
+extern \"C\" { \n
+  int f(int); \n
+}\n
+  int s${s}(void) { return f(${s}); }\n"
+  )
+
+  list(APPEND srcs "${src_cu}")
+  list(APPEND srcs "${src_cxx}")
+endforeach()
+
+
+
+add_library(tgt SHARED ${srcs})
+
+set_target_properties(tgt PROPERTIES UNITY_BUILD ON
+                          UNITY_BUILD_MODE GROUP
+                          #UNITY_BUILD_BATCH_SIZE will be ignored
+                          UNITY_BUILD_BATCH_SIZE 2)
+
+set_source_files_properties(s1.cu s2.cu s3.cu s4.cu
+                            s1.cxx s2.cxx s3.cxx s4.cxx
+                            PROPERTIES UNITY_GROUP "a"
+                            )
+set_source_files_properties(s5.cu s6.cu s7.cu s8.cu
+                            s5.cxx s6.cxx s7.cxx s8.cxx
+                            PROPERTIES UNITY_GROUP "b"
+                            )