From 355eea887c2c5ffeb23c86f45e0a800ab060dd74 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Mon, 6 Mar 2017 09:41:40 -0500 Subject: [PATCH] Get the vtkm cont cuda object to compile properly. --- vtkm/cont/CMakeLists.txt | 18 +++++++--- .../internal/ArrayManagerExecutionCuda.cu | 2 +- .../cuda/internal/ArrayManagerExecutionCuda.h | 21 +++++------- vtkm/cont/cuda/internal/CMakeLists.txt | 33 ++++++++++++------- .../cuda/internal/ArrayPortalFromThrust.h | 7 ++-- 5 files changed, 48 insertions(+), 33 deletions(-) diff --git a/vtkm/cont/CMakeLists.txt b/vtkm/cont/CMakeLists.txt index 486a1eeac..2a323b926 100644 --- a/vtkm/cont/CMakeLists.txt +++ b/vtkm/cont/CMakeLists.txt @@ -109,10 +109,15 @@ set(sources ${sources} $) add_subdirectory(cuda) if (VTKm_ENABLE_CUDA) - # Need to come back to this at some point, but can't make OBJECT libraries - # with cuda_add_library: - # - # set(sources ${sources} $) + get_property(vtkm_cont_cuda_object_files GLOBAL + PROPERTY vtkm_cont_cuda_object_files ) + # mark the file as generated, this needs to be done in this directory + # for 'reasons'. + set_source_files_properties(${vtkm_cont_cuda_object_files} + PROPERTIES GENERATED TRUE) + + set(sources ${sources} ${vtkm_cont_cuda_object_files}) + endif() add_subdirectory(tbb) @@ -120,11 +125,16 @@ if (VTKm_ENABLE_TBB) set(sources ${sources} $) endif() + vtkm_library( SOURCES ${sources} WRAP_FOR_CUDA ${device_sources} HEADERS ${header_impls}) target_include_directories(vtkm_cont PRIVATE ${VTKm_BACKEND_INCLUDE_DIRS}) +if(VTKm_ENABLE_CUDA) + add_dependencies(vtkm_cont vtkm_cont_cuda) +endif() + #----------------------------------------------------------------------------- add_subdirectory(testing) diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu index 9b1962bdc..dc3371adc 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.cu @@ -18,7 +18,7 @@ // this software. //============================================================================ -#include "ArrayManagerExecutionCuda.h" +#include namespace vtkm { namespace cont { diff --git a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h index dcbdab359..590a6cdfb 100644 --- a/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h +++ b/vtkm/cont/cuda/internal/ArrayManagerExecutionCuda.h @@ -22,7 +22,6 @@ #include -#include #include #include #include @@ -118,18 +117,16 @@ public: } // namespace internal -// Disabled for now -- see comment in cont/cuda/internal/CMakeLists.txt about -// cuda_add_library and OBJECT libraries. -//#ifdef VTKM_BUILD_PREPARE_FOR_DEVICE -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(char, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int8, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::UInt8, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int32, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int64, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Float32, DeviceAdapterTagCuda) -//EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Float64, DeviceAdapterTagCuda) -//#endif // VTKM_BUILD_PREPARE_FOR_DEVICE +#ifdef VTKM_BUILD_PREPARE_FOR_DEVICE +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(char, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int8, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::UInt8, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int32, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Int64, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Float32, DeviceAdapterTagCuda) +EXPORT_ARRAYHANDLE_DEVICE_ADAPTER(vtkm::Float64, DeviceAdapterTagCuda) +#endif // VTKM_BUILD_PREPARE_FOR_DEVICE } } // namespace vtkm::cont diff --git a/vtkm/cont/cuda/internal/CMakeLists.txt b/vtkm/cont/cuda/internal/CMakeLists.txt index ee246d801..54d96296b 100644 --- a/vtkm/cont/cuda/internal/CMakeLists.txt +++ b/vtkm/cont/cuda/internal/CMakeLists.txt @@ -34,15 +34,26 @@ vtkm_declare_headers(CUDA ${headers} TESTABLE ${VTKm_ENABLE_CUDA}) if (VTKm_ENABLE_CUDA) add_subdirectory(testing) - # Need to sort this out still...cuda_add_library doesn't support OBJECT - # libraries: - # -# cuda_add_library(vtkm_cont_cuda OBJECT ArrayManagerExecutionCuda.cu) -# set_property(TARGET vtkm_cont_cuda PROPERTY POSITION_INDEPENDENT_CODE ON) -# target_compile_features(vtkm_cont_cuda PRIVATE cxx_auto_type) -# target_include_directories(vtkm_cont_cuda PRIVATE -# "${VTKm_SOURCE_DIR}" -# "${VTKm_BINARY_DIR}/include" -# ${VTKm_BACKEND_INCLUDE_DIRS} -# ) + #todo we need to add a custom target and feed that as a dependency + #for vtkm_cont + cuda_include_directories(${VTKm_SOURCE_DIR} + ${VTKm_BINARY_DIR}/include + ${VTKm_BACKEND_INCLUDE_DIRS} + ) + + set(compile_options ) + if(BUILD_SHARED_LIBS AND NOT WIN32) + list(APPEND compile_options -Xcompiler=${CMAKE_CXX_COMPILE_OPTIONS_VISIBILITY}hidden) + list(APPEND compile_options -Xcompiler=-fPIC) + endif() + + cuda_compile(vtkm_cont_cuda_object_files ArrayManagerExecutionCuda.cu + OPTIONS "${compile_options}") + + #Setup the dependency chain for the custom object build so that + add_custom_target(vtkm_cont_cuda DEPENDS ${vtkm_cont_cuda_object_files}) + + set_property( GLOBAL + PROPERTY vtkm_cont_cuda_object_files ${vtkm_cont_cuda_object_files}) + endif() diff --git a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h index 1ff69ec4d..bf5ceac9c 100644 --- a/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h +++ b/vtkm/exec/cuda/internal/ArrayPortalFromThrust.h @@ -226,7 +226,6 @@ struct load_through_texture( (this->EndIterator - this->BeginIterator) ); } - __host__ __device__ + VTKM_EXEC_CONT ValueType Get(vtkm::Id index) const { typedef typename ::thrust::iterator_traits::difference_type SizeType; return *(this->BeginIterator + static_cast(index)); } - __host__ __device__ + VTKM_EXEC_CONT void Set(vtkm::Id index, ValueType value) const { typedef typename ::thrust::iterator_traits::difference_type SizeType; @@ -343,12 +342,10 @@ public: } #else - __host__ ValueType Get(vtkm::Id vtkmNotUsed(index) ) const { return ValueType(); } - __host__ void Set(vtkm::Id vtkmNotUsed(index), ValueType vtkmNotUsed(value)) const { #if ! (defined(VTKM_MSVC) && defined(VTKM_CUDA)) VTKM_ASSERT(true && "Cannot set to const array.");