Цель ROCm / HIP в cmake с несколькими исходными файлами

Question

Как мы можем скомпилировать один исполняемый файл из нескольких исходных файлов.cu?

CMakeLists.txt:

cmake_minimum_required(VERSION 3.0.0)
project(test VERSION 0.1 LANGUAGES CXX)


include_directories(${CMAKE_CURRENT_SOURCE_DIR})

####
## checking for HIP 
####

# setup HIP
if(NOT DEFINED ROCM_FOLDER_NAME)
    set(ROCM_FOLDER_NAME rocm)
endif()

if(NOT DEFINED HIP_PATH)
    if(NOT DEFINED ENV{HIP_PATH})
        set(HIP_PATH "/opt/${ROCM_FOLDER_NAME}/hip" CACHE PATH "Path to which HIP has been installed")
    else()
        set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed")
    endif()
endif()
set(hip_DIR "/opt/rocm/hip/lib/cmake/hip")
set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
set(CMAKE_MODULE_PATH "${HIP_PATH}/lib/cmake/hip" ${CMAKE_MODULE_PATH})
set(CMAKE_PREFIX_PATH "/opt/${ROCM_FOLDER_NAME}/lib/cmake/AMDDeviceLibs/" "/opt/${ROCM_FOLDER_NAME}/lib/cmake/amd_comgr/" "/opt/${ROCM_FOLDER_NAME}/lib/cmake/hsa-runtime64/" ${CMAKE_PREFIX_PATH})


find_package(HIP REQUIRED)
message(STATUS "Found HIP: " ${HIP_VERSION})


# list sources and names for targets
set(TARGET_SOURCE_FILES main.cu test.cu)
set(TARGET_NAME test_target)

# add HIP targets
set_source_files_properties(${TARGET_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
hip_add_executable(${TARGET_NAME} ${TARGET_SOURCE_FILES} HIPCC_OPTIONS ${HIPCC_OPTIONS} HCC_OPTIONS ${HCC_OPTIONS} NVCC_OPTIONS ${NVCC_OPTIONS})

test.h:

#pragma once 

template<typename T>
__host__ __device__ T addOne(const T t)
{
  return t + static_cast<T>(1.0);
}

extern template __host__ __device__ float addOne(const float t);

test.cu:

#include <test.h>
#include "hip/hip_runtime.h"
template __host__ __device__ float addOne(const float t);

main.cu:

#include <test.h>

#include "hip/hip_runtime.h"

#include <iostream>

template<typename T>
__global__ void myTestKernel(const T t)
{
  T t2 = addOne(t);
}


int main()
{
  std::cout << "Executing Kernel" << std::endl;
  const unsigned blocks = 512;
  const unsigned threadsPerBlock = 256;
  hipLaunchKernelGGL(myTestKernel, dim3(blocks), dim3(threadsPerBlock), 0, 0, 1.0f);
  hipDeviceSynchronize();
  std::cout << "All done!" << std::endl;
  return 0;
}

При попытке скомпилировать пример во время связывания я получаю следующее сообщение об ошибке:

lld: error: undefined hidden symbol: float addOne<float>(float)
>>> referenced by /tmp/main-gfx803-74431b.o:(void myTestKernel<float>(float))
>>> referenced by /tmp/main-gfx803-74431b.o:(void myTestKernel<float>(float))
clang-12: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)

Похоже, что компоновщик не рассматривает / не находит явные экземпляры шаблона в test.cu. Как добавить в компоновщик все исходные файлы?

Спасибо за помощь!