Add Concurrency, support long trajectories, Minor fixes

CUDA/CMakeLists.txt

@@ -94,14 +94,15 @@ endif(GPU_DOUBLE_PREC)
 
 SET(FERMI_GPU OFF CACHE BOOL "Enable build for (old) Fermi architectures (Compute capability 2.0)")
 
-
+set(MY_NVCC_FLAGS -I${CUDA_INCLUDE_DIRS})
+set(CMAKE_CXX_FLAGS -I${CUDA_INCLUDE_DIRS})
 IF(FERMI_GPU)
-  set(MY_NVCC_FLAGS -gencode arch=compute_30,code=sm_30)
+  list(APPEND MY_NVCC_FLAGS -gencode arch=compute_30,code=sm_30)
   list(APPEND MY_NVCC_FLAGS -gencode arch=compute_50,code=sm_50)
   list(APPEND MY_NVCC_FLAGS -gencode=arch=compute_52,code=sm_52)
   list(APPEND MY_NVCC_FLAGS -gencode=arch=compute_52,code=compute_52)
+  list(APPEND MY_NVCC_FLAGS -gencode arch=compute_50,code=sm_50)
 ELSE(FERMI_GPU)
-  set(MY_NVCC_FLAGS -gencode arch=compute_50,code=sm_50)
   list(APPEND MY_NVCC_FLAGS -gencode=arch=compute_52,code=sm_52)
   list(APPEND MY_NVCC_FLAGS -gencode=arch=compute_52,code=compute_52)
 
@@ -126,7 +127,7 @@ ENDIF(FERMI_GPU)
 
 IF(CMAKE_BUILD_TYPE MATCHES Debug)
   MESSAGE("debug mode")
-  list(APPEND CUDA_NVCC_FLAGS ${MY_NVCC_FLAGS} --ptxas-options=-v)
+  list(APPEND CUDA_NVCC_FLAGS ${MY_NVCC_FLAGS} --ptxas-options=-v -G)
   SET(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -g -std=c++11")
 ELSE(CMAKE_BUILD_TYPE)
   list(APPEND CUDA_NVCC_FLAGS ${MY_NVCC_FLAGS})
@@ -150,6 +151,7 @@ CONFIGURE_FILE( ${CMAKE_SOURCE_DIR}/inc/cufft_config.hpp.cmake ${CMAKE_SOURCE_DI
 
 #Include dirs
 include_directories(inc)
+message(CUDA_INCLUDE_DIRS : ${CUDA_INCLUDE_DIRS})
 SET(GPUNUFFT_INC_DIR ${CMAKE_SOURCE_DIR}/inc)
 SET(GPUNUFFT_INCLUDE ${GPUNUFFT_INC_DIR}/cuda_utils.hpp 
 										 ${GPUNUFFT_INC_DIR}/cuda_utils.cuh
@@ -161,10 +163,8 @@ SET(GPUNUFFT_INCLUDE ${GPUNUFFT_INC_DIR}/cuda_utils.hpp
 										 ${GPUNUFFT_INC_DIR}/precomp_utils.hpp
                      ${GPUNUFFT_INC_DIR}/gpuNUFFT_operator.hpp
 										 ${GPUNUFFT_INC_DIR}/balanced_operator.hpp
-										 ${GPUNUFFT_INC_DIR}/texture_gpuNUFFT_operator.hpp
 										 ${GPUNUFFT_INC_DIR}/balanced_gpuNUFFT_operator.hpp
-                     ${GPUNUFFT_INC_DIR}/gpuNUFFT_operator_factory.hpp
-										 ${GPUNUFFT_INC_DIR}/balanced_texture_gpuNUFFT_operator.hpp)
+                     ${GPUNUFFT_INC_DIR}/gpuNUFFT_operator_factory.hpp)
 
 SET(MATLAB_HELPER_INCLUDE ${GPUNUFFT_INC_DIR}/matlab_helper.h)
 SET(CONFIG_INCLUDE ${GPUNUFFT_INC_DIR}/config.hpp ${GPUNUFFT_INC_DIR}/cufft_config.hpp)

CUDA/inc/balanced_gpuNUFFT_operator.hpp

@@ -29,7 +29,6 @@ class BalancedGpuNUFFTOperator : public GpuNUFFTOperator,
 
   ~BalancedGpuNUFFTOperator()
   {
-    if (!matlabSharedMem)
       freeLocalMemberArray(this->sectorProcessingOrder.data);
   }
 

CUDA/inc/config.hpp.cmake

@@ -29,8 +29,8 @@
   typedef cufftComplex CufftType;
 #endif
 
-typedef unsigned int SizeType;
-typedef unsigned int IndType;
+typedef unsigned long int SizeType;
+typedef unsigned long int IndType;
 typedef uint2 IndType2;
 typedef uint3 IndType3;
 

CUDA/inc/cuda_utils.cuh

@@ -7,70 +7,6 @@ __constant__ gpuNUFFT::GpuNUFFTInfo GI;
 
 __constant__ DType KERNEL[10000];
 
-texture<float, 1, cudaReadModeElementType> texKERNEL;
-texture<float, 2, cudaReadModeElementType> texKERNEL2D;
-texture<float, 3, cudaReadModeElementType> texKERNEL3D;
-
-texture<float2> texDATA;
-texture<cufftComplex> texGDATA;
-
-__inline__ __device__ float compute1DTextureLookup(float x, float y)
-{
-  return tex1D(texKERNEL, x) * tex1D(texKERNEL, y);
-}
-
-__inline__ __device__ float compute1DTextureLookup(float x, float y, float z)
-{
-  return tex1D(texKERNEL, x) * tex1D(texKERNEL, y) * tex1D(texKERNEL, z);
-}
-
-__inline__ __device__ float compute2DTextureLookup(float x, float y)
-{
-  return (float)tex2D(texKERNEL2D, (float)x, (float)y);
-}
-
-__inline__ __device__ float compute2DTextureLookup(float x, float y, float z)
-{
-  return (float)tex2D(texKERNEL2D, (float)x, (float)y) *
-         tex2D(texKERNEL2D, (float)z, 0);
-}
-
-__inline__ __device__ float compute3DTextureLookup(float x, float y)
-{
-  return tex3D(texKERNEL3D, x, y, 0);
-}
-
-__inline__ __device__ float compute3DTextureLookup(float x, float y, float z)
-{
-  return tex3D(texKERNEL3D, x, y, z);
-}
-
-__inline__ __device__ float computeTextureLookup(float x, float y)
-{
-  // wired to 2d
-  return compute2DTextureLookup((float)x, (float)y);
-  // switch(GI.interpolationType)
-  //{
-  //  case 1: return compute1DTextureLookup(x,y);
-  //  case 2: return compute2DTextureLookup(x,y);
-  //  case 3: return compute3DTextureLookup(x,y);
-  //  default: return (float)0.0;
-  //}
-}
-
-__inline__ __device__ float computeTextureLookup(float x, float y, float z)
-{
-  // wired to 2d
-  return compute2DTextureLookup(x, y, z);
-  // switch(GI.interpolationType)
-  //{
-  //  case 1: return compute1DTextureLookup(x,y,z);
-  //  case 2: return compute2DTextureLookup(x,y,z);
-  //  case 3: return compute3DTextureLookup(x,y,z);
-  //  default: return (float)0.0;
-  //}
-}
-
 #if __CUDA_ARCH__ < 200
 #define THREAD_BLOCK_SIZE 256
 #else

CUDA/inc/cuda_utils.hpp

@@ -59,7 +59,19 @@ inline void copyToDevice(TypeName *host_ptr, TypeName *device_ptr,
   HANDLE_ERROR(cudaMemcpy(device_ptr, host_ptr, num_elements * sizeof(TypeName),
                           cudaMemcpyHostToDevice));
 }
-
+/** \brief CUDA memcpy call to copy data from host to device
+ *
+ * @param host_ptr      host data pointer
+ * @param device_ptr    device pointer
+ * @param num_elements  amount of elements of size TypeName
+ */
+template <typename TypeName>
+inline void copyToDeviceAsync(TypeName *host_ptr, TypeName *device_ptr,
+                         IndType num_elements, cudaStream_t stream=0)
+{
+  HANDLE_ERROR(cudaMemcpyAsync(device_ptr, host_ptr, num_elements * sizeof(TypeName),
+                          cudaMemcpyHostToDevice, stream));
+}
 /** \brief CUDA memory allocation and memcpy call to copy data from host to
  *device
  *
@@ -98,13 +110,28 @@ inline void allocateAndSetMem(TypeName **device_ptr, IndType num_elements,
  */
 template <typename TypeName>
 inline void copyDeviceToDevice(TypeName *device_ptr_src,
-                               TypeName *device_ptr_dest, IndType num_elements)
+                               TypeName *device_ptr_dest, IndType num_elements
+                               )
 {
   HANDLE_ERROR(cudaMemcpy(device_ptr_dest, device_ptr_src,
                           num_elements * sizeof(TypeName),
                           cudaMemcpyDeviceToDevice));
 }
 
+/** \brief CUDA memcpy call to copy data from device ptr to device ptr
+ *
+ * @param device_ptr_src   source device pointer
+ * @param device_ptr_dest  destination device pointer
+ * @param num_elements     amount of elements of size TypeName
+ */
+template <typename TypeName>
+inline void copyDeviceToDeviceAsync(TypeName *device_ptr_src,
+                               TypeName *device_ptr_dest, IndType num_elements, cudaStream_t stream=0)
+{
+  HANDLE_ERROR(cudaMemcpyAsync(device_ptr_dest, device_ptr_src,
+                          num_elements * sizeof(TypeName),
+                          cudaMemcpyDeviceToDevice, stream));
+}
 /** \brief Copy CUDA memory from device to host
  *
  * @param device_ptr    device pointer
@@ -118,7 +145,19 @@ inline void copyFromDevice(TypeName *device_ptr, TypeName *host_ptr,
   HANDLE_ERROR(cudaMemcpy(host_ptr, device_ptr, num_elements * sizeof(TypeName),
                           cudaMemcpyDeviceToHost));
 }
-
+/** \brief Copy CUDA memory from device to host
+ *
+ * @param device_ptr    device pointer
+ * @param host_ptr      host pointer
+ * @param num_elements  amount of elements of size TypeName
+ */
+template <typename TypeName>
+inline void copyFromDeviceAsync(TypeName *device_ptr, TypeName *host_ptr,
+                           IndType num_elements, cudaStream_t stream=0)
+{
+  HANDLE_ERROR(cudaMemcpyAsync(host_ptr, device_ptr, num_elements * sizeof(TypeName),
+                          cudaMemcpyDeviceToHost, stream));
+}
 /** \brief Free variable list of device pointers. Use NULL as stopping element
  *
  * e.g.: freeTotalDeviceMemory(ptr1*, ptr2*,NULL);
@@ -181,7 +220,7 @@ inline void showMemoryInfo(bool force, FILE *stream)
   size_t total_mem = 0;
   cudaMemGetInfo(&free_mem, &total_mem);
   if (DEBUG || force)
-    fprintf(stream, "memory usage, free: %lu total: %lu\n", (SizeType)(free_mem),
+    printf("memory usage, free: %lu total: %lu\n", (SizeType)(free_mem),
     (SizeType)(total_mem));
 }
 
@@ -212,39 +251,6 @@ inline void showMemoryInfo()
  *
  * @param symbol Const symbol name
  */
-void initConstSymbol(const char *symbol, const void *src, IndType count);
-
-/** \brief Initialize texture memory on device
- *
- * CUDA Kernel function prototype.
- *
- * @param symbol Texture symbol name
- */
-void initTexture(const char *symbol, cudaArray **devicePtr,
-                 gpuNUFFT::Array<DType> hostTexture);
-
-/** \brief Bind to 1-d texture on device
- *
- * CUDA Kernel function prototype.
- *
- * @param symbol Texture symbol name
- */
-void bindTo1DTexture(const char *symbol, void *devicePtr, IndType count);
-
-/** \brief Unbind from device texture
- *
- * CUDA Kernel function prototype.
- *
- * @param symbol Texture symbol name
- */
-void unbindTexture(const char *symbol);
-
-/** \brief Free texture memory on device
- *
- * CUDA Kernel function prototype.
- *
- * @param symbol Texture symbol name
- */
-void freeTexture(const char *symbol, cudaArray *devicePtr);
+void initConstSymbol(const char *symbol, const void *src, IndType count, cudaStream_t stream=0);
 
 #endif