Reduced type set for Android

CMakeLists.txt

@@ -11,10 +11,13 @@ IF (NOT WIN32)
 SET(CMAKE_C_FLAGS "-std=c99 -Werror=implicit-function-declaration ${CMAKE_C_FLAGS}")
 ENDIF (NOT WIN32)
 IF(CUDA_HAS_FP16 OR NOT ${CUDA_VERSION} LESS 7.5)
-  ADD_DEFINITIONS(-DTH_GENERIC_USE_HALF=1)
   ADD_DEFINITIONS(-DCUDA_HAS_FP16=1)
 ENDIF()
 
+IF (ANDROID)
+ ADD_DEFINITIONS(-DTHC_MIN_MATH)
+ENDIF()
+
 INCLUDE_DIRECTORIES(${CUDA_INCLUDE_DIRS})
 
 ADD_SUBDIRECTORY(lib)

TensorMath.lua

@@ -551,6 +551,7 @@ local handledTypeaccreals = {
 
 for k, Tensor_ in pairs(handledTypenames) do
     Tensor = Tensor_
+
     if Tensor == 'CudaHalfTensor' then
         interface:print("#ifdef CUDA_HALF_TENSOR")
     end
@@ -592,6 +593,8 @@ for k, Tensor_ in pairs(handledTypenames) do
          cname("zero"),
          {{name=Tensor, returned=true}})
 
+    interface:print("#ifndef THC_MIN_MATH")
+
     wrap("zeros",
          cname("zeros"),
          {{name=Tensor, default=true, returned=true, method={default='nil'}},
@@ -1335,6 +1338,8 @@ void cutorch_%sMath_init(lua_State *L)
 }
 ]], Tensor, Tensor, Tensor, Tensor))
 
+    interface:print("#endif")
+
     if Tensor == 'CudaHalfTensor' then
         interface:print("#endif")
     end

generic/TensorOperator.c

@@ -2,6 +2,7 @@
 #define THC_GENERIC_FILE "generic/TensorOperator.c"
 #else
 
+#ifndef THC_GENERIC_NO_MATH
 static int cutorch_TensorOperator_(__add__)(lua_State *L)
 {
   THCTensor *tensor1 = luaT_toudata(L, 1, torch_Tensor);
@@ -242,13 +243,16 @@ static int cutorch_TensorOperator_(__div__)(lua_State *L)
 
   return 1;
 }
+#endif
 
 static const struct luaL_Reg cutorch_TensorOperator_(_) [] = {
+#ifndef THC_GENERIC_NO_MATH
   {"__add__", cutorch_TensorOperator_(__add__)},
   {"__sub__", cutorch_TensorOperator_(__sub__)},
   {"__unm__", cutorch_TensorOperator_(__unm__)},
   {"__mul__", cutorch_TensorOperator_(__mul__)},
   {"__div__", cutorch_TensorOperator_(__div__)},
+#endif
   {NULL, NULL}
 };
 

init.c

@@ -1056,26 +1056,30 @@ int luaopen_libcutorch(lua_State *L)
   cutorch_HalfTensorCopy_init(L);
 #endif
 
+  cutorch_CudaTensorOperator_init(L);
+#ifndef THC_MIN_MATH
   cutorch_CudaByteTensorOperator_init(L);
   cutorch_CudaCharTensorOperator_init(L);
   cutorch_CudaShortTensorOperator_init(L);
   cutorch_CudaIntTensorOperator_init(L);
   cutorch_CudaLongTensorOperator_init(L);
-  cutorch_CudaTensorOperator_init(L);
   cutorch_CudaDoubleTensorOperator_init(L);
 #ifdef CUDA_HALF_TENSOR
   cutorch_CudaHalfTensorOperator_init(L);
+#endif
 #endif
 
+  cutorch_CudaTensorMath_init(L);
+#ifndef THC_MIN_MATH
   cutorch_CudaByteTensorMath_init(L);
   cutorch_CudaCharTensorMath_init(L);
   cutorch_CudaShortTensorMath_init(L);
   cutorch_CudaIntTensorMath_init(L);
   cutorch_CudaLongTensorMath_init(L);
-  cutorch_CudaTensorMath_init(L);
   cutorch_CudaDoubleTensorMath_init(L);
 #ifdef CUDA_HALF_TENSOR
   cutorch_CudaHalfTensorMath_init(L);
+#endif
 #endif
 
   cutorch_Event_init(L);
@@ -1091,6 +1095,13 @@ int luaopen_libcutorch(lua_State *L)
 #endif
   lua_setfield(L, -2, "hasHalf");
 
+#ifdef THC_MIN_MATH 
+  lua_pushboolean(L, 1);
+#else
+  lua_pushboolean(L, 0);
+#endif
+  lua_setfield(L, -2, "minMath");
+
   /* store gpu driver version in field */
   int driverVersion;
   THCudaCheck(cudaDriverGetVersion(&driverVersion));

lib/THC/CMakeLists.txt

@@ -75,6 +75,10 @@ IF ($ENV{TH_BINARY_BUILD})
   SET(CMAKE_CXX_FLAGS "-static-libstdc++ ${CMAKE_CXX_FLAGS}")
 ENDIF()
 
+# IF (ARM)
+ ADD_DEFINITIONS(-DTHC_MIN_MATH)
+# ENDIF()
+
 IF(APPLE)
   IF(${CUDA_VERSION} LESS 6.0)
     # work around for mac os x bug:
@@ -206,14 +210,15 @@ IF(USE_MAGMA)
   TARGET_LINK_LIBRARIES(THC ${MAGMA_LIBRARIES} ${CUDA_cusparse_LIBRARY})
 ENDIF(USE_MAGMA)
 
+IF(NOT ANDROID)
 IF(NOT THC_SO_VERSION)
   SET(THC_SO_VERSION 0)
 ENDIF(NOT THC_SO_VERSION)
 MESSAGE(STATUS "THC_SO_VERSION: ${THC_SO_VERSION}")
 SET_TARGET_PROPERTIES(THC PROPERTIES
   VERSION   ${THC_SO_VERSION}
   SOVERSION ${THC_SO_VERSION})
