 THC_POINTWISE_APPLY2(THCudaTensor)
 THC_POINTWISE_APPLY2(THCudaDoubleTensor)
 THC_POINTWISE_APPLY2(THCudaHalfTensor)
+  
+
+template <typename TensorTypeA,
+          typename TensorTypeB,
+          typename TensorTypeC>
+bool THC_pointwiseApply3(THCState* state,
+                         TensorTypeA* a,
+                         TensorTypeB* b,
+                         TensorTypeC* c,
+                         const char* apply_header,
+                         const char* op_string) {
+  TensorArgType aType = ReadWrite;
+  TensorArgType bType = ReadWrite;
+  TensorArgType cType = ReadWrite;
+  cudaStream_t stream = state->currentStream;
+  long totalElements = TensorUtils<TensorTypeA>::getNumElements(state, a);
+
+  if (totalElements != TensorUtils<TensorTypeB>::getNumElements(state, b) ||
+      totalElements != TensorUtils<TensorTypeC>::getNumElements(state, c)) {
+    return false;
+  }
+
+  if (TensorUtils<TensorTypeA>::getDims(state, a) > MAX_CUTORCH_DIMS ||
+      TensorUtils<TensorTypeB>::getDims(state, b) > MAX_CUTORCH_DIMS ||
+      TensorUtils<TensorTypeC>::getDims(state, c) > MAX_CUTORCH_DIMS) {
+    return false;
+  }
+
+  if (TensorUtils<TensorTypeA>::getDims(state, a) == 0) {
+    // Zero-dim tensor; do nothing
+    return true;
+  }
+
+  const dim3 block = getApplyBlock();
+
+  dim3 grid;
+  if (!getApplyGrid(state, totalElements, grid)) {
+    return false;
+  }
+
+  // If tensor args have overlapping indices and are read/write, then
+  // we must expand the tensor to a contiguous form first, since
+  // otherwise there are conflicting writes. Upon copying back to the
+  // non-contiguous form, there will be conflicting writes, but at
+  // least with copy, one of the updaters will win atomically. This is
+  // a sketchy property of the old system as well (writing into all
+  // indices of a tensor with overlapping indices should probably be
+  // an error, since it is unclear which one should win), but we will
+  // preserve this last-writer-wins (in arbitrary copy order) behavior.
+  TensorTypeA* oldA = NULL;
+  TensorTypeB* oldB = NULL;
+  TensorTypeC* oldC = NULL;
+
+  if (aType == ReadWrite &&
+      TensorUtils<TensorTypeA>::overlappingIndices(state, a)) {
+    // Must perform in contiguous space
+    oldA = a;
+    a = TensorUtils<TensorTypeA>::newContiguous(state, a);
+  }
+  if (bType == ReadWrite &&
+      TensorUtils<TensorTypeB>::overlappingIndices(state, b)) {
+    // Must perform in contiguous space
+    oldB = b;
+    b = TensorUtils<TensorTypeB>::newContiguous(state, b);
+  }
+  if (cType == ReadWrite &&
+      TensorUtils<TensorTypeC>::overlappingIndices(state, c)) {
+    // Must perform in contiguous space
+    oldC = c;
+    c = TensorUtils<TensorTypeC>::newContiguous(state, c);
+  }
+
+#define HANDLE_CASE(TYPE, A, B, C)                                      \
+  kernelPointwiseApply3RTC<typename TensorUtils<TensorTypeA>::DataType, \
+                           typename TensorUtils<TensorTypeB>::DataType, \
+                           typename TensorUtils<TensorTypeC>::DataType, \
+                           TYPE>                                        \
+      (aInfo, bInfo, cInfo, apply_header, op_string, (TYPE) totalElements, grid, block, A, B, C, stream);
+
+#define HANDLE_C_CASE(TYPE, A, B, C)            \
+  {                                             \
+    if (cInfo.isContiguous()) {                 \
+      HANDLE_CASE(TYPE, A, B, -2);              \
+    } else {                                    \
+      switch (C) {                              \
+        case 1:                                 \
+        HANDLE_CASE(TYPE, A, B, 1);             \
+        break;                                  \
+        case 2:                                 \
+        HANDLE_CASE(TYPE, A, B, 2);             \
+        break;                                  \
+        default:                                \
+        HANDLE_CASE(TYPE, A, B, -1);            \
+        break;                                  \
+      }                                         \
+    }                                           \
+  }
+
+#define HANDLE_B_CASE(TYPE, A, B, C)            \
+  {                                             \
+    if (bInfo.isContiguous()) {                 \
+      HANDLE_C_CASE(TYPE, A, -2, C);            \
+    } else {                                    \
+      switch (B) {                              \
+        case 1:                                 \
+        HANDLE_C_CASE(TYPE, A, 1, C);           \
+        break;                                  \
+        case 2:                                 \
+        HANDLE_C_CASE(TYPE, A, 2, C);           \
+        break;                                  \
+        default:                                \
+        HANDLE_C_CASE(TYPE, A, -1, C);          \
+        break;                                  \
+      }                                         \
+    }                                           \
+  }
+
+#define HANDLE_A_CASE(TYPE, A, B, C)            \
+  {                                             \
+    if (aInfo.isContiguous()) {                 \
+      HANDLE_B_CASE(TYPE, -2, B, C);            \
+    } else {                                    \
+      switch (A) {                              \
+        case 1:                                 \
+        HANDLE_B_CASE(TYPE, 1, B, C);           \
+        break;                                  \
+        case 2:                                 \
+        HANDLE_B_CASE(TYPE, 2, B, C);           \
+        break;                                  \
+        default:                                \
+        HANDLE_B_CASE(TYPE, -1, B, C);          \
+        break;                                  \
+      }                                         \
+    }                                           \
+  }
+
+  if (TensorUtils<TensorTypeA>::canUse32BitIndexMath(state, a) &&
+      TensorUtils<TensorTypeB>::canUse32BitIndexMath(state, b) &&
+      TensorUtils<TensorTypeC>::canUse32BitIndexMath(state, c)) {
+    TensorInfo<typename TensorUtils<TensorTypeA>::DataType, unsigned int> aInfo =
+      getTensorInfo<TensorTypeA, unsigned int>(state, a);
+    aInfo.collapseDims();
+
+    TensorInfo<typename TensorUtils<TensorTypeB>::DataType, unsigned int> bInfo =
+      getTensorInfo<TensorTypeB, unsigned int>(state, b);
+    bInfo.collapseDims();
+
+    TensorInfo<typename TensorUtils<TensorTypeC>::DataType, unsigned int> cInfo =
+      getTensorInfo<TensorTypeC, unsigned int>(state, c);
+    cInfo.collapseDims();
+
+    HANDLE_A_CASE(unsigned int, aInfo.dims, bInfo.dims, cInfo.dims);
+  } else {
+    TensorInfo<typename TensorUtils<TensorTypeA>::DataType, unsigned long> aInfo =
+      getTensorInfo<TensorTypeA, unsigned long>(state, a);
+    aInfo.collapseDims();
+
+    TensorInfo<typename TensorUtils<TensorTypeB>::DataType, unsigned long> bInfo =
+      getTensorInfo<TensorTypeB, unsigned long>(state, b);
+    bInfo.collapseDims();
+
+    TensorInfo<typename TensorUtils<TensorTypeC>::DataType, unsigned long> cInfo =
+      getTensorInfo<TensorTypeC, unsigned long>(state, c);
+    cInfo.collapseDims();
+
+    // For large tensors, we only compile the completely contiguous
+    // version and the completely generic version, to reduce
+    // compilation time.
+    if (aInfo.isContiguous() && bInfo.isContiguous() && cInfo.isContiguous()) {
+      kernelPointwiseApply3RTC<typename TensorUtils<TensorTypeA>::DataType,
+                               typename TensorUtils<TensorTypeB>::DataType,
+                               typename TensorUtils<TensorTypeC>::DataType,
+                               unsigned long>
+        (aInfo, bInfo, cInfo, apply_header, op_string, totalElements, grid, block, -2, -2, -2, stream);
+    } else {
+      kernelPointwiseApply3RTC<typename TensorUtils<TensorTypeA>::DataType,
+                               typename TensorUtils<TensorTypeB>::DataType,
+                               typename TensorUtils<TensorTypeC>::DataType,
+                               unsigned long>
+        (aInfo, bInfo, cInfo, apply_header, op_string, (unsigned long) totalElements, grid, block, -1, -1, -1, stream);
+    }
+  }
+#undef HANDLE_CASE
+#undef HANDLE_C_CASE
+#undef HANDLE_B_CASE
+#undef HANDLE_A_CASE
+
+  if (oldA) {
+    // Ignore overlaps when copying back; if we use THCTensor_copy
+    // instead, it will recursively try and invoke ourselves to make
+    // oldA contiguous.
+    TensorUtils<TensorTypeA>::copyIgnoringOverlaps(state, oldA, a);
+    TensorUtils<TensorTypeA>::free(state, a);
+    a = oldA;
+  }
+
+  if (oldB) {
+    // Ignore overlaps when copying back; if we use THCTensor_copy
+    // instead, it will recursively try and invoke ourselves to make
+    // oldB contiguous.
+    TensorUtils<TensorTypeB>::copyIgnoringOverlaps(state, oldB, b);
+    TensorUtils<TensorTypeB>::free(state, b);
+    b = oldB;
+  }
+
+  if (oldC) {
+    // Ignore overlaps when copying back; if we use THCTensor_copy
+    // instead, it will recursively try and invoke ourselves to make
+    // oldC contiguous.
+    TensorUtils<TensorTypeC>::copyIgnoringOverlaps(state, oldC, c);
+    TensorUtils<TensorTypeC>::free(state, c);
+    c = oldC;
+  }
+
+  return true;
+}
+
+
+#define THC_POINTWISE_APPLY3(TYPE) \
+  extern "C" \
+  bool TH_CONCAT_2(TYPE, _pointwiseApply3)(THCState* state, \
+                                    TYPE* a, \
+                                    TYPE* b, \
+                                    TYPE* c, \
+                                    const char* apply_header, \
+                                    const char* op_string) { \
+    return THC_pointwiseApply3<TYPE>(state, a, b, c, apply_header, op_string); \
+  }
+
+THC_POINTWISE_APPLY3(THCudaTensor)
+THC_POINTWISE_APPLY3(THCudaDoubleTensor)
+THC_POINTWISE_APPLY3(THCudaHalfTensor)

  launch(ptx->data(), "kernel", args, grid, block, (CUstream)stream);
}

*/

template <typename TensorTypeA>
bool THC_pointwiseApply1(THCState* state,

}


// Example op: 'x = y*z'
const char* instanciate_apply3 = "                                      \n\
#include <header.h>                                                     \n\
typedef %s Ta;                                                          \n\
typedef %s Tb;                                                          \n\
typedef %s Tc;                                                          \n\
typedef %s IndexType; 							\n\
struct Op {                                                             \n\
  __device__ __forceinline__						\n\
  void operator()(Ta* a, Tb* b, Tc *c) {            			\n\
    float& x = *a;                                                      \n\
    float& y = *b;							\n\
    float& z = *c;							\n\
