Upload edit\Qwen3-TTS-test\.venv\Lib\site-packages\torch\include\ATen\native\mps\MetalShaderLibrary.h with huggingface_hub
Browse files
edit//Qwen3-TTS-test//.venv//Lib//site-packages//torch//include//ATen//native//mps//MetalShaderLibrary.h
ADDED
|
@@ -0,0 +1,159 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#pragma once
|
| 2 |
+
#ifdef __OBJC__
|
| 3 |
+
#include <Metal/Metal.h>
|
| 4 |
+
typedef id<MTLLibrary> MTLLibrary_t;
|
| 5 |
+
typedef id<MTLFunction> MTLFunction_t;
|
| 6 |
+
typedef id<MTLComputePipelineState> MTLComputePipelineState_t;
|
| 7 |
+
typedef id<MTLComputeCommandEncoder> MTLComputeCommandEncoder_t;
|
| 8 |
+
#else
|
| 9 |
+
typedef void MTLCompileOptions;
|
| 10 |
+
typedef void* MTLLibrary_t;
|
| 11 |
+
typedef void* MTLFunction_t;
|
| 12 |
+
typedef void* MTLComputePipelineState_t;
|
| 13 |
+
typedef void* MTLComputeCommandEncoder_t;
|
| 14 |
+
#endif
|
| 15 |
+
|
| 16 |
+
#include <c10/util/OptionalArrayRef.h>
|
| 17 |
+
#include <functional>
|
| 18 |
+
#include <optional>
|
| 19 |
+
#include <type_traits>
|
| 20 |
+
#include <unordered_map>
|
| 21 |
+
#include <vector>
|
| 22 |
+
|
| 23 |
+
// Forward declaration of TensorBase
|
| 24 |
+
namespace at {
|
| 25 |
+
class TensorBase;
|
| 26 |
+
}
|
| 27 |
+
|
| 28 |
+
namespace at::native::mps {
|
| 29 |
+
|
| 30 |
+
namespace detail {
|
| 31 |
+
template <typename T>
|
| 32 |
+
class has_size_type {
|
| 33 |
+
template <typename U>
|
| 34 |
+
static constexpr std::true_type check(typename U::size_type*);
|
| 35 |
+
template <typename>
|
| 36 |
+
static constexpr std::false_type check(...);
|
| 37 |
+
|
| 38 |
+
public:
|
| 39 |
+
static constexpr bool value = decltype(check<T>(nullptr))::value;
|
| 40 |
+
};
|
| 41 |
+
|
| 42 |
+
template <typename T>
|
| 43 |
+
constexpr bool has_size_type_v = has_size_type<T>::value;
|
| 44 |
+
|
| 45 |
+
} // namespace detail
|
| 46 |
+
|
| 47 |
+
class MetalKernelFunction {
|
| 48 |
+
public:
|
| 49 |
+
MetalKernelFunction(MTLComputePipelineState_t cps_);
|
| 50 |
+
~MetalKernelFunction();
|
| 51 |
+
MetalKernelFunction(MetalKernelFunction&) = delete;
|
| 52 |
+
// Shader properties
|
| 53 |
+
uint64_t getMaxThreadsPerThreadgroup() const;
|
| 54 |
+
uint64_t getThreadExecutionWidth() const;
|
| 55 |
+
uint64_t getStaticThreadGroupMemoryLength() const;
|
| 56 |
+
void runCommandBlock(std::function<void(void)> f);
|
| 57 |
+
// Methods below should be called from runCommandBlock functionT
|
| 58 |
+
void startEncoding();
|
| 59 |
+
void setArg(unsigned idx, const at::TensorBase& t);
|
| 60 |
+
void setArg(unsigned idx, const void* ptr, uint64_t size);
|
| 61 |
+
template <
|
| 62 |
+
typename T,
|
| 63 |
+
typename = std::enable_if_t<
|
| 64 |
+
std::is_integral_v<T> || std::is_same_v<T, float> ||
|
| 65 |
+
(std::is_class_v<T> && std::is_trivially_copyable_v<T> &&
|
| 66 |
+
!detail::has_size_type_v<T>)>>
|
| 67 |
+
inline void setArg(unsigned idx, const T val) {
|
| 68 |
+
setArg(idx, &val, sizeof(T));
|
| 69 |
+
}
|
| 70 |
+
|
| 71 |
+
template <
|
| 72 |
+
typename Container,
|
| 73 |
+
typename = std::enable_if_t<detail::has_size_type_v<Container>>>
|
| 74 |
+
inline void setArg(unsigned idx, const Container& values) {
|
| 75 |
+
setArg(
|
| 76 |
+
idx,
|
| 77 |
+
values.data(),
|
| 78 |
+
values.size() * sizeof(typename Container::value_type));
|
| 79 |
+
}
|
| 80 |
+
void dispatch(
|
| 81 |
+
uint64_t length,
|
| 82 |
+
std::optional<uint64_t> groupSize = std::nullopt);
|
| 83 |
+
void dispatch(
|
| 84 |
+
c10::ArrayRef<uint64_t> length,
|
| 85 |
+
c10::OptionalArrayRef<uint64_t> groupSize = std::nullopt);
|
| 86 |
+
|
| 87 |
+
private:
|
| 88 |
+
MTLComputePipelineState_t cps;
|
| 89 |
+
MTLComputeCommandEncoder_t encoder = nullptr;
|
| 90 |
+
};
|
| 91 |
+
|
| 92 |
+
class MetalShaderLibrary {
|
| 93 |
+
public:
|
| 94 |
+
MetalShaderLibrary(const std::string& src)
|
| 95 |
+
: shaderSource(src), nparams(0), compile_options(nullptr) {}
|
| 96 |
+
MetalShaderLibrary(const std::string& src, unsigned nparams_)
|
| 97 |
+
: shaderSource(src), nparams(nparams_), compile_options(nullptr) {}
|
| 98 |
+
MetalShaderLibrary(
|
| 99 |
+
const std::string& src,
|
| 100 |
+
unsigned nparams_,
|
| 101 |
+
MTLCompileOptions* compile_options_)
|
| 102 |
+
: shaderSource(src),
|
| 103 |
+
nparams(nparams_),
|
| 104 |
+
compile_options(compile_options_) {}
|
| 105 |
+
MetalShaderLibrary(const MetalShaderLibrary&) = delete;
|
| 106 |
+
virtual ~MetalShaderLibrary();
|
| 107 |
+
std::vector<std::string> getFunctionNames();
|
| 108 |
+
std::shared_ptr<MetalKernelFunction> getKernelFunction(
|
| 109 |
+
const std::string& name);
|
| 110 |
+
inline MTLComputePipelineState_t getPipelineStateForFunc(
|
| 111 |
+
const std::string& fname) {
|
| 112 |
+
return getLibraryPipelineState(getLibrary(), fname).first;
|
| 113 |
+
}
|
| 114 |
+
MTLComputePipelineState_t getPipelineStateForFunc(
|
| 115 |
+
const std::string& fname,
|
| 116 |
+
const std::initializer_list<std::string>& params) {
|
| 117 |
+
return getLibraryPipelineState(getLibrary(params), fname).first;
|
| 118 |
+
}
|
| 119 |
+
inline MTLFunction_t getMTLFunction(const std::string& fname) {
|
| 120 |
+
return getLibraryPipelineState(getLibrary(), fname).second;
|
| 121 |
+
}
|
| 122 |
+
MTLFunction_t getMTLFunction(
|
| 123 |
+
const std::string& fname,
|
| 124 |
+
const std::initializer_list<std::string>& params) {
|
| 125 |
+
return getLibraryPipelineState(getLibrary(params), fname).second;
|
| 126 |
+
}
|
| 127 |
+
static MetalShaderLibrary& getBundledLibrary();
|
| 128 |
+
|
| 129 |
+
protected:
|
| 130 |
+
virtual MTLLibrary_t getLibrary();
|
| 131 |
+
virtual MTLLibrary_t getLibrary(
|
| 132 |
+
const std::initializer_list<std::string>& params);
|
| 133 |
+
MTLLibrary_t library = nullptr;
|
| 134 |
+
|
| 135 |
+
private:
|
| 136 |
+
std::pair<MTLComputePipelineState_t, MTLFunction_t> getLibraryPipelineState(
|
| 137 |
+
MTLLibrary_t lib,
|
| 138 |
+
const std::string& fname);
|
| 139 |
+
MTLLibrary_t compileLibrary(const std::string& src);
|
| 140 |
+
std::string shaderSource;
|
| 141 |
+
unsigned nparams;
|
| 142 |
+
MTLCompileOptions* compile_options;
|
| 143 |
+
std::unordered_map<std::string, MTLLibrary_t> libMap;
|
| 144 |
+
std::unordered_map<
|
| 145 |
+
std::string,
|
| 146 |
+
std::pair<MTLComputePipelineState_t, MTLFunction_t>>
|
| 147 |
+
cplMap;
|
| 148 |
+
};
|
| 149 |
+
|
| 150 |
+
class DynamicMetalShaderLibrary : public MetalShaderLibrary {
|
| 151 |
+
public:
|
| 152 |
+
DynamicMetalShaderLibrary(const std::string& src) : MetalShaderLibrary(src) {
|
| 153 |
+
// Compile right away
|
| 154 |
+
getLibrary();
|
| 155 |
+
}
|
| 156 |
+
~DynamicMetalShaderLibrary();
|
| 157 |
+
};
|
| 158 |
+
|
| 159 |
+
} // namespace at::native::mps
|