Delete Deepspeed-Windows directory
This commit is contained in:
1 parent
f52ccd3931
commit
bf5d913cae
7 files changed
-4963
No files matched your search
@@ -1,285 +0,0 @@
|
||||
// Copyright (c) Microsoft Corporation.
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
|
||||
// DeepSpeed Team
|
||||
|
||||
#include <ATen/cuda/CUDAContext.h>
|
||||
#include <torch/extension.h>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
#include "quantization.h"
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_quantize(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if ((((size / groups) - 1) / 4096 + 1) <= 256) {
|
||||
launch_fake_quantize_kernel(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_sr_quantize(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if (((size / groups) / 4 / 1024) <= 256) {
|
||||
launch_sr_fake_quantize_kernel(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_quantize_asym(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if ((((size / groups) - 1) / 4096 + 1) <= 256) {
|
||||
launch_fake_quantize_kernel_asym(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_sr_quantize_asym(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if (((size / groups) / 4 / 1024) <= 256) {
|
||||
launch_sr_fake_quantize_kernel_asym(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> quantize_kernel(at::Tensor& input_vals,
|
||||
int groups,
|
||||
int numBits,
|
||||
quantize::Type quantType)
|
||||
{
|
||||
auto dtype = at::kFloat;
|
||||
auto params_options = at::TensorOptions()
|
||||
.dtype(dtype)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int param_elems = (quantize::requires_offset(quantType)) ? 2 : 1;
|
||||
auto params = torch::empty({groups, param_elems}, params_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
auto output_sizes = input_vals.sizes().vec();
|
||||
output_sizes[output_sizes.size() - 1] /= numBits == 8 ? 1 : 2;
|
||||
auto output = torch::empty(output_sizes, output_options);
|
||||
|
||||
const int elems_per_group = at::numel(input_vals) / groups;
|
||||
|
||||
launch_quant((int8_t*)output.data_ptr(),
|
||||
(float*)params.data_ptr(),
|
||||
(__half*)input_vals.data_ptr(),
|
||||
groups,
|
||||
elems_per_group,
|
||||
numBits,
|
||||
quantType,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return {output, params};
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor dequantize(at::Tensor& quantized_data,
|
||||
at::Tensor& params,
|
||||
int groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type)
|
||||
{
|
||||
auto dtype = (std::is_same<T, float>::value) ? torch::kFloat32 : torch::kFloat16;
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(dtype)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
auto output_sizes = quantized_data.sizes().vec();
|
||||
output_sizes[output_sizes.size() - 1] *= num_bits == 8 ? 1 : 2;
|
||||
auto output = torch::empty(output_sizes, output_options);
|
||||
|
||||
const int total_elems = at::numel(output);
|
||||
const int elems_per_group = total_elems / groups;
|
||||
|
||||
launch_dequantize_kernel((T*)output.data_ptr(),
|
||||
(const int8_t*)quantized_data.data_ptr(),
|
||||
(const float*)params.data_ptr(),
|
||||
quant_type,
|
||||
num_bits,
|
||||
elems_per_group,
|
||||
total_elems,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return output;
|
||||
}
|
||||
|
||||
at::Tensor dequantize_int4_to_half_experimental(at::Tensor& data_in,
|
||||
at::Tensor& scale_buffer,
|
||||
at::Tensor& min_val_buffer,
|
||||
int num_group,
|
||||
int group_size)
|
||||
{
|
||||
auto output_options = at::TensorOptions().dtype(at::kHalf).device(at::kCUDA);
|
||||
auto output = torch::empty({num_group, group_size}, output_options);
|
||||
|
||||
launch_dequantize_int4_to_half_experimental((uint8_t*)data_in.data_ptr(),
|
||||
(half*)output.data_ptr(),
|
||||
(half*)scale_buffer.data_ptr(),
|
||||
(half*)min_val_buffer.data_ptr(),
|
||||
num_group,
|
||||
group_size,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return output;
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> ds_swizzle_quant(at::Tensor& input_vals,
|
||||
int groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type,
|
||||
int pipeline_size,
|
||||
int nodes,
|
||||
int devices_per_node)
|
||||
{
|
||||
auto scales_options = at::TensorOptions()
|
||||
.dtype(at::kFloat)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int scales_elems = (quantize::requires_offset(quant_type)) ? 2 : 1;
|
||||
auto scales = torch::empty({groups, scales_elems}, scales_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
const int quantization_scalar = 8 / num_bits;
|
||||
const int compressed_vals = at::numel(input_vals) / quantization_scalar;
|
||||
|
||||
auto output = torch::empty({compressed_vals}, output_options);
|
||||
const int elems_per_group = at::numel(input_vals) / groups;
|
||||
|
||||
launch_swizzled_quant((int8_t*)output.data_ptr(),
|
||||
(float*)scales.data_ptr(),
|
||||
(__half*)input_vals.data_ptr(),
|
||||
num_bits,
|
||||
quant_type,
|
||||
groups,
|
||||
elems_per_group,
|
||||
pipeline_size,
|
||||
nodes,
|
||||
devices_per_node,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return {output, scales};
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> quantized_reduction(at::Tensor& input_vals,
|
||||
at::Tensor& input_scales,
|
||||
int in_groups,
|
||||
int out_groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type,
|
||||
int devices_per_node)
|
||||
{
|
||||
auto scales_options = at::TensorOptions()
|
||||
.dtype(at::kFloat)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int scales_elems = (quantize::requires_offset(quant_type)) ? 2 : 1;
|
||||
auto scales = torch::empty({out_groups, scales_elems}, scales_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
std::vector<int64_t> sz_vector(input_vals.sizes().begin(), input_vals.sizes().end());
|
||||
sz_vector[sz_vector.size() - 1] = sz_vector.back() / devices_per_node; // num of GPU per nodes
|
||||
at::IntArrayRef sz(sz_vector);
|
||||
auto output = torch::empty(sz, output_options);
|
||||
|
||||
const int elems_per_in_tensor = at::numel(input_vals) / devices_per_node;
|
||||
const int elems_per_in_group = elems_per_in_tensor / (in_groups / devices_per_node);
|
||||
const int elems_per_out_group = elems_per_in_tensor / out_groups;
|
||||
|
||||
// old code:
|
||||
// std::vector<long int> sz(input_vals.sizes().begin(), input_vals.sizes().end());
|
||||
// sz[sz.size() - 1] = sz.back() / devices_per_node; // num of GPU per nodes
|
||||
// const int elems_per_in_tensor = at::numel(input_vals) / devices_per_node;
|
||||
// auto output = torch::empty(sz, output_options);
|
||||
|
||||
// const int elems_per_in_group = elems_per_in_tensor / (in_groups / devices_per_node);
|
||||
// const int elems_per_out_group = elems_per_in_tensor / out_groups;
|
||||
|
||||
launch_dequant_reduce((int8_t*)output.data_ptr(),
|
||||
(float*)scales.data_ptr(),
|
||||
(const int8_t*)input_vals.data_ptr(),
|
||||
(const float*)input_scales.data_ptr(),
|
||||
devices_per_node,
|
||||
num_bits,
|
||||
quant_type,
|
||||
out_groups,
|
||||
elems_per_out_group,
|
||||
elems_per_in_tensor,
|
||||
in_groups / devices_per_node,
|
||||
elems_per_in_group,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
return {output, scales};
|
||||
}
|
||||
|
||||
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m)
|
||||
{
|
||||
m.def("ds_quantize_fp32", &ds_quantize<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_quantize_fp16", &ds_quantize<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_sr_quantize_fp32", &ds_sr_quantize<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_sr_quantize_fp16", &ds_sr_quantize<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_quantize_asym_fp32", &ds_quantize_asym<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def(
|
||||
"ds_quantize_asym_fp16", &ds_quantize_asym<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_sr_quantize_asym_fp32",
|
||||
&ds_sr_quantize_asym<float>,
|
||||
"DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_sr_quantize_asym_fp16",
|
||||
&ds_sr_quantize_asym<__half>,
|
||||
"DeepSpeed Quantize with fp16 (CUDA)");
|
||||
pybind11::enum_<quantize::Type>(m, "QuantizationType")
|
||||
.value("Symmetric", quantize::Type::Symmetric)
|
||||
.value("Asymmetric", quantize::Type::Asymmetric)
|
||||
.export_values();
|
||||
m.def("quantize", &quantize_kernel);
|
||||
m.def("dequantize", &dequantize<__half>);
|
||||
m.def("dequantize_fp32", &dequantize<float>);
|
||||
m.def("dequantize_int4_to_half_experimental",
|
||||
&dequantize_int4_to_half_experimental,
|
||||
"Dequantize int4 to half (experimental)");
|
||||
m.def("swizzle_quant", &ds_swizzle_quant);
|
||||
m.def("quantized_reduction", &quantized_reduction);
|
||||
}
|
||||
File diff suppressed because it is too large.
Load diff
@@ -1,160 +0,0 @@
|
||||
# Copyright (c) Microsoft Corporation.
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
|
||||
# DeepSpeed Team
|
||||
|
||||
import os
|
||||
import torch
|
||||
import deepspeed
|
||||
import subprocess
|
||||
import argparse
|
||||
import psutil
|
||||
from .ops.op_builder.all_ops import ALL_OPS
|
||||
from .git_version_info import installed_ops, torch_info
|
||||
from deepspeed.accelerator import get_accelerator
|
||||
|
||||
GREEN = '\033[92m'
|
||||
RED = '\033[91m'
|
||||
YELLOW = '\033[93m'
|
||||
END = '\033[0m'
|
||||
SUCCESS = f"{GREEN} [SUCCESS] {END}"
|
||||
OKAY = f"{GREEN}[OKAY]{END}"
|
||||
WARNING = f"{YELLOW}[WARNING]{END}"
|
||||
FAIL = f'{RED}[FAIL]{END}'
|
||||
INFO = '[INFO]'
|
||||
|
||||
color_len = len(GREEN) + len(END)
|
||||
okay = f"{GREEN}[OKAY]{END}"
|
||||
warning = f"{YELLOW}[WARNING]{END}"
|
||||
|
||||
|
||||
def op_report(verbose=True):
|
||||
max_dots = 23
|
||||
max_dots2 = 11
|
||||
h = ["op name", "installed", "compatible"]
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print("DeepSpeed C++/CUDA extension op report")
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
|
||||
print("NOTE: Ops not installed will be just-in-time (JIT) compiled at\n"
|
||||
" runtime if needed. Op compatibility means that your system\n"
|
||||
" meet the required dependencies to JIT install the op.")
|
||||
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print("JIT compiled ops requires ninja")
|
||||
ninja_status = OKAY if ninja_installed() else FAIL
|
||||
print('ninja', "." * (max_dots - 5), ninja_status)
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print(h[0], "." * (max_dots - len(h[0])), h[1], "." * (max_dots2 - len(h[1])), h[2])
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
installed = f"{GREEN}[YES]{END}"
|
||||
no = f"{YELLOW}[NO]{END}"
|
||||
for op_name, builder in ALL_OPS.items():
|
||||
dots = "." * (max_dots - len(op_name))
|
||||
is_compatible = OKAY if builder.is_compatible(verbose) else no
|
||||
is_installed = installed if installed_ops.get(op_name, False) else no
|
||||
dots2 = '.' * ((len(h[1]) + (max_dots2 - len(h[1]))) - (len(is_installed) - color_len))
|
||||
print(op_name, dots, is_installed, dots2, is_compatible)
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
|
||||
|
||||
def ninja_installed():
|
||||
try:
|
||||
import ninja # noqa: F401 # type: ignore
|
||||
except ImportError:
|
||||
return False
|
||||
return True
|
||||
|
||||
|
||||
def nvcc_version():
|
||||
import torch.utils.cpp_extension
|
||||
cuda_home = torch.utils.cpp_extension.CUDA_HOME
|
||||
if cuda_home is None:
|
||||
return f"{RED} [FAIL] cannot find CUDA_HOME via torch.utils.cpp_extension.CUDA_HOME={torch.utils.cpp_extension.CUDA_HOME} {END}"
|
||||
try:
|
||||
output = subprocess.check_output([cuda_home + "/bin/nvcc", "-V"], universal_newlines=True)
|
||||
except FileNotFoundError:
|
||||
return f"{RED} [FAIL] nvcc missing {END}"
|
||||
output_split = output.split()
|
||||
release_idx = output_split.index("release")
|
||||
release = output_split[release_idx + 1].replace(',', '').split(".")
|
||||
return ".".join(release)
|
||||
|
||||
|
||||
def get_shm_size():
|
||||
try:
|
||||
temp_dir = os.getenv('TEMP') or os.getenv('TMP') or os.path.join(os.path.expanduser('~'), 'tmp')
|
||||
shm_stats = psutil.disk_usage(temp_dir)
|
||||
shm_size = shm_stats.total
|
||||
shm_hbytes = human_readable_size(shm_size)
|
||||
warn = []
|
||||
if shm_size < 512 * 1024**2:
|
||||
warn.append(
|
||||
f" {YELLOW} [WARNING] Shared memory size might be too small, consider increasing it. {END}"
|
||||
)
|
||||
# Add additional warnings specific to your use case if needed.
|
||||
return shm_hbytes, warn
|
||||
except Exception as e:
|
||||
return "UNKNOWN", [f"Error getting shared memory size: {e}"]
|
||||
|
||||
def human_readable_size(size):
|
||||
units = ['B', 'KB', 'MB', 'GB', 'TB']
|
||||
i = 0
|
||||
while size >= 1024 and i < len(units) - 1:
|
||||
size /= 1024
|
||||
i += 1
|
||||
return f'{size:.2f} {units[i]}'
|
||||
|
||||
|
||||
def debug_report():
|
||||
max_dots = 33
|
||||
|
||||
report = [("torch install path", torch.__path__), ("torch version", torch.__version__),
|
||||
("deepspeed install path", deepspeed.__path__),
|
||||
("deepspeed info", f"{deepspeed.__version__}, {deepspeed.__git_hash__}, {deepspeed.__git_branch__}")]
|
||||
if get_accelerator().device_name() == 'cuda':
|
||||
hip_version = getattr(torch.version, "hip", None)
|
||||
report.extend([("torch cuda version", torch.version.cuda), ("torch hip version", hip_version),
|
||||
("nvcc version", (None if hip_version else nvcc_version())),
|
||||
("deepspeed wheel compiled w.", f"torch {torch_info['version']}, " +
|
||||
(f"hip {torch_info['hip_version']}" if hip_version else f"cuda {torch_info['cuda_version']}"))
|
||||
])
|
||||
else:
|
||||
report.extend([("deepspeed wheel compiled w.", f"torch {torch_info['version']} ")])
|
||||
|
||||
report.append(("shared memory (/dev/shm) size", get_shm_size()))
|
||||
|
||||
print("DeepSpeed general environment info:")
|
||||
for name, value in report:
|
||||
warns = []
|
||||
if isinstance(value, tuple):
|
||||
value, warns = value
|
||||
print(name, "." * (max_dots - len(name)), value)
|
||||
if warns:
|
||||
for warn in warns:
|
||||
print(warn)
|
||||
|
||||
|
||||
def parse_arguments():
|
||||
parser = argparse.ArgumentParser()
|
||||
parser.add_argument('--hide_operator_status',
|
||||
action='store_true',
|
||||
help='Suppress display of installation and compatibility statuses of DeepSpeed operators. ')
|
||||
parser.add_argument('--hide_errors_and_warnings', action='store_true', help='Suppress warning and error messages.')
|
||||
args = parser.parse_args()
|
||||
return args
|
||||
|
||||
|
||||
def main(hide_operator_status=False, hide_errors_and_warnings=False):
|
||||
if not hide_operator_status:
|
||||
op_report(verbose=not hide_errors_and_warnings)
|
||||
debug_report()
|
||||
|
||||
|
||||
def cli_main():
|
||||
args = parse_arguments()
|
||||
main(hide_operator_status=args.hide_operator_status, hide_errors_and_warnings=args.hide_errors_and_warnings)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
@@ -1,299 +0,0 @@
|
||||
// Copyright (c) Microsoft Corporation.
|
||||
// SPDX-License-Identifier: Apache-2.0
|
||||
|
||||
// DeepSpeed Team
|
||||
|
||||
#include <ATen/cuda/CUDAContext.h>
|
||||
#include <torch/extension.h>
|
||||
#include <cassert>
|
||||
#include <vector>
|
||||
#include "quantization.h"
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_quantize(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if ((((size / groups) - 1) / 4096 + 1) <= 256) {
|
||||
launch_fake_quantize_kernel(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_sr_quantize(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if (((size / groups) / 4 / 1024) <= 256) {
|
||||
launch_sr_fake_quantize_kernel(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_quantize_asym(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if ((((size / groups) - 1) / 4096 + 1) <= 256) {
|
||||
launch_fake_quantize_kernel_asym(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor ds_sr_quantize_asym(at::Tensor& vals, int groups, int bits)
|
||||
{
|
||||
auto t_size = vals.sizes();
|
||||
int size = 1;
|
||||
for (auto dim : t_size) size *= dim;
|
||||
|
||||
if (((size / groups) / 4 / 1024) <= 256) {
|
||||
launch_sr_fake_quantize_kernel_asym(
|
||||
(T*)vals.data_ptr(), size, groups, bits, at::cuda::getCurrentCUDAStream());
|
||||
}
|
||||
return vals;
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> quantize_kernel(at::Tensor& input_vals,
|
||||
int groups,
|
||||
int numBits,
|
||||
quantize::Type quantType)
|
||||
{
|
||||
auto dtype = at::kFloat;
|
||||
auto params_options = at::TensorOptions()
|
||||
.dtype(dtype)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int param_elems = (quantize::requires_offset(quantType)) ? 2 : 1;
|
||||
auto params = torch::empty({groups, param_elems}, params_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
auto output_sizes = input_vals.sizes().vec();
|
||||
output_sizes[output_sizes.size() - 1] /= numBits == 8 ? 1 : 2;
|
||||
auto output = torch::empty(output_sizes, output_options);
|
||||
|
||||
const int elems_per_group = at::numel(input_vals) / groups;
|
||||
|
||||
launch_quant((int8_t*)output.data_ptr(),
|
||||
(float*)params.data_ptr(),
|
||||
(__half*)input_vals.data_ptr(),
|
||||
groups,
|
||||
elems_per_group,
|
||||
numBits,
|
||||
quantType,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return {output, params};
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
at::Tensor dequantize(at::Tensor& quantized_data,
|
||||
at::Tensor& params,
|
||||
int groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type)
|
||||
{
|
||||
auto dtype = (std::is_same<T, float>::value) ? torch::kFloat32 : torch::kFloat16;
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(dtype)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
auto output_sizes = quantized_data.sizes().vec();
|
||||
output_sizes[output_sizes.size() - 1] *= num_bits == 8 ? 1 : 2;
|
||||
auto output = torch::empty(output_sizes, output_options);
|
||||
|
||||
const int total_elems = at::numel(output);
|
||||
const int elems_per_group = total_elems / groups;
|
||||
|
||||
launch_dequantize_kernel((T*)output.data_ptr(),
|
||||
(const int8_t*)quantized_data.data_ptr(),
|
||||
(const float*)params.data_ptr(),
|
||||
quant_type,
|
||||
num_bits,
|
||||
elems_per_group,
|
||||
total_elems,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return output;
|
||||
}
|
||||
|
||||
at::Tensor dequantize_int4_to_half_experimental(at::Tensor& data_in,
|
||||
at::Tensor& scale_buffer,
|
||||
at::Tensor& min_val_buffer,
|
||||
int num_group,
|
||||
int group_size)
|
||||
{
|
||||
auto output_options = at::TensorOptions().dtype(at::kHalf).device(at::kCUDA);
|
||||
auto output = torch::empty({num_group, group_size}, output_options);
|
||||
|
||||
launch_dequantize_int4_to_half_experimental((uint8_t*)data_in.data_ptr(),
|
||||
(half*)output.data_ptr(),
|
||||
(half*)scale_buffer.data_ptr(),
|
||||
(half*)min_val_buffer.data_ptr(),
|
||||
num_group,
|
||||
group_size,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return output;
|
||||
}
|
||||
|
||||
at::Tensor dequantize_int8_to_half_experimental(at::Tensor& data_in,
|
||||
at::Tensor& scale_buffer,
|
||||
at::Tensor& min_val_buffer,
|
||||
int num_group,
|
||||
int group_size)
|
||||
{
|
||||
auto output_options = at::TensorOptions().dtype(at::kHalf).device(at::kCUDA);
|
||||
auto output = torch::empty({num_group, group_size}, output_options);
|
||||
|
||||
launch_dequantize_int8_to_half_experimental((uint8_t*)data_in.data_ptr(),
|
||||
(half*)output.data_ptr(),
|
||||
(half*)scale_buffer.data_ptr(),
|
||||
(half*)min_val_buffer.data_ptr(),
|
||||
num_group,
|
||||
group_size,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return output;
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> ds_swizzle_quant(at::Tensor& input_vals,
|
||||
int groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type,
|
||||
int pipeline_size,
|
||||
int nodes,
|
||||
int devices_per_node)
|
||||
{
|
||||
auto scales_options = at::TensorOptions()
|
||||
.dtype(at::kFloat)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int scales_elems = (quantize::requires_offset(quant_type)) ? 2 : 1;
|
||||
auto scales = torch::empty({groups, scales_elems}, scales_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
const int quantization_scalar = 8 / num_bits;
|
||||
const int compressed_vals = at::numel(input_vals) / quantization_scalar;
|
||||
|
||||
auto output = torch::empty({compressed_vals}, output_options);
|
||||
const int elems_per_group = at::numel(input_vals) / groups;
|
||||
|
||||
launch_swizzled_quant((int8_t*)output.data_ptr(),
|
||||
(float*)scales.data_ptr(),
|
||||
(__half*)input_vals.data_ptr(),
|
||||
num_bits,
|
||||
quant_type,
|
||||
groups,
|
||||
elems_per_group,
|
||||
pipeline_size,
|
||||
nodes,
|
||||
devices_per_node,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
|
||||
return {output, scales};
|
||||
}
|
||||
|
||||
std::vector<at::Tensor> quantized_reduction(at::Tensor& input_vals,
|
||||
at::Tensor& input_scales,
|
||||
int in_groups,
|
||||
int out_groups,
|
||||
int num_bits,
|
||||
quantize::Type quant_type,
|
||||
int devices_per_node)
|
||||
{
|
||||
auto scales_options = at::TensorOptions()
|
||||
.dtype(at::kFloat)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
const int scales_elems = (quantize::requires_offset(quant_type)) ? 2 : 1;
|
||||
auto scales = torch::empty({out_groups, scales_elems}, scales_options);
|
||||
|
||||
auto output_options = at::TensorOptions()
|
||||
.dtype(at::kChar)
|
||||
.layout(at::kStrided)
|
||||
.device(at::kCUDA)
|
||||
.requires_grad(false);
|
||||
|
||||
std::vector<int64_t> sz_vector(input_vals.sizes().begin(), input_vals.sizes().end());
|
||||
sz_vector[sz_vector.size() - 1] = sz_vector.back() / devices_per_node; // num of GPU per nodes
|
||||
at::IntArrayRef sz(sz_vector);
|
||||
auto output = torch::empty(sz, output_options);
|
||||
|
||||
const int elems_per_in_tensor = at::numel(input_vals) / devices_per_node;
|
||||
const int elems_per_in_group = elems_per_in_tensor / (in_groups / devices_per_node);
|
||||
const int elems_per_out_group = elems_per_in_tensor / out_groups;
|
||||
|
||||
launch_dequant_reduce((int8_t*)output.data_ptr(),
|
||||
(float*)scales.data_ptr(),
|
||||
(const int8_t*)input_vals.data_ptr(),
|
||||
(const float*)input_scales.data_ptr(),
|
||||
devices_per_node,
|
||||
num_bits,
|
||||
quant_type,
|
||||
out_groups,
|
||||
elems_per_out_group,
|
||||
elems_per_in_tensor,
|
||||
in_groups / devices_per_node,
|
||||
elems_per_in_group,
|
||||
at::cuda::getCurrentCUDAStream());
|
||||
return {output, scales};
|
||||
}
|
||||
|
||||
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m)
|
||||
{
|
||||
m.def("ds_quantize_fp32", &ds_quantize<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_quantize_fp16", &ds_quantize<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_sr_quantize_fp32", &ds_sr_quantize<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_sr_quantize_fp16", &ds_sr_quantize<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_quantize_asym_fp32", &ds_quantize_asym<float>, "DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def(
|
||||
"ds_quantize_asym_fp16", &ds_quantize_asym<__half>, "DeepSpeed Quantize with fp16 (CUDA)");
|
||||
m.def("ds_sr_quantize_asym_fp32",
|
||||
&ds_sr_quantize_asym<float>,
|
||||
"DeepSpeed Quantize with fp32 (CUDA)");
|
||||
m.def("ds_sr_quantize_asym_fp16",
|
||||
&ds_sr_quantize_asym<__half>,
|
||||
"DeepSpeed Quantize with fp16 (CUDA)");
|
||||
pybind11::enum_<quantize::Type>(m, "QuantizationType")
|
||||
.value("Symmetric", quantize::Type::Symmetric)
|
||||
.value("Asymmetric", quantize::Type::Asymmetric)
|
||||
.export_values();
|
||||
m.def("quantize", &quantize_kernel);
|
||||
m.def("dequantize", &dequantize<__half>);
|
||||
m.def("dequantize_fp32", &dequantize<float>);
|
||||
m.def("dequantize_int4_to_half_experimental",
|
||||
&dequantize_int4_to_half_experimental,
|
||||
"Dequantize int4 to half (experimental)");
|
||||
m.def("dequantize_int8_to_half_experimental",
|
||||
&dequantize_int8_to_half_experimental,
|
||||
"Dequantize int8 to half (experimental)");
|
||||
m.def("swizzle_quant", &ds_swizzle_quant);
|
||||
m.def("quantized_reduction", &quantized_reduction);
|
||||
}
|
||||
File diff suppressed because it is too large.
Load diff
@@ -1,161 +0,0 @@
|
||||
# Copyright (c) Microsoft Corporation.
|
||||
# SPDX-License-Identifier: Apache-2.0
|
||||
|
||||
# DeepSpeed Team
|
||||
|
||||
import os
|
||||
import torch
|
||||
import deepspeed
|
||||
import subprocess
|
||||
import argparse
|
||||
import psutil
|
||||
from .ops.op_builder.all_ops import ALL_OPS
|
||||
from .git_version_info import installed_ops, torch_info
|
||||
from deepspeed.accelerator import get_accelerator
|
||||
|
||||
GREEN = '\033[92m'
|
||||
RED = '\033[91m'
|
||||
YELLOW = '\033[93m'
|
||||
END = '\033[0m'
|
||||
SUCCESS = f"{GREEN} [SUCCESS] {END}"
|
||||
OKAY = f"{GREEN}[OKAY]{END}"
|
||||
WARNING = f"{YELLOW}[WARNING]{END}"
|
||||
FAIL = f'{RED}[FAIL]{END}'
|
||||
INFO = '[INFO]'
|
||||
|
||||
color_len = len(GREEN) + len(END)
|
||||
okay = f"{GREEN}[OKAY]{END}"
|
||||
warning = f"{YELLOW}[WARNING]{END}"
|
||||
|
||||
|
||||
def op_report(verbose=True):
|
||||
max_dots = 23
|
||||
max_dots2 = 11
|
||||
h = ["op name", "installed", "compatible"]
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print("DeepSpeed C++/CUDA extension op report")
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
|
||||
print("NOTE: Ops not installed will be just-in-time (JIT) compiled at\n"
|
||||
" runtime if needed. Op compatibility means that your system\n"
|
||||
" meet the required dependencies to JIT install the op.")
|
||||
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print("JIT compiled ops requires ninja")
|
||||
ninja_status = OKAY if ninja_installed() else FAIL
|
||||
print('ninja', "." * (max_dots - 5), ninja_status)
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
print(h[0], "." * (max_dots - len(h[0])), h[1], "." * (max_dots2 - len(h[1])), h[2])
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
installed = f"{GREEN}[YES]{END}"
|
||||
no = f"{YELLOW}[NO]{END}"
|
||||
for op_name, builder in ALL_OPS.items():
|
||||
dots = "." * (max_dots - len(op_name))
|
||||
is_compatible = OKAY if builder.is_compatible(verbose) else no
|
||||
is_installed = installed if installed_ops.get(op_name, False) else no
|
||||
dots2 = '.' * ((len(h[1]) + (max_dots2 - len(h[1]))) - (len(is_installed) - color_len))
|
||||
print(op_name, dots, is_installed, dots2, is_compatible)
|
||||
print("-" * (max_dots + max_dots2 + len(h[0]) + len(h[1])))
|
||||
|
||||
|
||||
def ninja_installed():
|
||||
try:
|
||||
import ninja # noqa: F401 # type: ignore
|
||||
except ImportError:
|
||||
return False
|
||||
return True
|
||||
|
||||
|
||||
def nvcc_version():
|
||||
import torch.utils.cpp_extension
|
||||
cuda_home = torch.utils.cpp_extension.CUDA_HOME
|
||||
if cuda_home is None:
|
||||
return f"{RED} [FAIL] cannot find CUDA_HOME via torch.utils.cpp_extension.CUDA_HOME={torch.utils.cpp_extension.CUDA_HOME} {END}"
|
||||
try:
|
||||
output = subprocess.check_output([cuda_home + "/bin/nvcc", "-V"], universal_newlines=True)
|
||||
except FileNotFoundError:
|
||||
return f"{RED} [FAIL] nvcc missing {END}"
|
||||
output_split = output.split()
|
||||
release_idx = output_split.index("release")
|
||||
release = output_split[release_idx + 1].replace(',', '').split(".")
|
||||
return ".".join(release)
|
||||
|
||||
|
||||
def get_shm_size():
|
||||
try:
|
||||
temp_dir = os.getenv('TEMP') or os.getenv('TMP') or os.path.join(os.path.expanduser('~'), 'tmp')
|
||||
shm_stats = psutil.disk_usage(temp_dir)
|
||||
shm_size = shm_stats.total
|
||||
shm_hbytes = human_readable_size(shm_size)
|
||||
warn = []
|
||||
if shm_size < 512 * 1024**2:
|
||||
warn.append(
|
||||
f" {YELLOW} [WARNING] Shared memory size might be too small, consider increasing it. {END}"
|
||||
)
|
||||
# Add additional warnings specific to your use case if needed.
|
||||
return shm_hbytes, warn
|
||||
except Exception as e:
|
||||
return "UNKNOWN", [f"Error getting shared memory size: {e}"]
|
||||
|
||||
|
||||
def human_readable_size(size):
|
||||
units = ['B', 'KB', 'MB', 'GB', 'TB']
|
||||
i = 0
|
||||
while size >= 1024 and i < len(units) - 1:
|
||||
size /= 1024
|
||||
i += 1
|
||||
return f'{size:.2f} {units[i]}'
|
||||
|
||||
|
||||
def debug_report():
|
||||
max_dots = 33
|
||||
|
||||
report = [("torch install path", torch.__path__), ("torch version", torch.__version__),
|
||||
("deepspeed install path", deepspeed.__path__),
|
||||
("deepspeed info", f"{deepspeed.__version__}, {deepspeed.__git_hash__}, {deepspeed.__git_branch__}")]
|
||||
if get_accelerator().device_name() == 'cuda':
|
||||
hip_version = getattr(torch.version, "hip", None)
|
||||
report.extend([("torch cuda version", torch.version.cuda), ("torch hip version", hip_version),
|
||||
("nvcc version", (None if hip_version else nvcc_version())),
|
||||
("deepspeed wheel compiled w.", f"torch {torch_info['version']}, " +
|
||||
(f"hip {torch_info['hip_version']}" if hip_version else f"cuda {torch_info['cuda_version']}"))
|
||||
])
|
||||
else:
|
||||
report.extend([("deepspeed wheel compiled w.", f"torch {torch_info['version']} ")])
|
||||
|
||||
report.append(("shared memory (/dev/shm) size", get_shm_size()))
|
||||
|
||||
print("DeepSpeed general environment info:")
|
||||
for name, value in report:
|
||||
warns = []
|
||||
if isinstance(value, tuple):
|
||||
value, warns = value
|
||||
print(name, "." * (max_dots - len(name)), value)
|
||||
if warns:
|
||||
for warn in warns:
|
||||
print(warn)
|
||||
|
||||
|
||||
def parse_arguments():
|
||||
parser = argparse.ArgumentParser()
|
||||
parser.add_argument('--hide_operator_status',
|
||||
action='store_true',
|
||||
help='Suppress display of installation and compatibility statuses of DeepSpeed operators. ')
|
||||
parser.add_argument('--hide_errors_and_warnings', action='store_true', help='Suppress warning and error messages.')
|
||||
args = parser.parse_args()
|
||||
return args
|
||||
|
||||
|
||||
def main(hide_operator_status=False, hide_errors_and_warnings=False):
|
||||
if not hide_operator_status:
|
||||
op_report(verbose=not hide_errors_and_warnings)
|
||||
debug_report()
|
||||
|
||||
|
||||
def cli_main():
|
||||
args = parse_arguments()
|
||||
main(hide_operator_status=args.hide_operator_status, hide_errors_and_warnings=args.hide_errors_and_warnings)
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
@@ -1,18 +0,0 @@
|
||||
These are the modified files you would need for a MANUAL build of DeepSpeed 11.1 or 11.2.
|
||||
|
||||
They were created by a mix of instructions from erew123 and S95Sedan.
|
||||
|
||||
|
||||
- You still have to download the DeepSpeed v11.1 or v11.2 installation files from Microsoft to use these files.
|
||||
|
||||
https://github.com/microsoft/DeepSpeed/releases/tag/v0.11.2
|
||||
|
||||
|
||||
- These files were provided by S95Sedan here:
|
||||
|
||||
https://github.com/S95Sedan/Deepspeed-Windows
|
||||
|
||||
|
||||
- They are related to the manual build of DeepSpeed for AllTalk, which is linked here:
|
||||
|
||||
https://github.com/erew123/alltalk_tts/tree/main?tab=readme-ov-file#-option-2---a-bit-more-complicated
|
||||
Reference in new issue
Block a user