     cudaStream_t stream)
 {
   // using c++11 std::is_same here
-  const char* type;
-  if (std::is_same<IndexType, unsigned long>::value)
-    type = "unsigned long";
-  else if(std::is_same<IndexType, unsigned int>::value)
-    type = "unsigned int";
-
-  char src[2048];
-  sprintf(src, instanciate_apply3, type, op, A, B, C);
+  const char* index_type = getTypeString<IndexType>();
+  const char* ta_type = getTypeString<Ta>();
+  const char* tb_type = getTypeString<Tb>();
+  const char* tc_type = getTypeString<Tc>();
+
+  char src[4096];
+  sprintf(src, instanciate_apply3, ta_type, tb_type, tc_type, index_type, op, A, B, C);
   const char *headers[] = {apply_header};
   const char *includeNames[] = {"header.h"};
 
   PTXPtr ptx;
-  Apply3Hash hash(op, type, A, B, C);
-  auto found_hash = apply3cache.find(hash);
-  if(found_hash == apply3cache.end())
+  ApplyHash hash(op, std::string(ta_type)+"_"+tb_type+"_"+tc_type, index_type, A, B, C);
+  auto found_hash = applycache.find(hash);
+  if(found_hash == applycache.end())
   {
     ptx = PTXPtr(new PTX());
     compilePTX(src, headers, includeNames, *ptx);
-    apply3cache.emplace(hash, ptx);
+    applycache.emplace(hash, ptx);
   }
   else
     ptx = found_hash->second;

}


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

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)
