mps_utils.mm 2.19 KB
Newer Older
Leyuan Wang committed
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80
/*!
 *  Copyright (c) 2017 by Contributors
 * \file Use external mps utils function
 */
#include "mps_utils.h"

namespace tvm {
namespace contrib {

// MPS Data Type
MPSDataType MPSType::DLTypeToMPSType(const DLDataType &dtype) {
  switch (dtype.code) {
  case kDLInt:
    if (dtype.bits == 8 && dtype.lanes == 1)
      return MPSDataTypeInt8;
    else if (dtype.bits == 16 && dtype.lanes == 1)
      return MPSDataTypeInt16;
    else
      LOG(FATAL) << "Unsupported type";
    break;
  case kDLUInt:
    if (dtype.bits == 8 && dtype.lanes == 1)
      return MPSDataTypeUInt8;
    else if (dtype.bits == 16 && dtype.lanes == 1)
      return MPSDataTypeUInt16;
    else if (dtype.bits == 32 && dtype.lanes == 1)
      return MPSDataTypeUInt32;
    LOG(FATAL) << "Unsupported type";
    break;
  case kDLFloat:
    if (dtype.bits == 16 && dtype.lanes == 1)
      return MPSDataTypeFloat16;
    else if (dtype.bits == 32 && dtype.lanes == 1)
      return MPSDataTypeFloat32;
    else
      LOG(FATAL) << "Unsupported type";
    break;
  default:
    LOG(FATAL) << "Unsupported type";
  }
  return MPSDataTypeFloat32;
}

// MetalThreadEntry

MPSImage *MetalThreadEntry::AllocMPSImage(id<MTLDevice> dev,
                                          MPSImageDescriptor *desc) {
  MPSImage *mpsimg = [[MPSImage alloc] initWithDevice:dev imageDescriptor:desc];
  img_table.push_back(mpsimg);
  return mpsimg;
}

MPSTemporaryImage *MetalThreadEntry::AllocTempImage(id<MTLCommandBuffer> cb,
                                                    MPSImageDescriptor *desc) {
  MPSTemporaryImage *mpsimg =
      [MPSTemporaryImage temporaryImageWithCommandBuffer:cb
                                         imageDescriptor:desc];
  return mpsimg;
}

MetalThreadEntry::MetalThreadEntry() {
  auto func = runtime::Registry::Get("device_api.metal");
  void *ret = (*func)();
  metal_api = static_cast<runtime::metal::MetalWorkspace *>(ret);
}

MetalThreadEntry::~MetalThreadEntry() {
  for (int i = 0; i < img_table.size(); ++i) {
    [img_table[i] dealloc];
  }
}

typedef dmlc::ThreadLocalStore<MetalThreadEntry> MetalThreadStore;

MetalThreadEntry *MetalThreadEntry::ThreadLocal() {
  return MetalThreadStore::Get();
}

} // namespace contrib
} // namespace tvm