-
+ENDIF()
 
 INSTALL(TARGETS THC
           RUNTIME DESTINATION "${THC_INSTALL_BIN_SUBDIR}"

lib/THC/THCGenerateByteType.h

@@ -7,13 +7,17 @@
 #define Real Byte
 #define CReal CudaByte
 #define THC_REAL_IS_BYTE
+# ifdef THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
 #undef CReal
 #undef THC_REAL_IS_BYTE
+#undef THC_GENERIC_NO_MATH 
 
 #ifndef THCGenerateAllTypes
 #undef THC_GENERIC_FILE

lib/THC/THCGenerateCharType.h

@@ -7,14 +7,17 @@
 #define Real Char
 #define CReal CudaChar
 #define THC_REAL_IS_CHAR
+# ifdef THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
 #undef CReal
 #undef THC_REAL_IS_CHAR
-
+#undef THC_GENERIC_NO_MATH
 #ifndef THCGenerateAllTypes
 #undef THC_GENERIC_FILE
 #endif

lib/THC/THCGenerateDoubleType.h

@@ -7,14 +7,17 @@
 #define Real Double
 #define CReal CudaDouble
 #define THC_REAL_IS_DOUBLE
+# ifdef THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
 #undef CReal
 #undef THC_REAL_IS_DOUBLE
-
+#undef THC_GENERIC_NO_MATH
 #ifndef THCGenerateAllTypes
 #ifndef THCGenerateFloatTypes
 #undef THC_GENERIC_FILE

lib/THC/THCGenerateHalfType.h

@@ -15,14 +15,16 @@
 #ifdef CUDA_HALF_TENSOR
 #define CReal CudaHalf
 #endif
-
+# if defined THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #define THC_REAL_IS_HALF
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
-
+#undef THC_GENERIC_NO_MATH
 #ifdef CUDA_HALF_TENSOR
 #undef CReal
 #endif

lib/THC/THCGenerateIntType.h

@@ -7,14 +7,17 @@
 #define Real Int
 #define CReal CudaInt
 #define THC_REAL_IS_INT
+# ifdef THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
 #undef CReal
 #undef THC_REAL_IS_INT
-
+#undef THC_GENERIC_NO_MATH
 #ifndef THCGenerateAllTypes
 #undef THC_GENERIC_FILE
 #endif

lib/THC/THCGenerateShortType.h

@@ -7,14 +7,17 @@
 #define Real Short
 #define CReal CudaShort
 #define THC_REAL_IS_SHORT
+# ifdef THC_MIN_MATH
+#  define THC_GENERIC_NO_MATH 1
+# endif
 #line 1 THC_GENERIC_FILE
 #include THC_GENERIC_FILE
 #undef real
 #undef accreal
 #undef Real
 #undef CReal
 #undef THC_REAL_IS_SHORT
-
+#undef THC_GENERIC_NO_MATH
 #ifndef THCGenerateAllTypes
 #undef THC_GENERIC_FILE
 #endif

lib/THC/THCHalf.h

@@ -3,10 +3,12 @@
 
 #include "THCGeneral.h"
 
+#ifndef THC_MIN_MATH
 /* We compile with CudaHalfTensor support if we have this: */
 #if CUDA_VERSION >= 7050 || CUDA_HAS_FP16
 #define CUDA_HALF_TENSOR 1
 #endif
+#endif
 
 #ifdef CUDA_HALF_TENSOR
 

lib/THC/generic/THCTensorMath.cu

@@ -390,5 +390,4 @@ accreal THCTensor_(trace)(THCState *state, THCTensor *src_) {
   THCTensor_(free)(state, diag);
   return trace;
 }
-
 #endif

lib/THC/generic/THCTensorMathBlas.cu

@@ -1,7 +1,7 @@
 #ifndef THC_GENERIC_FILE
 #define THC_GENERIC_FILE "generic/THCTensorMathBlas.cu"
 #else
-
+# ifndef THC_GENERIC_NO_MATH
 THC_API accreal
 THCTensor_(dot)(THCState *state, THCTensor *self, THCTensor *src)
 {
@@ -596,5 +596,6 @@ THCTensor_(baddbmm)(THCState *state, THCTensor *result, real beta, THCTensor *t,
   THError("unimplemented data type");
 #endif
 }
+# endif /* THC_GENERIC_NO_MATH */
 
 #endif

lib/THC/generic/THCTensorMathCompareT.cu

@@ -1,7 +1,7 @@
 #ifndef THC_GENERIC_FILE
 #define THC_GENERIC_FILE "generic/THCTensorMathCompareT.cu"
 #else
-
+# ifndef THC_GENERIC_NO_MATH
 THC_API void
 THCTensor_(ltTensor)(THCState *state, THCudaByteTensor *self_, THCTensor *src1, THCTensor *src2)
 {
@@ -109,5 +109,5 @@ THCTensor_(neTensorT)(THCState *state, THCTensor *self_, THCTensor *src1, THCTen
                     TensorNEOp<typename TensorUtils<THCTensor>::DataType,
                     typename TensorUtils<THCTensor>::DataType>());
 }
-
+# endif /* THC_GENERIC_NO_MATH */
 #endif

lib/THC/generic/THCTensorMathMagma.cu

@@ -1,7 +1,7 @@
 #ifndef THC_GENERIC_FILE
 #define THC_GENERIC_FILE "generic/THCTensorMathMagma.cu"
 #else
-
+# ifndef THC_GENERIC_NO_MATH
 #if defined(THC_REAL_IS_FLOAT) || defined(THC_REAL_IS_DOUBLE)
 
 #ifdef USE_MAGMA
@@ -646,5 +646,5 @@ THC_API void THCTensor_(qr)(THCState *state, THCTensor *rq_, THCTensor *rr_, THC
 }
 
 #endif
-
+# endif /* THC_GENERIC_NO_MATH */
 #endif

lib/THC/generic/THCTensorMathPairwise.cu

@@ -1,7 +1,7 @@
 #ifndef THC_GENERIC_FILE
 #define THC_GENERIC_FILE "generic/THCTensorMathPairwise.cu"
 #else
-
+# ifndef THC_GENERIC_NO_MATH
 THC_API void
 THCTensor_(add)(THCState *state, THCTensor *self_, THCTensor *src_, real value)
 {
@@ -209,5 +209,5 @@ THC_API int THCTensor_(equal)(THCState *state, THCTensor *self_, THCTensor *src_
 
   return min != 0;
 }
-
+# endif /* THC_GENERIC_NO_MATH */
 #endif

lib/THC/generic/THCTensorMathPointwise.cu

@@ -1,7 +1,7 @@
 #ifndef THC_GENERIC_FILE
 #define THC_GENERIC_FILE "generic/THCTensorMathPointwise.cu"
 #else
-
+# ifndef THC_GENERIC_NO_MATH
 #define IMPLEMENT_CUDA_TENSOR_BASIC_FUNC_(NAME, CFUNC, REAL)             \
   struct Tensor_##NAME##_##REAL##_Op {                                  \
     __device__ __forceinline__ void operator()(real* out, real* in) const { \
@@ -518,5 +518,5 @@ THCTensor_(addcdiv)(THCState *state, THCTensor *self_, THCTensor *t, real value,
 
   THCudaCheck(cudaGetLastError());
 }
-
+# endif /* THC_GENERIC_NO_MATH */
 #endif