FM / THC

/usr/local/lib64/python3.6/site-packages/torch/include/THC/

/usr/local/lib64/python3.6/site-packages/torch/include/THC

Name	Size	Mode	Actions
generic/	-	0755	rm
THC.h	328	0644	edit dl rm
THCAllocator.h	370	0644	edit dl rm
THCAsmUtils.cuh	3444	0644	edit dl rm
THCAtomics.cuh	13094	0644	edit dl rm
THCCachingHostAllocator.h	1254	0644	edit dl rm
THCDeviceTensor-inl.cuh	11515	0644	edit dl rm
THCDeviceTensor.cuh	16160	0644	edit dl rm
THCDeviceTensorUtils-inl.cuh	4497	0644	edit dl rm
THCDeviceTensorUtils.cuh	2720	0644	edit dl rm
THCDeviceUtils.cuh	942	0644	edit dl rm
THCGeneral.h	2741	0644	edit dl rm
THCGeneral.hpp	743	0644	edit dl rm
THCGenerateAllTypes.h	958	0644	edit dl rm
THCGenerateBFloat16Type.h	515	0644	edit dl rm
THCGenerateBoolType.h	449	0644	edit dl rm
THCGenerateByteType.h	419	0644	edit dl rm
THCGenerateCharType.h	418	0644	edit dl rm
THCGenerateComplexDoubleType.h	533	0644	edit dl rm
THCGenerateComplexFloatType.h	526	0644	edit dl rm
THCGenerateComplexTypes.h	298	0644	edit dl rm
THCGenerateDoubleType.h	464	0644	edit dl rm
THCGenerateFloatType.h	550	0644	edit dl rm
THCGenerateFloatTypes.h	779	0644	edit dl rm
THCGenerateHalfType.h	481	0644	edit dl rm
THCGenerateIntType.h	414	0644	edit dl rm
THCGenerateLongType.h	419	0644	edit dl rm
THCGenerateShortType.h	424	0644	edit dl rm
THCIntegerDivider.cuh	4095	0644	edit dl rm
THCNumerics.cuh	19813	0644	edit dl rm
THCScanUtils.cuh	4789	0644	edit dl rm
THCSleep.h	232	0644	edit dl rm
THCStorage.h	487	0644	edit dl rm
THCStorage.hpp	846	0644	edit dl rm
THCStorageCopy.h	466	0644	edit dl rm
THCTensor.h	623	0644	edit dl rm
THCTensor.hpp	1070	0644	edit dl rm
THCTensorCopy.h	467	0644	edit dl rm
THCTensorCopy.hpp	600	0644	edit dl rm
THCTensorMathReduce.cuh	664	0644	edit dl rm
THCThrustAllocator.cuh	618	0644	edit dl rm

Edit: /usr/local/lib64/python3.6/site-packages/torch/include/THC/THCAsmUtils.cuh (3444B)

#ifndef THC_ASM_UTILS_INC
#define THC_ASM_UTILS_INC

// Collection of direct PTX functions

template <typename T>
struct Bitfield {};

template <>
struct Bitfield<unsigned int> {
  static __device__ __forceinline__
  unsigned int getBitfield(unsigned int val, int pos, int len) {
#if defined(__HIP_PLATFORM_HCC__)
    pos &= 0xff;
    len &= 0xff;

unsigned int m = (1u << len) - 1u;
    return (val >> pos) & m;
#else
    unsigned int ret;
    asm("bfe.u32 %0, %1, %2, %3;" : "=r"(ret) : "r"(val), "r"(pos), "r"(len));
    return ret;
#endif
  }

static __device__ __forceinline__
  unsigned int setBitfield(unsigned int val, unsigned int toInsert, int pos, int len) {
#if defined(__HIP_PLATFORM_HCC__)
    pos &= 0xff;
    len &= 0xff;

unsigned int m = (1u << len) - 1u;
    toInsert &= m;
    toInsert <<= pos;
    m <<= pos;

return (val & ~m) | toInsert;
#else
    unsigned int ret;
    asm("bfi.b32 %0, %1, %2, %3, %4;" :
        "=r"(ret) : "r"(toInsert), "r"(val), "r"(pos), "r"(len));
    return ret;
#endif
  }
};

template <>
struct Bitfield<uint64_t> {
  static __device__ __forceinline__
  uint64_t getBitfield(uint64_t val, int pos, int len) {
#if defined(__HIP_PLATFORM_HCC__)
    pos &= 0xff;
    len &= 0xff;

uint64_t m = (1u << len) - 1u;
    return (val >> pos) & m;
#else
    uint64_t ret;
    asm("bfe.u64 %0, %1, %2, %3;" : "=l"(ret) : "l"(val), "r"(pos), "r"(len));
    return ret;
#endif
  }

static __device__ __forceinline__
  uint64_t setBitfield(uint64_t val, uint64_t toInsert, int pos, int len) {
#if defined(__HIP_PLATFORM_HCC__)
    pos &= 0xff;
    len &= 0xff;

uint64_t m = (1u << len) - 1u;
    toInsert &= m;
    toInsert <<= pos;
    m <<= pos;

return (val & ~m) | toInsert;
#else
    uint64_t ret;
    asm("bfi.b64 %0, %1, %2, %3, %4;" :
        "=l"(ret) : "l"(toInsert), "l"(val), "r"(pos), "r"(len));
    return ret;
#endif
  }
};

__device__ __forceinline__ int getLaneId() {
#if defined(__HIP_PLATFORM_HCC__)
  return __lane_id();
#else
  int laneId;
  asm("mov.s32 %0, %%laneid;" : "=r"(laneId) );
  return laneId;
#endif
}

#if defined(__HIP_PLATFORM_HCC__)
__device__ __forceinline__ unsigned long long int getLaneMaskLt() {
  const std::uint64_t m = (1ull << getLaneId()) - 1ull;
  return m;
}
#else
__device__ __forceinline__ unsigned getLaneMaskLt() {
  unsigned mask;
  asm("mov.u32 %0, %%lanemask_lt;" : "=r"(mask));
  return mask;
}
#endif

#if defined (__HIP_PLATFORM_HCC__)
__device__ __forceinline__ unsigned long long int getLaneMaskLe() {
  std::uint64_t m = UINT64_MAX >> (sizeof(std::uint64_t) * CHAR_BIT - (getLaneId() + 1));
  return m;
}
#else
__device__ __forceinline__ unsigned getLaneMaskLe() {
  unsigned mask;
  asm("mov.u32 %0, %%lanemask_le;" : "=r"(mask));
  return mask;
}
#endif

#if defined(__HIP_PLATFORM_HCC__)
__device__ __forceinline__ unsigned long long int getLaneMaskGt() {
  const std::uint64_t m = getLaneMaskLe();
  return m ? ~m : m;
}
#else
__device__ __forceinline__ unsigned getLaneMaskGt() {
  unsigned mask;
  asm("mov.u32 %0, %%lanemask_gt;" : "=r"(mask));
  return mask;
}
#endif

#if defined(__HIP_PLATFORM_HCC__)
__device__ __forceinline__ unsigned long long int getLaneMaskGe() {
  const std::uint64_t m = getLaneMaskLt();
  return ~m;
}
#else
__device__ __forceinline__ unsigned getLaneMaskGe() {
  unsigned mask;
  asm("mov.u32 %0, %%lanemask_ge;" : "=r"(mask));
  return mask;
}
#endif

#endif // THC_ASM_UTILS_INC

cmd: