Compare commits
24 Commits
master-1ce
...
master-9c5
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
9c51d8787f | ||
|
|
f9f0d4685b | ||
|
|
8d2050a5cf | ||
|
|
08f5b41956 | ||
|
|
b6daf5c55b | ||
|
|
be6cd1a4bf | ||
|
|
e1384defca | ||
|
|
814280343c | ||
|
|
1d2af5ca3f | ||
|
|
ce1bcc74a6 | ||
|
|
760cfaa618 | ||
|
|
6d16f6853e | ||
|
|
036ba9e6d8 | ||
|
|
ec82d5279a | ||
|
|
afea457eda | ||
|
|
646e77638e | ||
|
|
3ac48ea1a7 | ||
|
|
607e39489f | ||
|
|
ccae95aec9 | ||
|
|
90e9178d18 | ||
|
|
48bcce493f | ||
|
|
a469688e30 | ||
|
|
61980171a1 | ||
|
|
583cc5bba2 |
75
.github/workflows/build.yml
vendored
@@ -4,17 +4,36 @@ on:
|
||||
workflow_dispatch: # allows manual triggering
|
||||
inputs:
|
||||
create_release:
|
||||
description: 'Create new release'
|
||||
description: "Create new release"
|
||||
required: true
|
||||
type: boolean
|
||||
push:
|
||||
branches:
|
||||
- master
|
||||
- ci
|
||||
paths: ['.github/workflows/**', '**/CMakeLists.txt', '**/Makefile', '**/*.h', '**/*.hpp', '**/*.c', '**/*.cpp', '**/*.cu']
|
||||
paths:
|
||||
[
|
||||
".github/workflows/**",
|
||||
"**/CMakeLists.txt",
|
||||
"**/Makefile",
|
||||
"**/*.h",
|
||||
"**/*.hpp",
|
||||
"**/*.c",
|
||||
"**/*.cpp",
|
||||
"**/*.cu",
|
||||
]
|
||||
pull_request:
|
||||
types: [opened, synchronize, reopened]
|
||||
paths: ['**/CMakeLists.txt', '**/Makefile', '**/*.h', '**/*.hpp', '**/*.c', '**/*.cpp', '**/*.cu']
|
||||
paths:
|
||||
[
|
||||
"**/CMakeLists.txt",
|
||||
"**/Makefile",
|
||||
"**/*.h",
|
||||
"**/*.hpp",
|
||||
"**/*.c",
|
||||
"**/*.cpp",
|
||||
"**/*.cu",
|
||||
]
|
||||
|
||||
env:
|
||||
BRANCH_NAME: ${{ github.head_ref || github.ref_name }}
|
||||
@@ -67,12 +86,12 @@ jobs:
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v3
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
|
||||
|
||||
|
||||
macOS-latest-cmake:
|
||||
runs-on: macos-latest
|
||||
|
||||
@@ -120,29 +139,29 @@ jobs:
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v3
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-${{ steps.system-info.outputs.OS_TYPE }}-${{ steps.system-info.outputs.OS_NAME }}-${{ steps.system-info.outputs.OS_VERSION }}-${{ steps.system-info.outputs.CPU_ARCH }}.zip
|
||||
|
||||
|
||||
windows-latest-cmake:
|
||||
runs-on: windows-latest
|
||||
runs-on: windows-2019
|
||||
|
||||
strategy:
|
||||
matrix:
|
||||
include:
|
||||
- build: 'noavx'
|
||||
defines: '-DGGML_AVX=OFF -DGGML_AVX2=OFF -DGGML_FMA=OFF -DSD_BUILD_SHARED_LIBS=ON'
|
||||
- build: 'avx2'
|
||||
defines: '-DGGML_AVX2=ON -DSD_BUILD_SHARED_LIBS=ON'
|
||||
- build: 'avx'
|
||||
defines: '-DGGML_AVX2=OFF -DSD_BUILD_SHARED_LIBS=ON'
|
||||
- build: 'avx512'
|
||||
defines: '-DGGML_AVX512=ON -DSD_BUILD_SHARED_LIBS=ON'
|
||||
- build: 'cuda12'
|
||||
defines: '-DSD_CUBLAS=ON -DSD_BUILD_SHARED_LIBS=ON'
|
||||
- build: 'rocm5.5'
|
||||
- build: "noavx"
|
||||
defines: "-DGGML_AVX=OFF -DGGML_AVX2=OFF -DGGML_FMA=OFF -DSD_BUILD_SHARED_LIBS=ON"
|
||||
- build: "avx2"
|
||||
defines: "-DGGML_AVX2=ON -DSD_BUILD_SHARED_LIBS=ON"
|
||||
- build: "avx"
|
||||
defines: "-DGGML_AVX2=OFF -DSD_BUILD_SHARED_LIBS=ON"
|
||||
- build: "avx512"
|
||||
defines: "-DGGML_AVX512=ON -DSD_BUILD_SHARED_LIBS=ON"
|
||||
- build: "cuda12"
|
||||
defines: "-DSD_CUBLAS=ON -DSD_BUILD_SHARED_LIBS=ON"
|
||||
- build: "rocm5.5"
|
||||
defines: '-G Ninja -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++ -DSD_HIPBLAS=ON -DCMAKE_BUILD_TYPE=Release -DAMDGPU_TARGETS="gfx1100;gfx1102;gfx1030" -DSD_BUILD_SHARED_LIBS=ON'
|
||||
steps:
|
||||
- name: Clone
|
||||
@@ -156,8 +175,8 @@ jobs:
|
||||
if: ${{ matrix.build == 'cuda12' }}
|
||||
uses: Jimver/cuda-toolkit@v0.2.11
|
||||
with:
|
||||
cuda: '12.2.0'
|
||||
method: 'network'
|
||||
cuda: "12.2.0"
|
||||
method: "network"
|
||||
sub-packages: '["nvcc", "cudart", "cublas", "cublas_dev", "thrust", "visual_studio_integration"]'
|
||||
|
||||
- name: Install rocm-toolkit
|
||||
@@ -165,7 +184,7 @@ jobs:
|
||||
if: ${{ matrix.build == 'rocm5.5' }}
|
||||
uses: Cyberhan123/rocm-toolkit@v0.1.0
|
||||
with:
|
||||
rocm: '5.5.0'
|
||||
rocm: "5.5.0"
|
||||
|
||||
- name: Install Ninja
|
||||
id: install-ninja
|
||||
@@ -231,15 +250,17 @@ jobs:
|
||||
|
||||
- name: Upload Cuda runtime
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' && matrix.build == 'cuda12' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v3
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-cudart-sd-bin-win-cu12-x64.zip
|
||||
path: |
|
||||
cudart-sd-bin-win-cu12-x64.zip
|
||||
|
||||
- name: Upload artifacts
|
||||
if: ${{ ( github.event_name == 'push' && github.ref == 'refs/heads/master' ) || github.event.inputs.create_release == 'true' }}
|
||||
uses: actions/upload-artifact@v3
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip
|
||||
path: |
|
||||
sd-${{ env.BRANCH_NAME }}-${{ steps.commit.outputs.short }}-bin-win-${{ matrix.build }}-x64.zip
|
||||
|
||||
@@ -256,7 +277,11 @@ jobs:
|
||||
steps:
|
||||
- name: Download artifacts
|
||||
id: download-artifact
|
||||
uses: actions/download-artifact@v3
|
||||
uses: actions/download-artifact@v4
|
||||
with:
|
||||
path: ./artifact
|
||||
pattern: sd-*
|
||||
merge-multiple: true
|
||||
|
||||
- name: Get commit hash
|
||||
id: commit
|
||||
|
||||
2
.gitmodules
vendored
@@ -1,3 +1,3 @@
|
||||
[submodule "ggml"]
|
||||
path = ggml
|
||||
url = https://github.com/leejet/ggml.git
|
||||
url = https://github.com/ggerganov/ggml.git
|
||||
|
||||
@@ -34,7 +34,7 @@ option(SD_BUILD_SHARED_LIBS "sd: build shared libs" OFF)
|
||||
|
||||
if(SD_CUBLAS)
|
||||
message("Use CUBLAS as backend stable-diffusion")
|
||||
set(GGML_CUBLAS ON)
|
||||
set(GGML_CUDA ON)
|
||||
add_definitions(-DSD_USE_CUBLAS)
|
||||
endif()
|
||||
|
||||
|
||||
60
README.md
@@ -14,6 +14,7 @@ Inference of [Stable Diffusion](https://github.com/CompVis/stable-diffusion) in
|
||||
- !!!The VAE in SDXL encounters NaN issues under FP16, but unfortunately, the ggml_conv_2d only operates under FP16. Hence, a parameter is needed to specify the VAE that has fixed the FP16 NaN issue. You can find it here: [SDXL VAE FP16 Fix](https://huggingface.co/madebyollin/sdxl-vae-fp16-fix/blob/main/sdxl_vae.safetensors).
|
||||
|
||||
- [SD-Turbo](https://huggingface.co/stabilityai/sd-turbo) and [SDXL-Turbo](https://huggingface.co/stabilityai/sdxl-turbo) support
|
||||
- [PhotoMaker](https://github.com/TencentARC/PhotoMaker) support.
|
||||
- 16-bit, 32-bit float support
|
||||
- 4-bit, 5-bit and 8-bit integer quantization support
|
||||
- Accelerated memory-efficient CPU inference
|
||||
@@ -60,6 +61,9 @@ Inference of [Stable Diffusion](https://github.com/CompVis/stable-diffusion) in
|
||||
|
||||
## Usage
|
||||
|
||||
For most users, you can download the built executable program from the latest [release](https://github.com/leejet/stable-diffusion.cpp/releases/latest).
|
||||
If the built product does not meet your requirements, you can choose to build it manually.
|
||||
|
||||
### Get the Code
|
||||
|
||||
```
|
||||
@@ -148,7 +152,7 @@ cmake --build . --config Release
|
||||
### Run
|
||||
|
||||
```
|
||||
usage: ./build/bin/sd [arguments]
|
||||
usage: ./bin/sd [arguments]
|
||||
|
||||
arguments:
|
||||
-h, --help show this help message and exit
|
||||
@@ -160,6 +164,9 @@ arguments:
|
||||
--taesd [TAESD_PATH] path to taesd. Using Tiny AutoEncoder for fast decoding (low quality)
|
||||
--control-net [CONTROL_PATH] path to control net model
|
||||
--embd-dir [EMBEDDING_PATH] path to embeddings.
|
||||
--stacked-id-embd-dir [DIR] path to PHOTOMAKER stacked id embeddings.
|
||||
--input-id-images-dir [DIR] path to PHOTOMAKER input id images dir.
|
||||
--normalize-input normalize PHOTOMAKER input id images
|
||||
--upscale-model [ESRGAN_PATH] path to esrgan model. Upscale images after generate, just RealESRGAN_x4plus_anime_6B supported by now.
|
||||
--upscale-repeats Run the ESRGAN upscaler this many times (default 1)
|
||||
--type [TYPE] weight type (f32, f16, q4_0, q4_1, q5_0, q5_1, q8_0)
|
||||
@@ -172,6 +179,7 @@ arguments:
|
||||
-n, --negative-prompt PROMPT the negative prompt (default: "")
|
||||
--cfg-scale SCALE unconditional guidance scale: (default: 7.0)
|
||||
--strength STRENGTH strength for noising/unnoising (default: 0.75)
|
||||
--style-ratio STYLE-RATIO strength for keeping input identity (default: 20%)
|
||||
--control-strength STRENGTH strength to apply Control Net (default: 0.9)
|
||||
1.0 corresponds to full destruction of information in init image
|
||||
-H, --height H image height, in pixel space (default: 512)
|
||||
@@ -182,12 +190,13 @@ arguments:
|
||||
--rng {std_default, cuda} RNG (default: cuda)
|
||||
-s SEED, --seed SEED RNG seed (default: 42, use random seed for < 0)
|
||||
-b, --batch-count COUNT number of images to generate.
|
||||
--schedule {discrete, karras} Denoiser sigma schedule (default: discrete)
|
||||
--schedule {discrete, karras, ays} Denoiser sigma schedule (default: discrete)
|
||||
--clip-skip N ignore last layers of CLIP network; 1 ignores none, 2 ignores one layer (default: -1)
|
||||
<= 0 represents unspecified, will be 1 for SD1.x, 2 for SD2.x
|
||||
--vae-tiling process vae in tiles to reduce memory usage
|
||||
--control-net-cpu keep controlnet in cpu (for low vram)
|
||||
--canny apply canny preprocessor (edge detection)
|
||||
--color colors the logging tags according to level
|
||||
-v, --verbose print extra info
|
||||
```
|
||||
|
||||
@@ -296,6 +305,39 @@ You can use ESRGAN to upscale the generated images. At the moment, only the [Rea
|
||||
sd -m ../models/v1-5-pruned-emaonly.safetensors -p "a lovely cat" --upscale-model ../models/RealESRGAN_x4plus_anime_6B.pth
|
||||
```
|
||||
|
||||
#### Using PhotoMaker to personalize image generation
|
||||
|
||||
You can use [PhotoMaker](https://github.com/TencentARC/PhotoMaker) to personalize generated images with your own ID.
|
||||
|
||||
**NOTE**, currently PhotoMaker **ONLY** works with **SDXL** (any SDXL model files will work).
|
||||
|
||||
Download PhotoMaker model file (in safetensor format) [here](https://huggingface.co/bssrdf/PhotoMaker). The official release of the model file (in .bin format) does not work with ```stablediffusion.cpp```.
|
||||
|
||||
- Specify the PhotoMaker model path using the `--stacked-id-embd-dir PATH` parameter.
|
||||
- Specify the input images path using the `--input-id-images-dir PATH` parameter.
|
||||
- input images **must** have the same width and height for preprocessing (to be improved)
|
||||
|
||||
In prompt, make sure you have a class word followed by the trigger word ```"img"``` (hard-coded for now). The class word could be one of ```"man, woman, girl, boy"```. If input ID images contain asian faces, add ```Asian``` before the class
|
||||
word.
|
||||
|
||||
Another PhotoMaker specific parameter:
|
||||
|
||||
- ```--style-ratio (0-100)%```: default is 20 and 10-20 typically gets good results. Lower ratio means more faithfully following input ID (not necessarily better quality).
|
||||
|
||||
Other parameters recommended for running Photomaker:
|
||||
|
||||
- ```--cfg-scale 5.0```
|
||||
- ```-H 1024```
|
||||
- ```-W 1024```
|
||||
|
||||
If on low memory GPUs (<= 8GB), recommend running with ```--vae-on-cpu``` option to get artifact free images.
|
||||
|
||||
Example:
|
||||
|
||||
```bash
|
||||
bin/sd -m ../models/sdxlUnstableDiffusers_v11.safetensors --vae ../models/sdxl_vae.safetensors --stacked-id-embd-dir ../models/photomaker-v1.safetensors --input-id-images-dir ../assets/examples/scarletthead_woman -p "a girl img, retro futurism, retro game art style but extremely beautiful, intricate details, masterpiece, best quality, space-themed, cosmic, celestial, stars, galaxies, nebulas, planets, science fiction, highly detailed" -n "realistic, photo-realistic, worst quality, greyscale, bad anatomy, bad hands, error, text" --cfg-scale 5.0 --sampling-method euler -H 1024 -W 1024 --style-ratio 10 --vae-on-cpu -o output.png
|
||||
```
|
||||
|
||||
### Docker
|
||||
|
||||
#### Building using Docker
|
||||
@@ -319,6 +361,19 @@ docker run -v /path/to/models:/models -v /path/to/output/:/output sd [args...]
|
||||
| **Memory** (txt2img - 512 x 512) | ~2.8G | ~2.3G | ~2.1G | ~2.0G | ~2.0G | ~2.0G | ~2.0G |
|
||||
| **Memory** (txt2img - 512 x 512) *with Flash Attention* | ~2.4G | ~1.9G | ~1.6G | ~1.5G | ~1.5G | ~1.5G | ~1.5G |
|
||||
|
||||
## Bindings
|
||||
|
||||
These projects wrap `stable-diffusion.cpp` for easier use in other languages/frameworks.
|
||||
|
||||
* Golang: [seasonjs/stable-diffusion](https://github.com/seasonjs/stable-diffusion)
|
||||
* C#: [DarthAffe/StableDiffusion.NET](https://github.com/DarthAffe/StableDiffusion.NET)
|
||||
|
||||
## UIs
|
||||
|
||||
These projects use `stable-diffusion.cpp` as a backend for their image generation.
|
||||
|
||||
- [Jellybox](https://jellybox.com)
|
||||
|
||||
## Contributors
|
||||
|
||||
Thank you to all the people who have already contributed to stable-diffusion.cpp!
|
||||
@@ -335,3 +390,4 @@ Thank you to all the people who have already contributed to stable-diffusion.cpp
|
||||
- [k-diffusion](https://github.com/crowsonkb/k-diffusion)
|
||||
- [latent-consistency-model](https://github.com/luosiallen/latent-consistency-model)
|
||||
- [generative-models](https://github.com/Stability-AI/generative-models/)
|
||||
- [PhotoMaker](https://github.com/TencentARC/PhotoMaker)
|
||||
|
||||
BIN
assets/photomaker_examples/lenna_woman/lenna.jpg
Normal file
|
After Width: | Height: | Size: 39 KiB |
BIN
assets/photomaker_examples/newton_man/newton_0.jpg
Normal file
|
After Width: | Height: | Size: 311 KiB |
BIN
assets/photomaker_examples/newton_man/newton_1.jpg
Normal file
|
After Width: | Height: | Size: 53 KiB |
BIN
assets/photomaker_examples/newton_man/newton_2.png
Normal file
|
After Width: | Height: | Size: 1.4 MiB |
BIN
assets/photomaker_examples/newton_man/newton_3.jpg
Normal file
|
After Width: | Height: | Size: 26 KiB |
BIN
assets/photomaker_examples/scarletthead_woman/scarlett_0.jpg
Normal file
|
After Width: | Height: | Size: 88 KiB |
BIN
assets/photomaker_examples/scarletthead_woman/scarlett_1.jpg
Normal file
|
After Width: | Height: | Size: 26 KiB |
BIN
assets/photomaker_examples/scarletthead_woman/scarlett_2.jpg
Normal file
|
After Width: | Height: | Size: 72 KiB |
BIN
assets/photomaker_examples/scarletthead_woman/scarlett_3.jpg
Normal file
|
After Width: | Height: | Size: 107 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_1.jpg
Normal file
|
After Width: | Height: | Size: 54 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_2.jpeg
Normal file
|
After Width: | Height: | Size: 68 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_3.jpg
Normal file
|
After Width: | Height: | Size: 54 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_4.jpg
Normal file
|
After Width: | Height: | Size: 48 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_5.jpg
Normal file
|
After Width: | Height: | Size: 28 KiB |
BIN
assets/photomaker_examples/yangmi_woman/yangmi_6.jpg
Normal file
|
After Width: | Height: | Size: 30 KiB |
430
clip.hpp
@@ -75,9 +75,13 @@ class CLIPTokenizer {
|
||||
private:
|
||||
SDVersion version = VERSION_1_x;
|
||||
std::map<int, std::u32string> byte_encoder;
|
||||
std::map<std::u32string, int> byte_decoder;
|
||||
std::map<std::u32string, int> encoder;
|
||||
std::map<int, std::u32string> decoder;
|
||||
std::map<std::pair<std::u32string, std::u32string>, int> bpe_ranks;
|
||||
std::regex pat;
|
||||
int encoder_len;
|
||||
int bpe_len;
|
||||
|
||||
static std::string strip(const std::string& str) {
|
||||
std::string::size_type start = str.find_first_not_of(" \t\n\r\v\f");
|
||||
@@ -118,7 +122,11 @@ public:
|
||||
|
||||
void load_from_merges(const std::string& merges_utf8_str) {
|
||||
auto byte_unicode_pairs = bytes_to_unicode();
|
||||
byte_encoder = std::map<int, std::u32string>(byte_unicode_pairs.begin(), byte_unicode_pairs.end());
|
||||
// printf("byte_unicode_pairs have %lu pairs \n", byte_unicode_pairs.size());
|
||||
byte_encoder = std::map<int, std::u32string>(byte_unicode_pairs.begin(), byte_unicode_pairs.end());
|
||||
for (auto& pair : byte_unicode_pairs) {
|
||||
byte_decoder[pair.second] = pair.first;
|
||||
}
|
||||
// for (auto & pair: byte_unicode_pairs) {
|
||||
// std::cout << pair.first << ": " << pair.second << std::endl;
|
||||
// }
|
||||
@@ -138,6 +146,8 @@ public:
|
||||
size_t space_pos = merge.find(' ');
|
||||
merge_pairs.emplace_back(merge.substr(0, space_pos), merge.substr(space_pos + 1));
|
||||
// LOG_DEBUG("%s", utf32_to_utf8(merge.substr(space_pos + 1)).c_str());
|
||||
// printf("%s :: %s | %s \n", utf32_to_utf8(merge).c_str(), utf32_to_utf8(merge.substr(0, space_pos)).c_str(),
|
||||
// utf32_to_utf8(merge.substr(space_pos + 1)).c_str());
|
||||
}
|
||||
std::vector<std::u32string> vocab;
|
||||
for (const auto& pair : byte_unicode_pairs) {
|
||||
@@ -154,15 +164,36 @@ public:
|
||||
LOG_DEBUG("vocab size: %llu", vocab.size());
|
||||
int i = 0;
|
||||
for (const auto& token : vocab) {
|
||||
encoder[token] = i++;
|
||||
encoder[token] = i;
|
||||
decoder[i] = token;
|
||||
i++;
|
||||
}
|
||||
encoder_len = i;
|
||||
|
||||
auto it = encoder.find(utf8_to_utf32("img</w>"));
|
||||
if (it != encoder.end()) {
|
||||
LOG_DEBUG(" trigger word img already in vocab");
|
||||
} else {
|
||||
LOG_DEBUG(" trigger word img not in vocab yet");
|
||||
}
|
||||
|
||||
int rank = 0;
|
||||
for (const auto& merge : merge_pairs) {
|
||||
bpe_ranks[merge] = rank++;
|
||||
}
|
||||
bpe_len = rank;
|
||||
};
|
||||
|
||||
void add_token(const std::string& text) {
|
||||
std::u32string token = utf8_to_utf32(text);
|
||||
auto it = encoder.find(token);
|
||||
if (it != encoder.end()) {
|
||||
encoder[token] = encoder_len;
|
||||
decoder[encoder_len] = token;
|
||||
encoder_len++;
|
||||
}
|
||||
}
|
||||
|
||||
std::u32string bpe(const std::u32string& token) {
|
||||
std::vector<std::u32string> word;
|
||||
|
||||
@@ -243,6 +274,7 @@ public:
|
||||
size_t max_length = 0,
|
||||
bool padding = false) {
|
||||
std::vector<int32_t> tokens = encode(text, on_new_token_cb);
|
||||
|
||||
tokens.insert(tokens.begin(), BOS_TOKEN_ID);
|
||||
if (max_length > 0) {
|
||||
if (tokens.size() > max_length - 1) {
|
||||
@@ -259,9 +291,34 @@ public:
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return tokens;
|
||||
}
|
||||
|
||||
std::string decode(const std::vector<int>& tokens) {
|
||||
std::string text = "";
|
||||
for (int t : tokens) {
|
||||
if (t == 49406 || t == 49407)
|
||||
continue;
|
||||
std::u32string ts = decoder[t];
|
||||
// printf("%d, %s \n", t, utf32_to_utf8(ts).c_str());
|
||||
std::string s = utf32_to_utf8(ts);
|
||||
if (s.length() >= 4 && ends_with(s, "</w>")) {
|
||||
text += " " + s.replace(s.length() - 4, s.length() - 1, "");
|
||||
} else {
|
||||
text += " " + s;
|
||||
}
|
||||
}
|
||||
// std::vector<unsigned char> bytes;
|
||||
// for (auto c : text){
|
||||
// bytes.push_back(byte_decoder[c]);
|
||||
// }
|
||||
|
||||
// std::string s((char *)bytes.data());
|
||||
// std::string s = "";
|
||||
return trim(text);
|
||||
}
|
||||
|
||||
std::vector<int> encode(std::string text, on_new_token_cb_t on_new_token_cb) {
|
||||
std::string original_text = text;
|
||||
std::vector<int32_t> bpe_tokens;
|
||||
@@ -308,7 +365,8 @@ public:
|
||||
ss << "\"" << token << "\", ";
|
||||
}
|
||||
ss << "]";
|
||||
LOG_DEBUG("split prompt \"%s\" to tokens %s", original_text.c_str(), ss.str().c_str());
|
||||
// LOG_DEBUG("split prompt \"%s\" to tokens %s", original_text.c_str(), ss.str().c_str());
|
||||
// printf("split prompt \"%s\" to tokens %s \n", original_text.c_str(), ss.str().c_str());
|
||||
return bpe_tokens;
|
||||
}
|
||||
};
|
||||
@@ -469,7 +527,8 @@ public:
|
||||
: d_model(d_model),
|
||||
n_head(n_head),
|
||||
intermediate_size(intermediate_size) {
|
||||
blocks["self_attn"] = std::shared_ptr<GGMLBlock>(new MultiheadAttention(d_model, n_head));
|
||||
blocks["self_attn"] = std::shared_ptr<GGMLBlock>(new MultiheadAttention(d_model, n_head, true));
|
||||
|
||||
blocks["layer_norm1"] = std::shared_ptr<GGMLBlock>(new LayerNorm(d_model));
|
||||
blocks["layer_norm2"] = std::shared_ptr<GGMLBlock>(new LayerNorm(d_model));
|
||||
|
||||
@@ -508,7 +567,7 @@ public:
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* x, int clip_skip = -1, bool mask = true) {
|
||||
// x: [N, n_token, d_model]
|
||||
int layer_idx = n_layer - 1;
|
||||
LOG_DEBUG("clip_skip %d", clip_skip);
|
||||
// LOG_DEBUG("clip_skip %d", clip_skip);
|
||||
if (clip_skip > 0) {
|
||||
layer_idx = n_layer - clip_skip;
|
||||
}
|
||||
@@ -520,7 +579,7 @@ public:
|
||||
}
|
||||
std::string name = "layers." + std::to_string(i);
|
||||
auto layer = std::dynamic_pointer_cast<CLIPLayer>(blocks[name]);
|
||||
x = layer->forward(ctx, x); // [N, n_token, d_model]
|
||||
x = layer->forward(ctx, x, mask); // [N, n_token, d_model]
|
||||
// LOG_DEBUG("layer %d", i);
|
||||
}
|
||||
return x;
|
||||
@@ -620,8 +679,8 @@ public:
|
||||
class_embedding = ggml_repeat(ctx, class_embed_weight, class_embedding); // [N, embed_dim]
|
||||
class_embedding = ggml_reshape_4d(ctx, class_embedding, 1, embed_dim, 1, N); // [N, 1, embed_dim, 1]
|
||||
|
||||
struct ggml_tensor* x = ggml_concat(ctx, class_embedding, patch_embedding); // [N, num_positions, embed_dim, 1]
|
||||
x = ggml_reshape_3d(ctx, x, embed_dim, num_positions, N); // [N, num_positions, embed_dim]
|
||||
struct ggml_tensor* x = ggml_concat(ctx, class_embedding, patch_embedding, 2); // [N, num_positions, embed_dim, 1]
|
||||
x = ggml_reshape_3d(ctx, x, embed_dim, num_positions, N); // [N, num_positions, embed_dim]
|
||||
x = ggml_add(ctx, x, position_embed_weight);
|
||||
return x; // [N, num_positions, embed_dim]
|
||||
}
|
||||
@@ -703,7 +762,7 @@ public:
|
||||
auto final_layer_norm = std::dynamic_pointer_cast<LayerNorm>(blocks["final_layer_norm"]);
|
||||
|
||||
auto x = embeddings->forward(ctx, input_ids, tkn_embeddings); // [N, n_token, hidden_size]
|
||||
x = encoder->forward(ctx, x, return_pooled ? -1 : clip_skip, true);
|
||||
x = encoder->forward(ctx, x, return_pooled ? -1 : clip_skip, true);
|
||||
if (return_pooled || with_final_ln) {
|
||||
x = final_layer_norm->forward(ctx, x);
|
||||
}
|
||||
@@ -720,11 +779,6 @@ public:
|
||||
};
|
||||
|
||||
class CLIPVisionModel : public GGMLBlock {
|
||||
protected:
|
||||
void init_params(struct ggml_context* ctx, ggml_type wtype) {
|
||||
params["visual_projection"] = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, projection_dim, hidden_size);
|
||||
}
|
||||
|
||||
public:
|
||||
// network hparams
|
||||
int32_t num_channels = 3;
|
||||
@@ -735,16 +789,14 @@ public:
|
||||
int32_t intermediate_size = 4096;
|
||||
int32_t n_head = 16;
|
||||
int32_t n_layer = 24;
|
||||
int32_t projection_dim = 768;
|
||||
|
||||
public:
|
||||
CLIPVisionModel(CLIPVersion version = OPEN_CLIP_VIT_H_14) {
|
||||
CLIPVisionModel(CLIPVersion version = OPENAI_CLIP_VIT_L_14) {
|
||||
if (version == OPEN_CLIP_VIT_H_14) {
|
||||
hidden_size = 1280;
|
||||
intermediate_size = 5120;
|
||||
n_head = 16;
|
||||
n_layer = 32;
|
||||
projection_dim = 1024;
|
||||
} else if (version == OPEN_CLIP_VIT_BIGG_14) {
|
||||
hidden_size = 1664;
|
||||
intermediate_size = 8192;
|
||||
@@ -758,9 +810,8 @@ public:
|
||||
blocks["post_layernorm"] = std::shared_ptr<GGMLBlock>(new LayerNorm(hidden_size));
|
||||
}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* pixel_values) {
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* pixel_values, bool return_pooled = true) {
|
||||
// pixel_values: [N, num_channels, image_size, image_size]
|
||||
// return: // [N, projection_dim]
|
||||
auto embeddings = std::dynamic_pointer_cast<CLIPVisionEmbeddings>(blocks["embeddings"]);
|
||||
auto pre_layernorm = std::dynamic_pointer_cast<LayerNorm>(blocks["pre_layernorm"]);
|
||||
auto encoder = std::dynamic_pointer_cast<CLIPEncoder>(blocks["encoder"]);
|
||||
@@ -768,26 +819,60 @@ public:
|
||||
|
||||
auto x = embeddings->forward(ctx, pixel_values); // [N, num_positions, embed_dim]
|
||||
x = pre_layernorm->forward(ctx, x);
|
||||
x = encoder->forward(ctx, x, -1, true);
|
||||
x = encoder->forward(ctx, x, -1, false);
|
||||
x = post_layernorm->forward(ctx, x); // [N, n_token, hidden_size]
|
||||
|
||||
GGML_ASSERT(x->ne[2] == 1);
|
||||
int64_t max_token_idx = 0;
|
||||
ggml_tensor* pooled = ggml_view_1d(ctx, x, x->ne[0], x->nb[1] * max_token_idx); // assert N == 1
|
||||
auto visual_projection = params["visual_projection"];
|
||||
pooled = ggml_mul_mat(ctx, ggml_cont(ctx, ggml_transpose(ctx, visual_projection)), pooled);
|
||||
return pooled; // [N, projection_dim]
|
||||
GGML_ASSERT(x->ne[3] == 1);
|
||||
if (return_pooled) {
|
||||
ggml_tensor* pooled = ggml_cont(ctx, ggml_view_2d(ctx, x, x->ne[0], x->ne[2], x->nb[2], 0));
|
||||
return pooled; // [N, hidden_size]
|
||||
} else {
|
||||
return x; // [N, n_token, hidden_size]
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
class CLIPProjection : public UnaryBlock {
|
||||
protected:
|
||||
int64_t in_features;
|
||||
int64_t out_features;
|
||||
bool transpose_weight;
|
||||
|
||||
void init_params(struct ggml_context* ctx, ggml_type wtype) {
|
||||
if (transpose_weight) {
|
||||
LOG_ERROR("transpose_weight");
|
||||
params["weight"] = ggml_new_tensor_2d(ctx, wtype, out_features, in_features);
|
||||
} else {
|
||||
params["weight"] = ggml_new_tensor_2d(ctx, wtype, in_features, out_features);
|
||||
}
|
||||
}
|
||||
|
||||
public:
|
||||
CLIPProjection(int64_t in_features,
|
||||
int64_t out_features,
|
||||
bool transpose_weight = false)
|
||||
: in_features(in_features),
|
||||
out_features(out_features),
|
||||
transpose_weight(transpose_weight) {}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* x) {
|
||||
struct ggml_tensor* w = params["weight"];
|
||||
if (transpose_weight) {
|
||||
w = ggml_cont(ctx, ggml_transpose(ctx, w));
|
||||
}
|
||||
return ggml_nn_linear(ctx, x, w, NULL);
|
||||
}
|
||||
};
|
||||
|
||||
class CLIPVisionModelProjection : public GGMLBlock {
|
||||
public:
|
||||
int32_t hidden_size = 1024;
|
||||
int32_t projection_dim = 1024;
|
||||
int32_t projection_dim = 768;
|
||||
int32_t image_size = 224;
|
||||
|
||||
public:
|
||||
CLIPVisionModelProjection(CLIPVersion version = OPEN_CLIP_VIT_H_14) {
|
||||
CLIPVisionModelProjection(CLIPVersion version = OPENAI_CLIP_VIT_L_14,
|
||||
bool transpose_proj_w = false) {
|
||||
if (version == OPEN_CLIP_VIT_H_14) {
|
||||
hidden_size = 1280;
|
||||
projection_dim = 1024;
|
||||
@@ -795,17 +880,17 @@ public:
|
||||
hidden_size = 1664;
|
||||
}
|
||||
|
||||
blocks["visual_model"] = std::shared_ptr<GGMLBlock>(new CLIPVisionModel(version));
|
||||
blocks["visual_projection"] = std::shared_ptr<GGMLBlock>(new Linear(hidden_size, projection_dim, false));
|
||||
blocks["vision_model"] = std::shared_ptr<GGMLBlock>(new CLIPVisionModel(version));
|
||||
blocks["visual_projection"] = std::shared_ptr<GGMLBlock>(new CLIPProjection(hidden_size, projection_dim, transpose_proj_w));
|
||||
}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* pixel_values) {
|
||||
// pixel_values: [N, num_channels, image_size, image_size]
|
||||
// return: [N, num_positions, projection_dim]
|
||||
auto visual_model = std::dynamic_pointer_cast<CLIPVisionModel>(blocks["visual_model"]);
|
||||
auto visual_projection = std::dynamic_pointer_cast<Linear>(blocks["visual_projection"]);
|
||||
// return: [N, projection_dim]
|
||||
auto vision_model = std::dynamic_pointer_cast<CLIPVisionModel>(blocks["vision_model"]);
|
||||
auto visual_projection = std::dynamic_pointer_cast<CLIPProjection>(blocks["visual_projection"]);
|
||||
|
||||
auto x = visual_model->forward(ctx, pixel_values); // [N, embed_dim]
|
||||
auto x = vision_model->forward(ctx, pixel_values); // [N, hidden_size]
|
||||
x = visual_projection->forward(ctx, x); // [N, projection_dim]
|
||||
|
||||
return x; // [N, projection_dim]
|
||||
@@ -854,22 +939,6 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
return "clip";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
size_t params_mem_size = text_model.get_params_mem_size();
|
||||
if (version == VERSION_XL) {
|
||||
params_mem_size += text_model2.get_params_mem_size();
|
||||
}
|
||||
return params_mem_size;
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
size_t params_num = text_model.get_params_num();
|
||||
if (version == VERSION_XL) {
|
||||
params_num += text_model2.get_params_num();
|
||||
}
|
||||
return params_num;
|
||||
}
|
||||
|
||||
void set_clip_skip(int clip_skip) {
|
||||
text_model.set_clip_skip(clip_skip);
|
||||
if (version == VERSION_XL) {
|
||||
@@ -967,7 +1036,7 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
hidden_states2->ne[3]);
|
||||
hidden_states2 = ggml_cont(ctx, ggml_permute(ctx, hidden_states2, 2, 0, 1, 3));
|
||||
|
||||
hidden_states = ggml_concat(ctx, hidden_states, hidden_states2); // [N, n_token, hidden_size + hidden_size2]
|
||||
hidden_states = ggml_concat(ctx, hidden_states, hidden_states2, 2); // [N, n_token, hidden_size + hidden_size2]
|
||||
|
||||
hidden_states = ggml_cont(ctx, ggml_permute(ctx, hidden_states, 1, 2, 0, 3));
|
||||
}
|
||||
@@ -1000,7 +1069,7 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
auto token_embed_weight = text_model.get_token_embed_weight();
|
||||
token_embed_weight = ggml_reshape_3d(compute_ctx, token_embed_weight, token_embed_weight->ne[0], 1, token_embed_weight->ne[1]);
|
||||
// concatenate custom embeddings
|
||||
embeddings = ggml_concat(compute_ctx, token_embed_weight, custom_embeddings);
|
||||
embeddings = ggml_concat(compute_ctx, token_embed_weight, custom_embeddings, 2);
|
||||
embeddings = ggml_reshape_2d(compute_ctx, embeddings, embeddings->ne[0], embeddings->ne[2]);
|
||||
}
|
||||
|
||||
@@ -1029,6 +1098,205 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
return tokenize(text, text_model.n_token, padding);
|
||||
}
|
||||
|
||||
std::tuple<std::vector<int>, std::vector<float>, std::vector<bool>>
|
||||
tokenize_with_trigger_token(std::string text,
|
||||
int num_input_imgs,
|
||||
int32_t image_token,
|
||||
bool padding = false) {
|
||||
return tokenize_with_trigger_token(text, num_input_imgs, image_token,
|
||||
text_model.n_token, padding);
|
||||
}
|
||||
|
||||
std::vector<int> convert_token_to_id(std::string text) {
|
||||
auto on_new_token_cb = [&](std::string& str, std::vector<int32_t>& bpe_tokens) -> bool {
|
||||
size_t word_end = str.find(",");
|
||||
std::string embd_name = word_end == std::string::npos ? str : str.substr(0, word_end);
|
||||
embd_name = trim(embd_name);
|
||||
std::string embd_path = get_full_path(embd_dir, embd_name + ".pt");
|
||||
if (embd_path.size() == 0) {
|
||||
embd_path = get_full_path(embd_dir, embd_name + ".ckpt");
|
||||
}
|
||||
if (embd_path.size() == 0) {
|
||||
embd_path = get_full_path(embd_dir, embd_name + ".safetensors");
|
||||
}
|
||||
if (embd_path.size() > 0) {
|
||||
if (load_embedding(embd_name, embd_path, bpe_tokens)) {
|
||||
if (word_end != std::string::npos) {
|
||||
str = str.substr(word_end);
|
||||
} else {
|
||||
str = "";
|
||||
}
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
};
|
||||
std::vector<int> curr_tokens = tokenizer.encode(text, on_new_token_cb);
|
||||
return curr_tokens;
|
||||
}
|
||||
|
||||
std::string decode(const std::vector<int>& tokens) {
|
||||
return tokenizer.decode(tokens);
|
||||
}
|
||||
|
||||
void pad_tokens(std::vector<int>& tokens,
|
||||
std::vector<float>& weights,
|
||||
size_t max_length = 0,
|
||||
bool padding = false) {
|
||||
if (max_length > 0 && padding) {
|
||||
size_t n = std::ceil(tokens.size() * 1.0 / (max_length - 2));
|
||||
if (n == 0) {
|
||||
n = 1;
|
||||
}
|
||||
size_t length = max_length * n;
|
||||
LOG_DEBUG("token length: %llu", length);
|
||||
std::vector<int> new_tokens;
|
||||
std::vector<float> new_weights;
|
||||
new_tokens.push_back(BOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
int token_idx = 0;
|
||||
for (int i = 1; i < length; i++) {
|
||||
if (token_idx >= tokens.size()) {
|
||||
break;
|
||||
}
|
||||
if (i % max_length == 0) {
|
||||
new_tokens.push_back(BOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
} else if (i % max_length == max_length - 1) {
|
||||
new_tokens.push_back(EOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
} else {
|
||||
new_tokens.push_back(tokens[token_idx]);
|
||||
new_weights.push_back(weights[token_idx]);
|
||||
token_idx++;
|
||||
}
|
||||
}
|
||||
|
||||
new_tokens.push_back(EOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
tokens = new_tokens;
|
||||
weights = new_weights;
|
||||
|
||||
if (padding) {
|
||||
int pad_token_id = PAD_TOKEN_ID;
|
||||
if (version == VERSION_2_x) {
|
||||
pad_token_id = 0;
|
||||
}
|
||||
tokens.insert(tokens.end(), length - tokens.size(), pad_token_id);
|
||||
weights.insert(weights.end(), length - weights.size(), 1.0);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::tuple<std::vector<int>, std::vector<float>, std::vector<bool>>
|
||||
tokenize_with_trigger_token(std::string text,
|
||||
int num_input_imgs,
|
||||
int32_t image_token,
|
||||
size_t max_length = 0,
|
||||
bool padding = false) {
|
||||
auto parsed_attention = parse_prompt_attention(text);
|
||||
|
||||
{
|
||||
std::stringstream ss;
|
||||
ss << "[";
|
||||
for (const auto& item : parsed_attention) {
|
||||
ss << "['" << item.first << "', " << item.second << "], ";
|
||||
}
|
||||
ss << "]";
|
||||
LOG_DEBUG("parse '%s' to %s", text.c_str(), ss.str().c_str());
|
||||
}
|
||||
|
||||
auto on_new_token_cb = [&](std::string& str, std::vector<int32_t>& bpe_tokens) -> bool {
|
||||
size_t word_end = str.find(",");
|
||||
std::string embd_name = word_end == std::string::npos ? str : str.substr(0, word_end);
|
||||
embd_name = trim(embd_name);
|
||||
std::string embd_path = get_full_path(embd_dir, embd_name + ".pt");
|
||||
if (embd_path.size() == 0) {
|
||||
embd_path = get_full_path(embd_dir, embd_name + ".ckpt");
|
||||
}
|
||||
if (embd_path.size() == 0) {
|
||||
embd_path = get_full_path(embd_dir, embd_name + ".safetensors");
|
||||
}
|
||||
if (embd_path.size() > 0) {
|
||||
if (load_embedding(embd_name, embd_path, bpe_tokens)) {
|
||||
if (word_end != std::string::npos) {
|
||||
str = str.substr(word_end);
|
||||
} else {
|
||||
str = "";
|
||||
}
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
};
|
||||
|
||||
std::vector<int> tokens;
|
||||
std::vector<float> weights;
|
||||
std::vector<bool> class_token_mask;
|
||||
int32_t class_idx = -1, tokens_acc = 0;
|
||||
for (const auto& item : parsed_attention) {
|
||||
std::vector<int> class_token_index;
|
||||
std::vector<int> clean_input_ids;
|
||||
const std::string& curr_text = item.first;
|
||||
float curr_weight = item.second;
|
||||
// printf(" %s: %f \n", curr_text.c_str(), curr_weight);
|
||||
std::vector<int> curr_tokens = tokenizer.encode(curr_text, on_new_token_cb);
|
||||
int32_t clean_index = 0;
|
||||
for (uint32_t i = 0; i < curr_tokens.size(); i++) {
|
||||
int token_id = curr_tokens[i];
|
||||
if (token_id == image_token)
|
||||
class_token_index.push_back(clean_index - 1);
|
||||
else {
|
||||
clean_input_ids.push_back(token_id);
|
||||
clean_index++;
|
||||
}
|
||||
}
|
||||
// GGML_ASSERT(class_token_index.size() == 1); // PhotoMaker currently does not support multiple
|
||||
// trigger words in a single prompt.
|
||||
if (class_token_index.size() == 1) {
|
||||
// Expand the class word token and corresponding mask
|
||||
int class_token = clean_input_ids[class_token_index[0]];
|
||||
class_idx = tokens_acc + class_token_index[0];
|
||||
std::vector<int> clean_input_ids_tmp;
|
||||
for (uint32_t i = 0; i < class_token_index[0]; i++)
|
||||
clean_input_ids_tmp.push_back(clean_input_ids[i]);
|
||||
for (uint32_t i = 0; i < num_input_imgs; i++)
|
||||
clean_input_ids_tmp.push_back(class_token);
|
||||
for (uint32_t i = class_token_index[0] + 1; i < clean_input_ids.size(); i++)
|
||||
clean_input_ids_tmp.push_back(clean_input_ids[i]);
|
||||
clean_input_ids.clear();
|
||||
clean_input_ids = clean_input_ids_tmp;
|
||||
}
|
||||
tokens_acc += clean_index;
|
||||
tokens.insert(tokens.end(), clean_input_ids.begin(), clean_input_ids.end());
|
||||
weights.insert(weights.end(), clean_input_ids.size(), curr_weight);
|
||||
}
|
||||
tokens.insert(tokens.begin(), BOS_TOKEN_ID);
|
||||
weights.insert(weights.begin(), 1.0);
|
||||
|
||||
pad_tokens(tokens, weights, max_length, padding);
|
||||
|
||||
for (uint32_t i = 0; i < tokens.size(); i++) {
|
||||
if (class_idx + 1 <= i && i < class_idx + 1 + num_input_imgs)
|
||||
class_token_mask.push_back(true);
|
||||
else
|
||||
class_token_mask.push_back(false);
|
||||
}
|
||||
|
||||
// printf("[");
|
||||
// for (int i = 0; i < tokens.size(); i++) {
|
||||
// printf("%d, ", class_token_mask[i] ? 1 : 0);
|
||||
// }
|
||||
// printf("]\n");
|
||||
|
||||
// for (int i = 0; i < tokens.size(); i++) {
|
||||
// std::cout << tokens[i] << ":" << weights[i] << ", ";
|
||||
// }
|
||||
// std::cout << std::endl;
|
||||
|
||||
return std::make_tuple(tokens, weights, class_token_mask);
|
||||
}
|
||||
|
||||
std::pair<std::vector<int>, std::vector<float>> tokenize(std::string text,
|
||||
size_t max_length = 0,
|
||||
bool padding = false) {
|
||||
@@ -1078,49 +1346,7 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
weights.insert(weights.end(), curr_tokens.size(), curr_weight);
|
||||
}
|
||||
|
||||
if (max_length > 0 && padding) {
|
||||
size_t n = std::ceil(tokens.size() * 1.0 / (max_length - 2));
|
||||
if (n == 0) {
|
||||
n = 1;
|
||||
}
|
||||
size_t length = max_length * n;
|
||||
LOG_DEBUG("token length: %llu", length);
|
||||
std::vector<int> new_tokens;
|
||||
std::vector<float> new_weights;
|
||||
new_tokens.push_back(BOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
int token_idx = 0;
|
||||
for (int i = 1; i < length; i++) {
|
||||
if (token_idx >= tokens.size()) {
|
||||
break;
|
||||
}
|
||||
if (i % max_length == 0) {
|
||||
new_tokens.push_back(BOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
} else if (i % max_length == max_length - 1) {
|
||||
new_tokens.push_back(EOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
} else {
|
||||
new_tokens.push_back(tokens[token_idx]);
|
||||
new_weights.push_back(weights[token_idx]);
|
||||
token_idx++;
|
||||
}
|
||||
}
|
||||
|
||||
new_tokens.push_back(EOS_TOKEN_ID);
|
||||
new_weights.push_back(1.0);
|
||||
tokens = new_tokens;
|
||||
weights = new_weights;
|
||||
|
||||
if (padding) {
|
||||
int pad_token_id = PAD_TOKEN_ID;
|
||||
if (version == VERSION_2_x) {
|
||||
pad_token_id = 0;
|
||||
}
|
||||
tokens.insert(tokens.end(), length - tokens.size(), pad_token_id);
|
||||
weights.insert(weights.end(), length - weights.size(), 1.0);
|
||||
}
|
||||
}
|
||||
pad_tokens(tokens, weights, max_length, padding);
|
||||
|
||||
// for (int i = 0; i < tokens.size(); i++) {
|
||||
// std::cout << tokens[i] << ":" << weights[i] << ", ";
|
||||
@@ -1132,10 +1358,10 @@ struct FrozenCLIPEmbedderWithCustomWords : public GGMLModule {
|
||||
};
|
||||
|
||||
struct FrozenCLIPVisionEmbedder : public GGMLModule {
|
||||
CLIPVisionModel vision_model;
|
||||
CLIPVisionModelProjection vision_model;
|
||||
|
||||
FrozenCLIPVisionEmbedder(ggml_backend_t backend, ggml_type wtype)
|
||||
: GGMLModule(backend, wtype) {
|
||||
: vision_model(OPEN_CLIP_VIT_H_14, true), GGMLModule(backend, wtype) {
|
||||
vision_model.init(params_ctx, wtype);
|
||||
}
|
||||
|
||||
@@ -1143,16 +1369,8 @@ struct FrozenCLIPVisionEmbedder : public GGMLModule {
|
||||
return "clip_vision";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return vision_model.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return vision_model.get_params_num();
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, struct ggml_tensor*>& tensors, const std::string prefix) {
|
||||
vision_model.get_param_tensors(tensors, prefix + "transformer.visual_model");
|
||||
vision_model.get_param_tensors(tensors, prefix + "transformer");
|
||||
}
|
||||
|
||||
struct ggml_cgraph* build_graph(struct ggml_tensor* pixel_values) {
|
||||
|
||||
@@ -369,14 +369,6 @@ struct ControlNet : public GGMLModule {
|
||||
return "control_net";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return control_net.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return control_net.get_params_num();
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, struct ggml_tensor*>& tensors, const std::string prefix) {
|
||||
control_net.get_param_tensors(tensors, prefix);
|
||||
}
|
||||
|
||||
538
denoiser.hpp
@@ -13,6 +13,7 @@ struct SigmaSchedule {
|
||||
float alphas_cumprod[TIMESTEPS];
|
||||
float sigmas[TIMESTEPS];
|
||||
float log_sigmas[TIMESTEPS];
|
||||
int version = 0;
|
||||
|
||||
virtual std::vector<float> get_sigmas(uint32_t n) = 0;
|
||||
|
||||
@@ -75,6 +76,144 @@ struct DiscreteSchedule : SigmaSchedule {
|
||||
}
|
||||
};
|
||||
|
||||
/*
|
||||
https://research.nvidia.com/labs/toronto-ai/AlignYourSteps/howto.html
|
||||
*/
|
||||
struct AYSSchedule : SigmaSchedule {
|
||||
/* interp and linear_interp adapted from dpilger26's NumCpp library:
|
||||
* https://github.com/dpilger26/NumCpp/tree/5e40aab74d14e257d65d3dc385c9ff9e2120c60e */
|
||||
constexpr double interp(double left, double right, double perc) noexcept {
|
||||
return (left * (1. - perc)) + (right * perc);
|
||||
}
|
||||
|
||||
/* This will make the assumption that the reference x and y values are
|
||||
* already sorted in ascending order because they are being generated as
|
||||
* such in the calling function */
|
||||
std::vector<double> linear_interp(std::vector<float> new_x,
|
||||
const std::vector<float> ref_x,
|
||||
const std::vector<float> ref_y) {
|
||||
const size_t len_x = new_x.size();
|
||||
size_t i = 0;
|
||||
size_t j = 0;
|
||||
std::vector<double> new_y(len_x);
|
||||
|
||||
if (ref_x.size() != ref_y.size()) {
|
||||
LOG_ERROR("Linear Interoplation Failed: length mismatch");
|
||||
return new_y;
|
||||
}
|
||||
|
||||
/* serves as the bounds checking for the below while loop */
|
||||
if ((new_x[0] < ref_x[0]) || (new_x[new_x.size() - 1] > ref_x[ref_x.size() - 1])) {
|
||||
LOG_ERROR("Linear Interpolation Failed: bad bounds");
|
||||
return new_y;
|
||||
}
|
||||
|
||||
while (i < len_x) {
|
||||
if ((ref_x[j] > new_x[i]) || (new_x[i] > ref_x[j + 1])) {
|
||||
j++;
|
||||
continue;
|
||||
}
|
||||
|
||||
const double perc = static_cast<double>(new_x[i] - ref_x[j]) / static_cast<double>(ref_x[j + 1] - ref_x[j]);
|
||||
|
||||
new_y[i] = interp(ref_y[j], ref_y[j + 1], perc);
|
||||
i++;
|
||||
}
|
||||
|
||||
return new_y;
|
||||
}
|
||||
|
||||
std::vector<float> linear_space(const float start, const float end, const size_t num_points) {
|
||||
std::vector<float> result(num_points);
|
||||
const float inc = (end - start) / (static_cast<float>(num_points - 1));
|
||||
|
||||
if (num_points > 0) {
|
||||
result[0] = start;
|
||||
|
||||
for (size_t i = 1; i < num_points; i++) {
|
||||
result[i] = result[i - 1] + inc;
|
||||
}
|
||||
}
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
std::vector<float> log_linear_interpolation(std::vector<float> sigma_in,
|
||||
const size_t new_len) {
|
||||
const size_t s_len = sigma_in.size();
|
||||
std::vector<float> x_vals = linear_space(0.f, 1.f, s_len);
|
||||
std::vector<float> y_vals(s_len);
|
||||
|
||||
/* Reverses the input array to be ascending instead of descending,
|
||||
* also hits it with a log, it is log-linear interpolation after all */
|
||||
for (size_t i = 0; i < s_len; i++) {
|
||||
y_vals[i] = std::log(sigma_in[s_len - i - 1]);
|
||||
}
|
||||
|
||||
std::vector<float> new_x_vals = linear_space(0.f, 1.f, new_len);
|
||||
std::vector<double> new_y_vals = linear_interp(new_x_vals, x_vals, y_vals);
|
||||
std::vector<float> results(new_len);
|
||||
|
||||
for (size_t i = 0; i < new_len; i++) {
|
||||
results[i] = static_cast<float>(std::exp(new_y_vals[new_len - i - 1]));
|
||||
}
|
||||
|
||||
return results;
|
||||
}
|
||||
|
||||
std::vector<float> get_sigmas(uint32_t len) {
|
||||
const std::vector<float> noise_levels[] = {
|
||||
/* SD1.5 */
|
||||
{14.6146412293f, 6.4745760956f, 3.8636745985f, 2.6946151520f,
|
||||
1.8841921177f, 1.3943805092f, 0.9642583904f, 0.6523686016f,
|
||||
0.3977456272f, 0.1515232662f, 0.0291671582f},
|
||||
/* SDXL */
|
||||
{14.6146412293f, 6.3184485287f, 3.7681790315f, 2.1811480769f,
|
||||
1.3405244945f, 0.8620721141f, 0.5550693289f, 0.3798540708f,
|
||||
0.2332364134f, 0.1114188177f, 0.0291671582f},
|
||||
/* SVD */
|
||||
{700.00f, 54.5f, 15.886f, 7.977f, 4.248f, 1.789f, 0.981f, 0.403f,
|
||||
0.173f, 0.034f, 0.002f},
|
||||
};
|
||||
|
||||
std::vector<float> inputs;
|
||||
std::vector<float> results(len + 1);
|
||||
|
||||
switch (version) {
|
||||
case VERSION_2_x: /* fallthrough */
|
||||
LOG_WARN("AYS not designed for SD2.X models");
|
||||
case VERSION_1_x:
|
||||
LOG_INFO("AYS using SD1.5 noise levels");
|
||||
inputs = noise_levels[0];
|
||||
break;
|
||||
case VERSION_XL:
|
||||
LOG_INFO("AYS using SDXL noise levels");
|
||||
inputs = noise_levels[1];
|
||||
break;
|
||||
case VERSION_SVD:
|
||||
LOG_INFO("AYS using SVD noise levels");
|
||||
inputs = noise_levels[2];
|
||||
break;
|
||||
default:
|
||||
LOG_ERROR("Version not compatable with AYS scheduler");
|
||||
return results;
|
||||
}
|
||||
|
||||
/* Stretches those pre-calculated reference levels out to the desired
|
||||
* size using log-linear interpolation */
|
||||
if ((len + 1) != inputs.size()) {
|
||||
results = log_linear_interpolation(inputs, len + 1);
|
||||
} else {
|
||||
results = inputs;
|
||||
}
|
||||
|
||||
/* Not sure if this is strictly neccessary */
|
||||
results[len] = 0.0f;
|
||||
|
||||
return results;
|
||||
}
|
||||
};
|
||||
|
||||
struct KarrasSchedule : SigmaSchedule {
|
||||
std::vector<float> get_sigmas(uint32_t n) {
|
||||
// These *COULD* be function arguments here,
|
||||
@@ -122,4 +261,401 @@ struct CompVisVDenoiser : public Denoiser {
|
||||
}
|
||||
};
|
||||
|
||||
#endif // __DENOISER_HPP__
|
||||
typedef std::function<ggml_tensor*(ggml_tensor*, float, int)> denoise_cb_t;
|
||||
|
||||
// k diffusion reverse ODE: dx = (x - D(x;\sigma)) / \sigma dt; \sigma(t) = t
|
||||
static void sample_k_diffusion(sample_method_t method,
|
||||
denoise_cb_t model,
|
||||
ggml_context* work_ctx,
|
||||
ggml_tensor* x,
|
||||
std::vector<float> sigmas,
|
||||
std::shared_ptr<RNG> rng) {
|
||||
size_t steps = sigmas.size() - 1;
|
||||
// sample_euler_ancestral
|
||||
switch (method) {
|
||||
case EULER_A: {
|
||||
struct ggml_tensor* noise = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
float sigma = sigmas[i];
|
||||
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigma, i + 1);
|
||||
|
||||
// d = (x - denoised) / sigma
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
for (int i = 0; i < ggml_nelements(d); i++) {
|
||||
vec_d[i] = (vec_x[i] - vec_denoised[i]) / sigma;
|
||||
}
|
||||
}
|
||||
|
||||
// get_ancestral_step
|
||||
float sigma_up = std::min(sigmas[i + 1],
|
||||
std::sqrt(sigmas[i + 1] * sigmas[i + 1] * (sigmas[i] * sigmas[i] - sigmas[i + 1] * sigmas[i + 1]) / (sigmas[i] * sigmas[i])));
|
||||
float sigma_down = std::sqrt(sigmas[i + 1] * sigmas[i + 1] - sigma_up * sigma_up);
|
||||
|
||||
// Euler method
|
||||
float dt = sigma_down - sigmas[i];
|
||||
// x = x + d * dt
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
|
||||
for (int i = 0; i < ggml_nelements(x); i++) {
|
||||
vec_x[i] = vec_x[i] + vec_d[i] * dt;
|
||||
}
|
||||
}
|
||||
|
||||
if (sigmas[i + 1] > 0) {
|
||||
// x = x + noise_sampler(sigmas[i], sigmas[i + 1]) * s_noise * sigma_up
|
||||
ggml_tensor_set_f32_randn(noise, rng);
|
||||
// noise = load_tensor_from_file(work_ctx, "./rand" + std::to_string(i+1) + ".bin");
|
||||
{
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_noise = (float*)noise->data;
|
||||
|
||||
for (int i = 0; i < ggml_nelements(x); i++) {
|
||||
vec_x[i] = vec_x[i] + vec_noise[i] * sigma_up;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case EULER: // Implemented without any sigma churn
|
||||
{
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
float sigma = sigmas[i];
|
||||
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigma, i + 1);
|
||||
|
||||
// d = (x - denoised) / sigma
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(d); j++) {
|
||||
vec_d[j] = (vec_x[j] - vec_denoised[j]) / sigma;
|
||||
}
|
||||
}
|
||||
|
||||
float dt = sigmas[i + 1] - sigma;
|
||||
// x = x + d * dt
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case HEUN: {
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* x2 = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigmas[i], -(i + 1));
|
||||
|
||||
// d = (x - denoised) / sigma
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_d[j] = (vec_x[j] - vec_denoised[j]) / sigmas[i];
|
||||
}
|
||||
}
|
||||
|
||||
float dt = sigmas[i + 1] - sigmas[i];
|
||||
if (sigmas[i + 1] == 0) {
|
||||
// Euler step
|
||||
// x = x + d * dt
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
} else {
|
||||
// Heun step
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_d2 = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_x2 = (float*)x2->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x2[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
|
||||
ggml_tensor* denoised = model(x2, sigmas[i + 1], i + 1);
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
float d2 = (vec_x2[j] - vec_denoised[j]) / sigmas[i + 1];
|
||||
vec_d[j] = (vec_d[j] + d2) / 2;
|
||||
vec_x[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case DPM2: {
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* x2 = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigmas[i], i + 1);
|
||||
|
||||
// d = (x - denoised) / sigma
|
||||
{
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_d[j] = (vec_x[j] - vec_denoised[j]) / sigmas[i];
|
||||
}
|
||||
}
|
||||
|
||||
if (sigmas[i + 1] == 0) {
|
||||
// Euler step
|
||||
// x = x + d * dt
|
||||
float dt = sigmas[i + 1] - sigmas[i];
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
} else {
|
||||
// DPM-Solver-2
|
||||
float sigma_mid = exp(0.5f * (log(sigmas[i]) + log(sigmas[i + 1])));
|
||||
float dt_1 = sigma_mid - sigmas[i];
|
||||
float dt_2 = sigmas[i + 1] - sigmas[i];
|
||||
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_x2 = (float*)x2->data;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x2[j] = vec_x[j] + vec_d[j] * dt_1;
|
||||
}
|
||||
|
||||
ggml_tensor* denoised = model(x2, sigma_mid, i + 1);
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
float d2 = (vec_x2[j] - vec_denoised[j]) / sigma_mid;
|
||||
vec_x[j] = vec_x[j] + d2 * dt_2;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
} break;
|
||||
case DPMPP2S_A: {
|
||||
struct ggml_tensor* noise = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* x2 = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigmas[i], i + 1);
|
||||
|
||||
// get_ancestral_step
|
||||
float sigma_up = std::min(sigmas[i + 1],
|
||||
std::sqrt(sigmas[i + 1] * sigmas[i + 1] * (sigmas[i] * sigmas[i] - sigmas[i + 1] * sigmas[i + 1]) / (sigmas[i] * sigmas[i])));
|
||||
float sigma_down = std::sqrt(sigmas[i + 1] * sigmas[i + 1] - sigma_up * sigma_up);
|
||||
auto t_fn = [](float sigma) -> float { return -log(sigma); };
|
||||
auto sigma_fn = [](float t) -> float { return exp(-t); };
|
||||
|
||||
if (sigma_down == 0) {
|
||||
// Euler step
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(d); j++) {
|
||||
vec_d[j] = (vec_x[j] - vec_denoised[j]) / sigmas[i];
|
||||
}
|
||||
|
||||
// TODO: If sigma_down == 0, isn't this wrong?
|
||||
// But
|
||||
// https://github.com/crowsonkb/k-diffusion/blob/master/k_diffusion/sampling.py#L525
|
||||
// has this exactly the same way.
|
||||
float dt = sigma_down - sigmas[i];
|
||||
for (int j = 0; j < ggml_nelements(d); j++) {
|
||||
vec_x[j] = vec_x[j] + vec_d[j] * dt;
|
||||
}
|
||||
} else {
|
||||
// DPM-Solver++(2S)
|
||||
float t = t_fn(sigmas[i]);
|
||||
float t_next = t_fn(sigma_down);
|
||||
float h = t_next - t;
|
||||
float s = t + 0.5f * h;
|
||||
|
||||
float* vec_d = (float*)d->data;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_x2 = (float*)x2->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
|
||||
// First half-step
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x2[j] = (sigma_fn(s) / sigma_fn(t)) * vec_x[j] - (exp(-h * 0.5f) - 1) * vec_denoised[j];
|
||||
}
|
||||
|
||||
ggml_tensor* denoised = model(x2, sigmas[i + 1], i + 1);
|
||||
|
||||
// Second half-step
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = (sigma_fn(t_next) / sigma_fn(t)) * vec_x[j] - (exp(-h) - 1) * vec_denoised[j];
|
||||
}
|
||||
}
|
||||
|
||||
// Noise addition
|
||||
if (sigmas[i + 1] > 0) {
|
||||
ggml_tensor_set_f32_randn(noise, rng);
|
||||
{
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_noise = (float*)noise->data;
|
||||
|
||||
for (int i = 0; i < ggml_nelements(x); i++) {
|
||||
vec_x[i] = vec_x[i] + vec_noise[i] * sigma_up;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case DPMPP2M: // DPM++ (2M) from Karras et al (2022)
|
||||
{
|
||||
struct ggml_tensor* old_denoised = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
auto t_fn = [](float sigma) -> float { return -log(sigma); };
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigmas[i], i + 1);
|
||||
|
||||
float t = t_fn(sigmas[i]);
|
||||
float t_next = t_fn(sigmas[i + 1]);
|
||||
float h = t_next - t;
|
||||
float a = sigmas[i + 1] / sigmas[i];
|
||||
float b = exp(-h) - 1.f;
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
float* vec_old_denoised = (float*)old_denoised->data;
|
||||
|
||||
if (i == 0 || sigmas[i + 1] == 0) {
|
||||
// Simpler step for the edge cases
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = a * vec_x[j] - b * vec_denoised[j];
|
||||
}
|
||||
} else {
|
||||
float h_last = t - t_fn(sigmas[i - 1]);
|
||||
float r = h_last / h;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
float denoised_d = (1.f + 1.f / (2.f * r)) * vec_denoised[j] - (1.f / (2.f * r)) * vec_old_denoised[j];
|
||||
vec_x[j] = a * vec_x[j] - b * denoised_d;
|
||||
}
|
||||
}
|
||||
|
||||
// old_denoised = denoised
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_old_denoised[j] = vec_denoised[j];
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case DPMPP2Mv2: // Modified DPM++ (2M) from https://github.com/AUTOMATIC1111/stable-diffusion-webui/discussions/8457
|
||||
{
|
||||
struct ggml_tensor* old_denoised = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
auto t_fn = [](float sigma) -> float { return -log(sigma); };
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigmas[i], i + 1);
|
||||
|
||||
float t = t_fn(sigmas[i]);
|
||||
float t_next = t_fn(sigmas[i + 1]);
|
||||
float h = t_next - t;
|
||||
float a = sigmas[i + 1] / sigmas[i];
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
float* vec_old_denoised = (float*)old_denoised->data;
|
||||
|
||||
if (i == 0 || sigmas[i + 1] == 0) {
|
||||
// Simpler step for the edge cases
|
||||
float b = exp(-h) - 1.f;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = a * vec_x[j] - b * vec_denoised[j];
|
||||
}
|
||||
} else {
|
||||
float h_last = t - t_fn(sigmas[i - 1]);
|
||||
float h_min = std::min(h_last, h);
|
||||
float h_max = std::max(h_last, h);
|
||||
float r = h_max / h_min;
|
||||
float h_d = (h_max + h_min) / 2.f;
|
||||
float b = exp(-h_d) - 1.f;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
float denoised_d = (1.f + 1.f / (2.f * r)) * vec_denoised[j] - (1.f / (2.f * r)) * vec_old_denoised[j];
|
||||
vec_x[j] = a * vec_x[j] - b * denoised_d;
|
||||
}
|
||||
}
|
||||
|
||||
// old_denoised = denoised
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_old_denoised[j] = vec_denoised[j];
|
||||
}
|
||||
}
|
||||
} break;
|
||||
case LCM: // Latent Consistency Models
|
||||
{
|
||||
struct ggml_tensor* noise = ggml_dup_tensor(work_ctx, x);
|
||||
struct ggml_tensor* d = ggml_dup_tensor(work_ctx, x);
|
||||
|
||||
for (int i = 0; i < steps; i++) {
|
||||
float sigma = sigmas[i];
|
||||
|
||||
// denoise
|
||||
ggml_tensor* denoised = model(x, sigma, i + 1);
|
||||
|
||||
// x = denoised
|
||||
{
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_denoised = (float*)denoised->data;
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = vec_denoised[j];
|
||||
}
|
||||
}
|
||||
|
||||
if (sigmas[i + 1] > 0) {
|
||||
// x += sigmas[i + 1] * noise_sampler(sigmas[i], sigmas[i + 1])
|
||||
ggml_tensor_set_f32_randn(noise, rng);
|
||||
// noise = load_tensor_from_file(res_ctx, "./rand" + std::to_string(i+1) + ".bin");
|
||||
{
|
||||
float* vec_x = (float*)x->data;
|
||||
float* vec_noise = (float*)noise->data;
|
||||
|
||||
for (int j = 0; j < ggml_nelements(x); j++) {
|
||||
vec_x[j] = vec_x[j] + sigmas[i + 1] * vec_noise[j];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
} break;
|
||||
|
||||
default:
|
||||
LOG_ERROR("Attempting to sample with nonexisting sample method %i", method);
|
||||
abort();
|
||||
}
|
||||
}
|
||||
|
||||
#endif // __DENOISER_HPP__
|
||||
|
||||
16
esrgan.hpp
@@ -42,13 +42,13 @@ public:
|
||||
auto conv5 = std::dynamic_pointer_cast<Conv2d>(blocks["conv5"]);
|
||||
|
||||
auto x1 = lrelu(ctx, conv1->forward(ctx, x));
|
||||
auto x_cat = ggml_concat(ctx, x, x1);
|
||||
auto x_cat = ggml_concat(ctx, x, x1, 2);
|
||||
auto x2 = lrelu(ctx, conv2->forward(ctx, x_cat));
|
||||
x_cat = ggml_concat(ctx, x_cat, x2);
|
||||
x_cat = ggml_concat(ctx, x_cat, x2, 2);
|
||||
auto x3 = lrelu(ctx, conv3->forward(ctx, x_cat));
|
||||
x_cat = ggml_concat(ctx, x_cat, x3);
|
||||
x_cat = ggml_concat(ctx, x_cat, x3, 2);
|
||||
auto x4 = lrelu(ctx, conv4->forward(ctx, x_cat));
|
||||
x_cat = ggml_concat(ctx, x_cat, x4);
|
||||
x_cat = ggml_concat(ctx, x_cat, x4, 2);
|
||||
auto x5 = conv5->forward(ctx, x_cat);
|
||||
|
||||
x5 = ggml_add(ctx, ggml_scale(ctx, x5, 0.2f), x);
|
||||
@@ -152,14 +152,6 @@ struct ESRGAN : public GGMLModule {
|
||||
return "esrgan";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return rrdb_net.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return rrdb_net.get_params_num();
|
||||
}
|
||||
|
||||
bool load_from_file(const std::string& file_path) {
|
||||
LOG_INFO("loading esrgan from '%s'", file_path.c_str());
|
||||
|
||||
|
||||
@@ -10,12 +10,17 @@
|
||||
#include "stable-diffusion.h"
|
||||
|
||||
#define STB_IMAGE_IMPLEMENTATION
|
||||
#define STB_IMAGE_STATIC
|
||||
#include "stb_image.h"
|
||||
|
||||
#define STB_IMAGE_WRITE_IMPLEMENTATION
|
||||
#define STB_IMAGE_WRITE_STATIC
|
||||
#include "stb_image_write.h"
|
||||
|
||||
#define STB_IMAGE_RESIZE_IMPLEMENTATION
|
||||
#define STB_IMAGE_RESIZE_STATIC
|
||||
#include "stb_image_resize.h"
|
||||
|
||||
const char* rng_type_to_str[] = {
|
||||
"std_default",
|
||||
"cuda",
|
||||
@@ -38,6 +43,7 @@ const char* schedule_str[] = {
|
||||
"default",
|
||||
"discrete",
|
||||
"karras",
|
||||
"ays",
|
||||
};
|
||||
|
||||
const char* modes_str[] = {
|
||||
@@ -65,6 +71,8 @@ struct SDParams {
|
||||
std::string esrgan_path;
|
||||
std::string controlnet_path;
|
||||
std::string embeddings_path;
|
||||
std::string stacked_id_embeddings_path;
|
||||
std::string input_id_images_path;
|
||||
sd_type_t wtype = SD_TYPE_COUNT;
|
||||
std::string lora_model_dir;
|
||||
std::string output_path = "output.png";
|
||||
@@ -73,12 +81,13 @@ struct SDParams {
|
||||
|
||||
std::string prompt;
|
||||
std::string negative_prompt;
|
||||
float min_cfg = 1.0f;
|
||||
float cfg_scale = 7.0f;
|
||||
int clip_skip = -1; // <= 0 represents unspecified
|
||||
int width = 512;
|
||||
int height = 512;
|
||||
int batch_count = 1;
|
||||
float min_cfg = 1.0f;
|
||||
float cfg_scale = 7.0f;
|
||||
float style_ratio = 20.f;
|
||||
int clip_skip = -1; // <= 0 represents unspecified
|
||||
int width = 512;
|
||||
int height = 512;
|
||||
int batch_count = 1;
|
||||
|
||||
int video_frames = 6;
|
||||
int motion_bucket_id = 127;
|
||||
@@ -95,7 +104,11 @@ struct SDParams {
|
||||
bool verbose = false;
|
||||
bool vae_tiling = false;
|
||||
bool control_net_cpu = false;
|
||||
bool normalize_input = false;
|
||||
bool clip_on_cpu = false;
|
||||
bool vae_on_cpu = false;
|
||||
bool canny_preprocess = false;
|
||||
bool color = false;
|
||||
int upscale_repeats = 1;
|
||||
};
|
||||
|
||||
@@ -110,10 +123,16 @@ void print_params(SDParams params) {
|
||||
printf(" esrgan_path: %s\n", params.esrgan_path.c_str());
|
||||
printf(" controlnet_path: %s\n", params.controlnet_path.c_str());
|
||||
printf(" embeddings_path: %s\n", params.embeddings_path.c_str());
|
||||
printf(" stacked_id_embeddings_path: %s\n", params.stacked_id_embeddings_path.c_str());
|
||||
printf(" input_id_images_path: %s\n", params.input_id_images_path.c_str());
|
||||
printf(" style ratio: %.2f\n", params.style_ratio);
|
||||
printf(" normzalize input image : %s\n", params.normalize_input ? "true" : "false");
|
||||
printf(" output_path: %s\n", params.output_path.c_str());
|
||||
printf(" init_img: %s\n", params.input_path.c_str());
|
||||
printf(" control_image: %s\n", params.control_image_path.c_str());
|
||||
printf(" clip on cpu: %s\n", params.clip_on_cpu ? "true" : "false");
|
||||
printf(" controlnet cpu: %s\n", params.control_net_cpu ? "true" : "false");
|
||||
printf(" vae decoder on cpu:%s\n", params.vae_on_cpu ? "true" : "false");
|
||||
printf(" strength(control): %.2f\n", params.control_strength);
|
||||
printf(" prompt: %s\n", params.prompt.c_str());
|
||||
printf(" negative_prompt: %s\n", params.negative_prompt.c_str());
|
||||
@@ -146,6 +165,9 @@ void print_usage(int argc, const char* argv[]) {
|
||||
printf(" --taesd [TAESD_PATH] path to taesd. Using Tiny AutoEncoder for fast decoding (low quality)\n");
|
||||
printf(" --control-net [CONTROL_PATH] path to control net model\n");
|
||||
printf(" --embd-dir [EMBEDDING_PATH] path to embeddings.\n");
|
||||
printf(" --stacked-id-embd-dir [DIR] path to PHOTOMAKER stacked id embeddings.\n");
|
||||
printf(" --input-id-images-dir [DIR] path to PHOTOMAKER input id images dir.\n");
|
||||
printf(" --normalize-input normalize PHOTOMAKER input id images\n");
|
||||
printf(" --upscale-model [ESRGAN_PATH] path to esrgan model. Upscale images after generate, just RealESRGAN_x4plus_anime_6B supported by now.\n");
|
||||
printf(" --upscale-repeats Run the ESRGAN upscaler this many times (default 1)\n");
|
||||
printf(" --type [TYPE] weight type (f32, f16, q4_0, q4_1, q5_0, q5_1, q8_0)\n");
|
||||
@@ -158,6 +180,7 @@ void print_usage(int argc, const char* argv[]) {
|
||||
printf(" -n, --negative-prompt PROMPT the negative prompt (default: \"\")\n");
|
||||
printf(" --cfg-scale SCALE unconditional guidance scale: (default: 7.0)\n");
|
||||
printf(" --strength STRENGTH strength for noising/unnoising (default: 0.75)\n");
|
||||
printf(" --style-ratio STYLE-RATIO strength for keeping input identity (default: 20%%)\n");
|
||||
printf(" --control-strength STRENGTH strength to apply Control Net (default: 0.9)\n");
|
||||
printf(" 1.0 corresponds to full destruction of information in init image\n");
|
||||
printf(" -H, --height H image height, in pixel space (default: 512)\n");
|
||||
@@ -168,12 +191,13 @@ void print_usage(int argc, const char* argv[]) {
|
||||
printf(" --rng {std_default, cuda} RNG (default: cuda)\n");
|
||||
printf(" -s SEED, --seed SEED RNG seed (default: 42, use random seed for < 0)\n");
|
||||
printf(" -b, --batch-count COUNT number of images to generate.\n");
|
||||
printf(" --schedule {discrete, karras} Denoiser sigma schedule (default: discrete)\n");
|
||||
printf(" --schedule {discrete, karras, ays} Denoiser sigma schedule (default: discrete)\n");
|
||||
printf(" --clip-skip N ignore last layers of CLIP network; 1 ignores none, 2 ignores one layer (default: -1)\n");
|
||||
printf(" <= 0 represents unspecified, will be 1 for SD1.x, 2 for SD2.x\n");
|
||||
printf(" --vae-tiling process vae in tiles to reduce memory usage\n");
|
||||
printf(" --control-net-cpu keep controlnet in cpu (for low vram)\n");
|
||||
printf(" --canny apply canny preprocessor (edge detection)\n");
|
||||
printf(" --color Colors the logging tags according to level\n");
|
||||
printf(" -v, --verbose print extra info\n");
|
||||
}
|
||||
|
||||
@@ -244,6 +268,18 @@ void parse_args(int argc, const char** argv, SDParams& params) {
|
||||
break;
|
||||
}
|
||||
params.embeddings_path = argv[i];
|
||||
} else if (arg == "--stacked-id-embd-dir") {
|
||||
if (++i >= argc) {
|
||||
invalid_arg = true;
|
||||
break;
|
||||
}
|
||||
params.stacked_id_embeddings_path = argv[i];
|
||||
} else if (arg == "--input-id-images-dir") {
|
||||
if (++i >= argc) {
|
||||
invalid_arg = true;
|
||||
break;
|
||||
}
|
||||
params.input_id_images_path = argv[i];
|
||||
} else if (arg == "--type") {
|
||||
if (++i >= argc) {
|
||||
invalid_arg = true;
|
||||
@@ -327,6 +363,12 @@ void parse_args(int argc, const char** argv, SDParams& params) {
|
||||
break;
|
||||
}
|
||||
params.strength = std::stof(argv[i]);
|
||||
} else if (arg == "--style-ratio") {
|
||||
if (++i >= argc) {
|
||||
invalid_arg = true;
|
||||
break;
|
||||
}
|
||||
params.style_ratio = std::stof(argv[i]);
|
||||
} else if (arg == "--control-strength") {
|
||||
if (++i >= argc) {
|
||||
invalid_arg = true;
|
||||
@@ -361,6 +403,12 @@ void parse_args(int argc, const char** argv, SDParams& params) {
|
||||
params.vae_tiling = true;
|
||||
} else if (arg == "--control-net-cpu") {
|
||||
params.control_net_cpu = true;
|
||||
} else if (arg == "--normalize-input") {
|
||||
params.normalize_input = true;
|
||||
} else if (arg == "--clip-on-cpu") {
|
||||
params.clip_on_cpu = true; // will slow down get_learned_condiotion but necessary for low MEM GPUs
|
||||
} else if (arg == "--vae-on-cpu") {
|
||||
params.vae_on_cpu = true; // will slow down latent decoding but necessary for low MEM GPUs
|
||||
} else if (arg == "--canny") {
|
||||
params.canny_preprocess = true;
|
||||
} else if (arg == "-b" || arg == "--batch-count") {
|
||||
@@ -428,6 +476,8 @@ void parse_args(int argc, const char** argv, SDParams& params) {
|
||||
exit(0);
|
||||
} else if (arg == "-v" || arg == "--verbose") {
|
||||
params.verbose = true;
|
||||
} else if (arg == "--color") {
|
||||
params.color = true;
|
||||
} else {
|
||||
fprintf(stderr, "error: unknown argument: %s\n", arg.c_str());
|
||||
print_usage(argc, argv);
|
||||
@@ -531,18 +581,47 @@ std::string get_image_params(SDParams params, int64_t seed) {
|
||||
return parameter_string;
|
||||
}
|
||||
|
||||
/* Enables Printing the log level tag in color using ANSI escape codes */
|
||||
void sd_log_cb(enum sd_log_level_t level, const char* log, void* data) {
|
||||
SDParams* params = (SDParams*)data;
|
||||
if (!params->verbose && level <= SD_LOG_DEBUG) {
|
||||
int tag_color;
|
||||
const char* level_str;
|
||||
FILE* out_stream = (level == SD_LOG_ERROR) ? stderr : stdout;
|
||||
|
||||
if (!log || (!params->verbose && level <= SD_LOG_DEBUG)) {
|
||||
return;
|
||||
}
|
||||
if (level <= SD_LOG_INFO) {
|
||||
fputs(log, stdout);
|
||||
fflush(stdout);
|
||||
} else {
|
||||
fputs(log, stderr);
|
||||
fflush(stderr);
|
||||
|
||||
switch (level) {
|
||||
case SD_LOG_DEBUG:
|
||||
tag_color = 37;
|
||||
level_str = "DEBUG";
|
||||
break;
|
||||
case SD_LOG_INFO:
|
||||
tag_color = 34;
|
||||
level_str = "INFO";
|
||||
break;
|
||||
case SD_LOG_WARN:
|
||||
tag_color = 35;
|
||||
level_str = "WARN";
|
||||
break;
|
||||
case SD_LOG_ERROR:
|
||||
tag_color = 31;
|
||||
level_str = "ERROR";
|
||||
break;
|
||||
default: /* Potential future-proofing */
|
||||
tag_color = 33;
|
||||
level_str = "?????";
|
||||
break;
|
||||
}
|
||||
|
||||
if (params->color == true) {
|
||||
fprintf(out_stream, "\033[%d;1m[%-5s]\033[0m ", tag_color, level_str);
|
||||
} else {
|
||||
fprintf(out_stream, "[%-5s] ", level_str);
|
||||
}
|
||||
fputs(log, out_stream);
|
||||
fflush(out_stream);
|
||||
}
|
||||
|
||||
int main(int argc, const char* argv[]) {
|
||||
@@ -579,32 +658,59 @@ int main(int argc, const char* argv[]) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
bool vae_decode_only = true;
|
||||
uint8_t* input_image_buffer = NULL;
|
||||
bool vae_decode_only = true;
|
||||
uint8_t* input_image_buffer = NULL;
|
||||
uint8_t* control_image_buffer = NULL;
|
||||
if (params.mode == IMG2IMG || params.mode == IMG2VID) {
|
||||
vae_decode_only = false;
|
||||
|
||||
int c = 0;
|
||||
input_image_buffer = stbi_load(params.input_path.c_str(), ¶ms.width, ¶ms.height, &c, 3);
|
||||
int width = 0;
|
||||
int height = 0;
|
||||
input_image_buffer = stbi_load(params.input_path.c_str(), &width, &height, &c, 3);
|
||||
if (input_image_buffer == NULL) {
|
||||
fprintf(stderr, "load image from '%s' failed\n", params.input_path.c_str());
|
||||
return 1;
|
||||
}
|
||||
if (c != 3) {
|
||||
fprintf(stderr, "input image must be a 3 channels RGB image, but got %d channels\n", c);
|
||||
if (c < 3) {
|
||||
fprintf(stderr, "the number of channels for the input image must be >= 3, but got %d channels\n", c);
|
||||
free(input_image_buffer);
|
||||
return 1;
|
||||
}
|
||||
if (params.width <= 0 || params.width % 64 != 0) {
|
||||
fprintf(stderr, "error: the width of image must be a multiple of 64\n");
|
||||
if (width <= 0) {
|
||||
fprintf(stderr, "error: the width of image must be greater than 0\n");
|
||||
free(input_image_buffer);
|
||||
return 1;
|
||||
}
|
||||
if (params.height <= 0 || params.height % 64 != 0) {
|
||||
fprintf(stderr, "error: the height of image must be a multiple of 64\n");
|
||||
if (height <= 0) {
|
||||
fprintf(stderr, "error: the height of image must be greater than 0\n");
|
||||
free(input_image_buffer);
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Resize input image ...
|
||||
if (params.height != height || params.width != width) {
|
||||
printf("resize input image from %dx%d to %dx%d\n", width, height, params.width, params.height);
|
||||
int resized_height = params.height;
|
||||
int resized_width = params.width;
|
||||
|
||||
uint8_t* resized_image_buffer = (uint8_t*)malloc(resized_height * resized_width * 3);
|
||||
if (resized_image_buffer == NULL) {
|
||||
fprintf(stderr, "error: allocate memory for resize input image\n");
|
||||
free(input_image_buffer);
|
||||
return 1;
|
||||
}
|
||||
stbir_resize(input_image_buffer, width, height, 0,
|
||||
resized_image_buffer, resized_width, resized_height, 0, STBIR_TYPE_UINT8,
|
||||
3 /*RGB channel*/, STBIR_ALPHA_CHANNEL_NONE, 0,
|
||||
STBIR_EDGE_CLAMP, STBIR_EDGE_CLAMP,
|
||||
STBIR_FILTER_BOX, STBIR_FILTER_BOX,
|
||||
STBIR_COLORSPACE_SRGB, nullptr);
|
||||
|
||||
// Save resized result
|
||||
free(input_image_buffer);
|
||||
input_image_buffer = resized_image_buffer;
|
||||
}
|
||||
}
|
||||
|
||||
sd_ctx_t* sd_ctx = new_sd_ctx(params.model_path.c_str(),
|
||||
@@ -613,6 +719,7 @@ int main(int argc, const char* argv[]) {
|
||||
params.controlnet_path.c_str(),
|
||||
params.lora_model_dir.c_str(),
|
||||
params.embeddings_path.c_str(),
|
||||
params.stacked_id_embeddings_path.c_str(),
|
||||
vae_decode_only,
|
||||
params.vae_tiling,
|
||||
true,
|
||||
@@ -620,38 +727,41 @@ int main(int argc, const char* argv[]) {
|
||||
params.wtype,
|
||||
params.rng_type,
|
||||
params.schedule,
|
||||
params.control_net_cpu);
|
||||
params.clip_on_cpu,
|
||||
params.control_net_cpu,
|
||||
params.vae_on_cpu);
|
||||
|
||||
if (sd_ctx == NULL) {
|
||||
printf("new_sd_ctx_t failed\n");
|
||||
return 1;
|
||||
}
|
||||
|
||||
sd_image_t* control_image = NULL;
|
||||
if (params.controlnet_path.size() > 0 && params.control_image_path.size() > 0) {
|
||||
int c = 0;
|
||||
control_image_buffer = stbi_load(params.control_image_path.c_str(), ¶ms.width, ¶ms.height, &c, 3);
|
||||
if (control_image_buffer == NULL) {
|
||||
fprintf(stderr, "load image from '%s' failed\n", params.control_image_path.c_str());
|
||||
return 1;
|
||||
}
|
||||
control_image = new sd_image_t{(uint32_t)params.width,
|
||||
(uint32_t)params.height,
|
||||
3,
|
||||
control_image_buffer};
|
||||
if (params.canny_preprocess) { // apply preprocessor
|
||||
control_image->data = preprocess_canny(control_image->data,
|
||||
control_image->width,
|
||||
control_image->height,
|
||||
0.08f,
|
||||
0.08f,
|
||||
0.8f,
|
||||
1.0f,
|
||||
false);
|
||||
}
|
||||
}
|
||||
|
||||
sd_image_t* results;
|
||||
if (params.mode == TXT2IMG) {
|
||||
sd_image_t* control_image = NULL;
|
||||
if (params.controlnet_path.size() > 0 && params.control_image_path.size() > 0) {
|
||||
int c = 0;
|
||||
input_image_buffer = stbi_load(params.control_image_path.c_str(), ¶ms.width, ¶ms.height, &c, 3);
|
||||
if (input_image_buffer == NULL) {
|
||||
fprintf(stderr, "load image from '%s' failed\n", params.control_image_path.c_str());
|
||||
return 1;
|
||||
}
|
||||
control_image = new sd_image_t{(uint32_t)params.width,
|
||||
(uint32_t)params.height,
|
||||
3,
|
||||
input_image_buffer};
|
||||
if (params.canny_preprocess) { // apply preprocessor
|
||||
control_image->data = preprocess_canny(control_image->data,
|
||||
control_image->width,
|
||||
control_image->height,
|
||||
0.08f,
|
||||
0.08f,
|
||||
0.8f,
|
||||
1.0f,
|
||||
false);
|
||||
}
|
||||
}
|
||||
results = txt2img(sd_ctx,
|
||||
params.prompt.c_str(),
|
||||
params.negative_prompt.c_str(),
|
||||
@@ -664,7 +774,10 @@ int main(int argc, const char* argv[]) {
|
||||
params.seed,
|
||||
params.batch_count,
|
||||
control_image,
|
||||
params.control_strength);
|
||||
params.control_strength,
|
||||
params.style_ratio,
|
||||
params.normalize_input,
|
||||
params.input_id_images_path.c_str());
|
||||
} else {
|
||||
sd_image_t input_image = {(uint32_t)params.width,
|
||||
(uint32_t)params.height,
|
||||
@@ -720,7 +833,12 @@ int main(int argc, const char* argv[]) {
|
||||
params.sample_steps,
|
||||
params.strength,
|
||||
params.seed,
|
||||
params.batch_count);
|
||||
params.batch_count,
|
||||
control_image,
|
||||
params.control_strength,
|
||||
params.style_ratio,
|
||||
params.normalize_input,
|
||||
params.input_id_images_path.c_str());
|
||||
}
|
||||
}
|
||||
|
||||
@@ -773,6 +891,8 @@ int main(int argc, const char* argv[]) {
|
||||
}
|
||||
free(results);
|
||||
free_sd_ctx(sd_ctx);
|
||||
free(control_image_buffer);
|
||||
free(input_image_buffer);
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
2
ggml
152
ggml_extend.hpp
@@ -80,8 +80,27 @@ __STATIC_INLINE__ ggml_fp16_t ggml_tensor_get_f16(const ggml_tensor* tensor, int
|
||||
return *(ggml_fp16_t*)((char*)(tensor->data) + i * tensor->nb[3] + j * tensor->nb[2] + k * tensor->nb[1] + l * tensor->nb[0]);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void print_ggml_tensor(struct ggml_tensor* tensor, bool shape_only = false) {
|
||||
printf("shape(%zu, %zu, %zu, %zu)\n", tensor->ne[0], tensor->ne[1], tensor->ne[2], tensor->ne[3]);
|
||||
static struct ggml_tensor* get_tensor_from_graph(struct ggml_cgraph* gf, const char* name) {
|
||||
struct ggml_tensor* res = NULL;
|
||||
for (int i = 0; i < gf->n_nodes; i++) {
|
||||
// printf("%d, %s \n", i, gf->nodes[i]->name);
|
||||
if (strcmp(ggml_get_name(gf->nodes[i]), name) == 0) {
|
||||
res = gf->nodes[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
for (int i = 0; i < gf->n_leafs; i++) {
|
||||
// printf("%d, %s \n", i, gf->leafs[i]->name);
|
||||
if (strcmp(ggml_get_name(gf->leafs[i]), name) == 0) {
|
||||
res = gf->leafs[i];
|
||||
break;
|
||||
}
|
||||
}
|
||||
return res;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void print_ggml_tensor(struct ggml_tensor* tensor, bool shape_only = false, const char* mark = "") {
|
||||
printf("%s (%s): shape(%zu, %zu, %zu, %zu)\n", mark, ggml_type_name(tensor->type), tensor->ne[0], tensor->ne[1], tensor->ne[2], tensor->ne[3]);
|
||||
fflush(stdout);
|
||||
if (shape_only) {
|
||||
return;
|
||||
@@ -217,6 +236,23 @@ __STATIC_INLINE__ uint8_t* sd_tensor_to_image(struct ggml_tensor* input) {
|
||||
return image_data;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ uint8_t* sd_tensor_to_mul_image(struct ggml_tensor* input, int idx) {
|
||||
int64_t width = input->ne[0];
|
||||
int64_t height = input->ne[1];
|
||||
int64_t channels = input->ne[2];
|
||||
GGML_ASSERT(channels == 3 && input->type == GGML_TYPE_F32);
|
||||
uint8_t* image_data = (uint8_t*)malloc(width * height * channels);
|
||||
for (int iy = 0; iy < height; iy++) {
|
||||
for (int ix = 0; ix < width; ix++) {
|
||||
for (int k = 0; k < channels; k++) {
|
||||
float value = ggml_tensor_get_f32(input, ix, iy, k, idx);
|
||||
*(image_data + iy * width * channels + ix * channels + k) = (uint8_t)(value * 255.0f);
|
||||
}
|
||||
}
|
||||
}
|
||||
return image_data;
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void sd_image_to_tensor(const uint8_t* image_data,
|
||||
struct ggml_tensor* output,
|
||||
bool scale = true) {
|
||||
@@ -237,6 +273,28 @@ __STATIC_INLINE__ void sd_image_to_tensor(const uint8_t* image_data,
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void sd_mul_images_to_tensor(const uint8_t* image_data,
|
||||
struct ggml_tensor* output,
|
||||
int idx,
|
||||
float* mean = NULL,
|
||||
float* std = NULL) {
|
||||
int64_t width = output->ne[0];
|
||||
int64_t height = output->ne[1];
|
||||
int64_t channels = output->ne[2];
|
||||
GGML_ASSERT(channels == 3 && output->type == GGML_TYPE_F32);
|
||||
for (int iy = 0; iy < height; iy++) {
|
||||
for (int ix = 0; ix < width; ix++) {
|
||||
for (int k = 0; k < channels; k++) {
|
||||
int value = *(image_data + iy * width * channels + ix * channels + k);
|
||||
float pixel_val = value / 255.0f;
|
||||
if (mean != NULL && std != NULL)
|
||||
pixel_val = (pixel_val - mean[k]) / std[k];
|
||||
ggml_tensor_set_f32(output, pixel_val, ix, iy, k, idx);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ void sd_image_f32_to_tensor(const float* image_data,
|
||||
struct ggml_tensor* output,
|
||||
bool scale = true) {
|
||||
@@ -247,7 +305,7 @@ __STATIC_INLINE__ void sd_image_f32_to_tensor(const float* image_data,
|
||||
for (int iy = 0; iy < height; iy++) {
|
||||
for (int ix = 0; ix < width; ix++) {
|
||||
for (int k = 0; k < channels; k++) {
|
||||
float value = *(image_data + iy * width * channels + ix * channels + k);
|
||||
int value = *(image_data + iy * width * channels + ix * channels + k);
|
||||
if (scale) {
|
||||
value /= 255.f;
|
||||
}
|
||||
@@ -396,7 +454,7 @@ __STATIC_INLINE__ void sd_tiling(ggml_tensor* input, ggml_tensor* output, const
|
||||
ggml_tensor* input_tile = ggml_new_tensor_4d(tiles_ctx, GGML_TYPE_F32, tile_size, tile_size, input->ne[2], 1);
|
||||
ggml_tensor* output_tile = ggml_new_tensor_4d(tiles_ctx, GGML_TYPE_F32, tile_size * scale, tile_size * scale, output->ne[2], 1);
|
||||
on_processing(input_tile, NULL, true);
|
||||
int num_tiles = (input_width * input_height) / (non_tile_overlap * non_tile_overlap);
|
||||
int num_tiles = ceil((float)input_width / non_tile_overlap) * ceil((float)input_height / non_tile_overlap);
|
||||
LOG_INFO("processing %i tiles", num_tiles);
|
||||
pretty_progress(1, num_tiles, 0.0f);
|
||||
int tile_count = 1;
|
||||
@@ -426,6 +484,7 @@ __STATIC_INLINE__ void sd_tiling(ggml_tensor* input, ggml_tensor* output, const
|
||||
if (tile_count < num_tiles) {
|
||||
pretty_progress(num_tiles, num_tiles, last_time);
|
||||
}
|
||||
ggml_free(tiles_ctx);
|
||||
}
|
||||
|
||||
__STATIC_INLINE__ struct ggml_tensor* ggml_group_norm_32(struct ggml_context* ctx,
|
||||
@@ -693,16 +752,17 @@ __STATIC_INLINE__ struct ggml_tensor* ggml_nn_timestep_embedding(
|
||||
return ggml_timestep_embedding(ctx, timesteps, dim, max_period);
|
||||
}
|
||||
|
||||
// struct GGMLComputeGraph {
|
||||
// virtual void init(struct ggml_context* ctx, ggml_type wtype) = 0;
|
||||
// virtual std::string get_desc() = 0;
|
||||
// virtual size_t get_params_mem_size() = 0;
|
||||
// virtual size_t get_params_num() = 0;
|
||||
// virtual struct ggml_cgraph* get_ggml_cgraph() = 0;
|
||||
// };
|
||||
__STATIC_INLINE__ size_t ggml_tensor_num(ggml_context* ctx) {
|
||||
size_t num = 0;
|
||||
for (ggml_tensor* t = ggml_get_first_tensor(ctx); t != nullptr; t = ggml_get_next_tensor(ctx, t)) {
|
||||
num++;
|
||||
}
|
||||
return num;
|
||||
}
|
||||
|
||||
#define MAX_PARAMS_TENSOR_NUM 10240
|
||||
#define MAX_GRAPH_SIZE 10240
|
||||
/* SDXL with LoRA requires more space */
|
||||
#define MAX_PARAMS_TENSOR_NUM 15360
|
||||
#define MAX_GRAPH_SIZE 15360
|
||||
|
||||
struct GGMLModule {
|
||||
protected:
|
||||
@@ -771,7 +831,10 @@ protected:
|
||||
|
||||
// compute the required memory
|
||||
size_t compute_buffer_size = ggml_gallocr_get_buffer_size(compute_allocr, 0);
|
||||
LOG_DEBUG("%s compute buffer size: %.2f MB", get_desc().c_str(), compute_buffer_size / 1024.0 / 1024.0);
|
||||
LOG_DEBUG("%s compute buffer size: %.2f MB(%s)",
|
||||
get_desc().c_str(),
|
||||
compute_buffer_size / 1024.0 / 1024.0,
|
||||
ggml_backend_is_cpu(backend) ? "RAM" : "VRAM");
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -787,9 +850,7 @@ protected:
|
||||
}
|
||||
|
||||
public:
|
||||
virtual size_t get_params_mem_size() = 0;
|
||||
virtual size_t get_params_num() = 0;
|
||||
virtual std::string get_desc() = 0;
|
||||
virtual std::string get_desc() = 0;
|
||||
|
||||
GGMLModule(ggml_backend_t backend, ggml_type wtype = GGML_TYPE_F32)
|
||||
: backend(backend), wtype(wtype) {
|
||||
@@ -809,15 +870,18 @@ public:
|
||||
}
|
||||
|
||||
bool alloc_params_buffer() {
|
||||
size_t num_tensors = get_params_num();
|
||||
size_t num_tensors = ggml_tensor_num(params_ctx);
|
||||
params_buffer = ggml_backend_alloc_ctx_tensors(params_ctx, backend);
|
||||
if (params_buffer == NULL) {
|
||||
LOG_ERROR("%s alloc params backend buffer failed", get_desc().c_str());
|
||||
return false;
|
||||
}
|
||||
size_t params_buffer_size = ggml_backend_buffer_get_size(params_buffer);
|
||||
LOG_DEBUG("%s params backend buffer size = % 6.2f MB (%i tensors)",
|
||||
get_desc().c_str(), params_buffer_size / (1024.0 * 1024.0), num_tensors);
|
||||
LOG_DEBUG("%s params backend buffer size = % 6.2f MB(%s) (%i tensors)",
|
||||
get_desc().c_str(),
|
||||
params_buffer_size / (1024.0 * 1024.0),
|
||||
ggml_backend_is_cpu(backend) ? "RAM" : "VRAM",
|
||||
num_tensors);
|
||||
return true;
|
||||
}
|
||||
|
||||
@@ -828,6 +892,13 @@ public:
|
||||
}
|
||||
}
|
||||
|
||||
size_t get_params_buffer_size() {
|
||||
if (params_buffer != NULL) {
|
||||
return ggml_backend_buffer_get_size(params_buffer);
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
void free_compute_buffer() {
|
||||
if (compute_allocr != NULL) {
|
||||
ggml_gallocr_free(compute_allocr);
|
||||
@@ -846,7 +917,7 @@ public:
|
||||
return NULL;
|
||||
}
|
||||
// it's performing a compute, check if backend isn't cpu
|
||||
if (!ggml_backend_is_cpu(backend) && tensor->backend == GGML_BACKEND_TYPE_CPU) {
|
||||
if (!ggml_backend_is_cpu(backend) && (tensor->buffer == NULL || ggml_backend_buffer_is_host(tensor->buffer))) {
|
||||
// pass input tensors to gpu memory
|
||||
auto backend_tensor = ggml_dup_tensor(compute_ctx, tensor);
|
||||
|
||||
@@ -865,11 +936,8 @@ public:
|
||||
alloc_compute_buffer(get_graph);
|
||||
reset_compute_ctx();
|
||||
struct ggml_cgraph* gf = get_graph();
|
||||
|
||||
GGML_ASSERT(ggml_gallocr_alloc_graph(compute_allocr, gf));
|
||||
|
||||
cpy_data_to_backend_tensor();
|
||||
|
||||
if (ggml_backend_is_cpu(backend)) {
|
||||
ggml_backend_cpu_set_n_threads(backend, n_threads);
|
||||
}
|
||||
@@ -879,13 +947,11 @@ public:
|
||||
ggml_backend_metal_set_n_cb(backend, n_threads);
|
||||
}
|
||||
#endif
|
||||
|
||||
ggml_backend_graph_compute(backend, gf);
|
||||
|
||||
#ifdef GGML_PERF
|
||||
ggml_graph_print(gf);
|
||||
#endif
|
||||
|
||||
if (output != NULL) {
|
||||
auto result = gf->nodes[gf->n_nodes - 1];
|
||||
if (*output == NULL && output_ctx != NULL) {
|
||||
@@ -903,19 +969,6 @@ public:
|
||||
};
|
||||
|
||||
class GGMLBlock {
|
||||
private:
|
||||
static char temp_buffer[1024 * 1024 * 10];
|
||||
ggml_context* get_temp_ctx() {
|
||||
struct ggml_init_params params;
|
||||
params.mem_size = sizeof(temp_buffer);
|
||||
params.mem_buffer = temp_buffer;
|
||||
params.no_alloc = true;
|
||||
|
||||
ggml_context* temp_ctx = ggml_init(params);
|
||||
GGML_ASSERT(temp_ctx != NULL);
|
||||
return temp_ctx;
|
||||
}
|
||||
|
||||
protected:
|
||||
typedef std::unordered_map<std::string, struct ggml_tensor*> ParameterMap;
|
||||
typedef std::unordered_map<std::string, std::shared_ptr<GGMLBlock>> GGMLBlockMap;
|
||||
@@ -938,14 +991,6 @@ public:
|
||||
init_params(ctx, wtype);
|
||||
}
|
||||
|
||||
std::tuple<size_t, size_t> get_params_info(ggml_type wtype) {
|
||||
ggml_context* temp_ctx = get_temp_ctx();
|
||||
init(temp_ctx, wtype);
|
||||
size_t num_tensors = get_params_num();
|
||||
size_t mem_size = get_params_mem_size();
|
||||
return {num_tensors, mem_size};
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
size_t num_tensors = params.size();
|
||||
for (auto& pair : blocks) {
|
||||
@@ -977,13 +1022,11 @@ public:
|
||||
}
|
||||
for (auto& pair : blocks) {
|
||||
auto& block = pair.second;
|
||||
|
||||
block->get_param_tensors(tensors, prefix + pair.first);
|
||||
}
|
||||
|
||||
for (auto& pair : params) {
|
||||
struct ggml_tensor* param = pair.second;
|
||||
|
||||
struct ggml_tensor* param = pair.second;
|
||||
tensors[prefix + pair.first] = pair.second;
|
||||
}
|
||||
}
|
||||
@@ -1243,13 +1286,12 @@ public:
|
||||
struct ggml_tensor* kqv = ggml_nn_attention(ctx, q, k, v, mask); // [N * n_head, n_token, d_head]
|
||||
|
||||
kqv = ggml_reshape_4d(ctx, kqv, d_head, n_token, n_head, N);
|
||||
kqv = ggml_cont(ctx, ggml_permute(ctx, kqv, 0, 2, 1, 3)); // [N, n_token, n_head, d_head]
|
||||
kqv = ggml_cont(ctx, ggml_permute(ctx, kqv, 0, 2, 1, 3)); // [N, n_token, n_head, d_head]
|
||||
x = ggml_reshape_3d(ctx, kqv, d_head * n_head, n_token, N); // [N, n_token, d_head * n_head]
|
||||
|
||||
x = ggml_reshape_3d(ctx, kqv, d_head * n_head, n_token, N); // [N * n_token, d_head * n_head]
|
||||
|
||||
x = out_proj->forward(ctx, x);
|
||||
x = out_proj->forward(ctx, x); // [N, n_token, embed_dim]
|
||||
return x;
|
||||
}
|
||||
};
|
||||
|
||||
#endif // __GGML_EXTEND__HPP__
|
||||
#endif // __GGML_EXTEND__HPP__
|
||||
|
||||
56
lora.hpp
@@ -11,12 +11,14 @@ struct LoraModel : public GGMLModule {
|
||||
std::string file_path;
|
||||
ModelLoader model_loader;
|
||||
bool load_failed = false;
|
||||
bool applied = false;
|
||||
|
||||
LoraModel(ggml_backend_t backend,
|
||||
ggml_type wtype,
|
||||
const std::string file_path = "")
|
||||
const std::string& file_path = "",
|
||||
const std::string& prefix = "")
|
||||
: file_path(file_path), GGMLModule(backend, wtype) {
|
||||
if (!model_loader.init_from_file(file_path)) {
|
||||
if (!model_loader.init_from_file(file_path, prefix)) {
|
||||
load_failed = true;
|
||||
}
|
||||
}
|
||||
@@ -25,16 +27,7 @@ struct LoraModel : public GGMLModule {
|
||||
return "lora";
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return LORA_GRAPH_SIZE;
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return model_loader.get_params_mem_size(NULL);
|
||||
}
|
||||
|
||||
|
||||
bool load_from_file() {
|
||||
bool load_from_file(bool filter_tensor = false) {
|
||||
LOG_INFO("loading LoRA from '%s'", file_path.c_str());
|
||||
|
||||
if (load_failed) {
|
||||
@@ -46,6 +39,11 @@ struct LoraModel : public GGMLModule {
|
||||
auto on_new_tensor_cb = [&](const TensorStorage& tensor_storage, ggml_tensor** dst_tensor) -> bool {
|
||||
const std::string& name = tensor_storage.name;
|
||||
|
||||
if (filter_tensor && !contains(name, "lora")) {
|
||||
// LOG_INFO("skipping LoRA tesnor '%s'", name.c_str());
|
||||
return true;
|
||||
}
|
||||
|
||||
if (dry_run) {
|
||||
struct ggml_tensor* real = ggml_new_tensor(params_ctx,
|
||||
tensor_storage.type,
|
||||
@@ -66,12 +64,11 @@ struct LoraModel : public GGMLModule {
|
||||
dry_run = false;
|
||||
model_loader.load_tensors(on_new_tensor_cb, backend);
|
||||
|
||||
|
||||
LOG_DEBUG("finished loaded lora");
|
||||
return true;
|
||||
}
|
||||
|
||||
struct ggml_cgraph* build_graph(std::map<std::string, struct ggml_tensor*> model_tensors) {
|
||||
struct ggml_cgraph* build_lora_graph(std::map<std::string, struct ggml_tensor*> model_tensors) {
|
||||
struct ggml_cgraph* gf = ggml_new_graph_custom(compute_ctx, LORA_GRAPH_SIZE, false);
|
||||
|
||||
std::set<std::string> applied_lora_tensors;
|
||||
@@ -85,7 +82,16 @@ struct LoraModel : public GGMLModule {
|
||||
}
|
||||
k_tensor = k_tensor.substr(0, k_pos);
|
||||
replace_all_chars(k_tensor, '.', '_');
|
||||
std::string lora_up_name = "lora." + k_tensor + ".lora_up.weight";
|
||||
// LOG_DEBUG("k_tensor %s", k_tensor.c_str());
|
||||
std::string lora_up_name = "lora." + k_tensor + ".lora_up.weight";
|
||||
if (lora_tensors.find(lora_up_name) == lora_tensors.end()) {
|
||||
if (k_tensor == "model_diffusion_model_output_blocks_2_2_conv") {
|
||||
// fix for some sdxl lora, like lcm-lora-xl
|
||||
k_tensor = "model_diffusion_model_output_blocks_2_1_conv";
|
||||
lora_up_name = "lora." + k_tensor + ".lora_up.weight";
|
||||
}
|
||||
}
|
||||
|
||||
std::string lora_down_name = "lora." + k_tensor + ".lora_down.weight";
|
||||
std::string alpha_name = "lora." + k_tensor + ".alpha";
|
||||
std::string scale_name = "lora." + k_tensor + ".scale";
|
||||
@@ -147,21 +153,37 @@ struct LoraModel : public GGMLModule {
|
||||
ggml_build_forward_expand(gf, final_weight);
|
||||
}
|
||||
|
||||
size_t total_lora_tensors_count = 0;
|
||||
size_t applied_lora_tensors_count = 0;
|
||||
|
||||
for (auto& kv : lora_tensors) {
|
||||
total_lora_tensors_count++;
|
||||
if (applied_lora_tensors.find(kv.first) == applied_lora_tensors.end()) {
|
||||
LOG_WARN("unused lora tensor %s", kv.first.c_str());
|
||||
} else {
|
||||
applied_lora_tensors_count++;
|
||||
}
|
||||
}
|
||||
/* Don't worry if this message shows up twice in the logs per LoRA,
|
||||
* this function is called once to calculate the required buffer size
|
||||
* and then again to actually generate a graph to be used */
|
||||
if (applied_lora_tensors_count != total_lora_tensors_count) {
|
||||
LOG_WARN("Only (%lu / %lu) LoRA tensors have been applied",
|
||||
applied_lora_tensors_count, total_lora_tensors_count);
|
||||
} else {
|
||||
LOG_DEBUG("(%lu / %lu) LoRA tensors applied successfully",
|
||||
applied_lora_tensors_count, total_lora_tensors_count);
|
||||
}
|
||||
|
||||
return gf;
|
||||
}
|
||||
|
||||
void apply(std::map<std::string, struct ggml_tensor*> model_tensors, int n_threads) {
|
||||
auto get_graph = [&]() -> struct ggml_cgraph* {
|
||||
return build_graph(model_tensors);
|
||||
return build_lora_graph(model_tensors);
|
||||
};
|
||||
GGMLModule::compute(get_graph, n_threads, true);
|
||||
}
|
||||
};
|
||||
|
||||
#endif // __LORA_HPP__
|
||||
#endif // __LORA_HPP__
|
||||
|
||||
116
model.cpp
@@ -108,14 +108,14 @@ std::unordered_map<std::string, std::string> open_clip_to_hf_clip_model = {
|
||||
{"model.positional_embedding", "transformer.text_model.embeddings.position_embedding.weight"},
|
||||
{"model.token_embedding.weight", "transformer.text_model.embeddings.token_embedding.weight"},
|
||||
{"model.text_projection", "transformer.text_model.text_projection"},
|
||||
{"model.visual.class_embedding", "transformer.visual_model.embeddings.class_embedding"},
|
||||
{"model.visual.conv1.weight", "transformer.visual_model.embeddings.patch_embedding.weight"},
|
||||
{"model.visual.ln_post.bias", "transformer.visual_model.post_layernorm.bias"},
|
||||
{"model.visual.ln_post.weight", "transformer.visual_model.post_layernorm.weight"},
|
||||
{"model.visual.ln_pre.bias", "transformer.visual_model.pre_layernorm.bias"},
|
||||
{"model.visual.ln_pre.weight", "transformer.visual_model.pre_layernorm.weight"},
|
||||
{"model.visual.positional_embedding", "transformer.visual_model.embeddings.position_embedding.weight"},
|
||||
{"model.visual.proj", "transformer.visual_model.visual_projection"},
|
||||
{"model.visual.class_embedding", "transformer.vision_model.embeddings.class_embedding"},
|
||||
{"model.visual.conv1.weight", "transformer.vision_model.embeddings.patch_embedding.weight"},
|
||||
{"model.visual.ln_post.bias", "transformer.vision_model.post_layernorm.bias"},
|
||||
{"model.visual.ln_post.weight", "transformer.vision_model.post_layernorm.weight"},
|
||||
{"model.visual.ln_pre.bias", "transformer.vision_model.pre_layernorm.bias"},
|
||||
{"model.visual.ln_pre.weight", "transformer.vision_model.pre_layernorm.weight"},
|
||||
{"model.visual.positional_embedding", "transformer.vision_model.embeddings.position_embedding.weight"},
|
||||
{"model.visual.proj", "transformer.visual_projection.weight"},
|
||||
};
|
||||
|
||||
std::unordered_map<std::string, std::string> open_clip_to_hk_clip_resblock = {
|
||||
@@ -157,6 +157,10 @@ std::string convert_open_clip_to_hf_clip(const std::string& name) {
|
||||
} else if (starts_with(new_name, "cond_stage_model.")) {
|
||||
prefix = "cond_stage_model.";
|
||||
new_name = new_name.substr(strlen("cond_stage_model."));
|
||||
} else if (ends_with(new_name, "vision_model.visual_projection.weight")) {
|
||||
prefix = new_name.substr(0, new_name.size() - strlen("vision_model.visual_projection.weight"));
|
||||
new_name = prefix + "visual_projection.weight";
|
||||
return new_name;
|
||||
} else {
|
||||
return new_name;
|
||||
}
|
||||
@@ -186,7 +190,7 @@ std::string convert_open_clip_to_hf_clip(const std::string& name) {
|
||||
replace_suffix();
|
||||
|
||||
open_clip_resblock_prefix = "model.visual.transformer.resblocks.";
|
||||
hf_clip_resblock_prefix = "transformer.visual_model.encoder.layers.";
|
||||
hf_clip_resblock_prefix = "transformer.vision_model.encoder.layers.";
|
||||
|
||||
replace_suffix();
|
||||
|
||||
@@ -200,6 +204,25 @@ std::string convert_vae_decoder_name(const std::string& name) {
|
||||
return name;
|
||||
}
|
||||
|
||||
/* If not a SDXL LoRA the unet" prefix will have already been replaced by this
|
||||
* point and "te2" and "te1" don't seem to appear in non-SDXL only "te_" */
|
||||
std::string convert_sdxl_lora_name(std::string tensor_name) {
|
||||
const std::pair<std::string, std::string> sdxl_lora_name_lookup[] = {
|
||||
{"unet", "model_diffusion_model"},
|
||||
{"te2", "cond_stage_model_1_transformer"},
|
||||
{"te1", "cond_stage_model_transformer"},
|
||||
{"text_encoder_2", "cond_stage_model_1_transformer"},
|
||||
{"text_encoder", "cond_stage_model_transformer"},
|
||||
};
|
||||
for (auto& pair_i : sdxl_lora_name_lookup) {
|
||||
if (tensor_name.compare(0, pair_i.first.length(), pair_i.first) == 0) {
|
||||
tensor_name = std::regex_replace(tensor_name, std::regex(pair_i.first), pair_i.second);
|
||||
break;
|
||||
}
|
||||
}
|
||||
return tensor_name;
|
||||
}
|
||||
|
||||
std::unordered_map<std::string, std::unordered_map<std::string, std::string>> suffix_conversion_underline = {
|
||||
{
|
||||
"attentions",
|
||||
@@ -248,7 +271,7 @@ std::unordered_map<std::string, std::unordered_map<std::string, std::string>> su
|
||||
},
|
||||
};
|
||||
|
||||
std::string convert_diffusers_name_to_compvis(const std::string& key, char seq) {
|
||||
std::string convert_diffusers_name_to_compvis(std::string key, char seq) {
|
||||
std::vector<std::string> m;
|
||||
|
||||
auto match = [](std::vector<std::string>& match_list, const std::regex& regex, const std::string& key) {
|
||||
@@ -282,6 +305,11 @@ std::string convert_diffusers_name_to_compvis(const std::string& key, char seq)
|
||||
return inner_key;
|
||||
};
|
||||
|
||||
// convert attn to out
|
||||
if (ends_with(key, "to_out")) {
|
||||
key += format("%c0", seq);
|
||||
}
|
||||
|
||||
// unet
|
||||
if (match(m, std::regex(format("unet%cconv_in(.*)", seq)), key)) {
|
||||
return format("model%cdiffusion_model%cinput_blocks%c0%c0", seq, seq, seq, seq) + m[0];
|
||||
@@ -391,8 +419,8 @@ std::string convert_diffusers_name_to_compvis(const std::string& key, char seq)
|
||||
}
|
||||
|
||||
std::string convert_tensor_name(const std::string& name) {
|
||||
std::string new_name;
|
||||
if (starts_with(name, "cond_stage_model.") || starts_with(name, "conditioner.embedders.")) {
|
||||
std::string new_name = name;
|
||||
if (starts_with(name, "cond_stage_model.") || starts_with(name, "conditioner.embedders.") || ends_with(name, ".vision_model.visual_projection.weight")) {
|
||||
new_name = convert_open_clip_to_hf_clip(name);
|
||||
} else if (starts_with(name, "first_stage_model.decoder")) {
|
||||
new_name = convert_vae_decoder_name(name);
|
||||
@@ -406,8 +434,12 @@ std::string convert_tensor_name(const std::string& name) {
|
||||
if (pos != std::string::npos) {
|
||||
std::string name_without_network_parts = name.substr(5, pos - 5);
|
||||
std::string network_part = name.substr(pos + 1);
|
||||
|
||||
// LOG_DEBUG("%s %s", name_without_network_parts.c_str(), network_part.c_str());
|
||||
std::string new_key = convert_diffusers_name_to_compvis(name_without_network_parts, '_');
|
||||
/* For dealing with the new SDXL LoRA tensor naming convention */
|
||||
new_key = convert_sdxl_lora_name(new_key);
|
||||
|
||||
if (new_key.empty()) {
|
||||
new_name = name;
|
||||
} else {
|
||||
@@ -416,6 +448,33 @@ std::string convert_tensor_name(const std::string& name) {
|
||||
} else {
|
||||
new_name = name;
|
||||
}
|
||||
} else if (contains(name, "lora_up") || contains(name, "lora_down") ||
|
||||
contains(name, "lora.up") || contains(name, "lora.down") ||
|
||||
contains(name, "lora_linear")) {
|
||||
size_t pos = new_name.find(".processor");
|
||||
if (pos != std::string::npos) {
|
||||
new_name.replace(pos, strlen(".processor"), "");
|
||||
}
|
||||
pos = new_name.rfind("lora");
|
||||
if (pos != std::string::npos) {
|
||||
std::string name_without_network_parts = new_name.substr(0, pos - 1);
|
||||
std::string network_part = new_name.substr(pos);
|
||||
// LOG_DEBUG("%s %s", name_without_network_parts.c_str(), network_part.c_str());
|
||||
std::string new_key = convert_diffusers_name_to_compvis(name_without_network_parts, '.');
|
||||
new_key = convert_sdxl_lora_name(new_key);
|
||||
replace_all_chars(new_key, '.', '_');
|
||||
size_t npos = network_part.rfind("_linear_layer");
|
||||
if (npos != std::string::npos) {
|
||||
network_part.replace(npos, strlen("_linear_layer"), "");
|
||||
}
|
||||
if (starts_with(network_part, "lora.")) {
|
||||
network_part = "lora_" + network_part.substr(5);
|
||||
}
|
||||
if (new_key.size() > 0) {
|
||||
new_name = "lora." + new_key + "." + network_part;
|
||||
}
|
||||
// LOG_DEBUG("new name: %s", new_name.c_str());
|
||||
}
|
||||
} else if (starts_with(name, "unet") || starts_with(name, "vae") || starts_with(name, "te")) { // for diffuser
|
||||
size_t pos = name.find_last_of('.');
|
||||
if (pos != std::string::npos) {
|
||||
@@ -512,10 +571,9 @@ void convert_tensor(void* src,
|
||||
if (dst_type == GGML_TYPE_F16) {
|
||||
ggml_fp32_to_fp16_row((float*)src, (ggml_fp16_t*)dst, n);
|
||||
} else {
|
||||
int64_t hist[16];
|
||||
std::vector<float> imatrix(n_per_row, 1.0f); // dummy importance matrix
|
||||
const float* im = imatrix.data();
|
||||
ggml_quantize_chunk(dst_type, (float*)src, dst, 0, nrows, n_per_row, hist, im);
|
||||
ggml_quantize_chunk(dst_type, (float*)src, dst, 0, nrows, n_per_row, im);
|
||||
}
|
||||
} else if (dst_type == GGML_TYPE_F32) {
|
||||
if (src_type == GGML_TYPE_F16) {
|
||||
@@ -543,10 +601,9 @@ void convert_tensor(void* src,
|
||||
if (dst_type == GGML_TYPE_F16) {
|
||||
ggml_fp32_to_fp16_row((float*)src_data_f32, (ggml_fp16_t*)dst, n);
|
||||
} else {
|
||||
int64_t hist[16];
|
||||
std::vector<float> imatrix(n_per_row, 1.0f); // dummy importance matrix
|
||||
const float* im = imatrix.data();
|
||||
ggml_quantize_chunk(dst_type, (float*)src_data_f32, dst, 0, nrows, n_per_row, hist, im);
|
||||
ggml_quantize_chunk(dst_type, (float*)src_data_f32, dst, 0, nrows, n_per_row, im);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -829,8 +886,12 @@ bool ModelLoader::init_from_safetensors_file(const std::string& file_path, const
|
||||
}
|
||||
}
|
||||
|
||||
TensorStorage tensor_storage(prefix + name, type, ne, n_dims, file_index, ST_HEADER_SIZE_LEN + header_size_ + begin);
|
||||
// ggml_n_dims returns 1 for scalars
|
||||
if (n_dims == 0) {
|
||||
n_dims = 1;
|
||||
}
|
||||
|
||||
TensorStorage tensor_storage(prefix + name, type, ne, n_dims, file_index, ST_HEADER_SIZE_LEN + header_size_ + begin);
|
||||
tensor_storage.reverse_ne();
|
||||
|
||||
size_t tensor_data_size = end - begin;
|
||||
@@ -1169,7 +1230,9 @@ bool ModelLoader::parse_data_pkl(uint8_t* buffer,
|
||||
if (reader.phase == PickleTensorReader::READ_DIMENS) {
|
||||
reader.tensor_storage.reverse_ne();
|
||||
reader.tensor_storage.file_index = file_index;
|
||||
reader.tensor_storage.name = prefix + reader.tensor_storage.name;
|
||||
// if(strcmp(prefix.c_str(), "scarlett") == 0)
|
||||
// printf(" got tensor %s \n ", reader.tensor_storage.name.c_str());
|
||||
reader.tensor_storage.name = prefix + reader.tensor_storage.name;
|
||||
tensor_storages.push_back(reader.tensor_storage);
|
||||
// LOG_DEBUG("%s", reader.tensor_storage.name.c_str());
|
||||
// reset
|
||||
@@ -1272,7 +1335,8 @@ std::string ModelLoader::load_merges() {
|
||||
return merges_utf8_str;
|
||||
}
|
||||
|
||||
void remove_duplicates(std::vector<TensorStorage>& vec) {
|
||||
std::vector<TensorStorage> remove_duplicates(const std::vector<TensorStorage>& vec) {
|
||||
std::vector<TensorStorage> res;
|
||||
std::unordered_map<std::string, size_t> name_to_index_map;
|
||||
|
||||
for (size_t i = 0; i < vec.size(); ++i) {
|
||||
@@ -1280,13 +1344,16 @@ void remove_duplicates(std::vector<TensorStorage>& vec) {
|
||||
auto it = name_to_index_map.find(current_name);
|
||||
|
||||
if (it != name_to_index_map.end()) {
|
||||
vec[it->second] = vec[i];
|
||||
res[it->second] = vec[i];
|
||||
} else {
|
||||
name_to_index_map[current_name] = i;
|
||||
res.push_back(vec[i]);
|
||||
}
|
||||
}
|
||||
|
||||
vec.resize(name_to_index_map.size());
|
||||
// vec.resize(name_to_index_map.size());
|
||||
|
||||
return res;
|
||||
}
|
||||
|
||||
bool ModelLoader::load_tensors(on_new_tensor_cb_t on_new_tensor_cb, ggml_backend_t backend) {
|
||||
@@ -1300,7 +1367,9 @@ bool ModelLoader::load_tensors(on_new_tensor_cb_t on_new_tensor_cb, ggml_backend
|
||||
|
||||
preprocess_tensor(tensor_storage, processed_tensor_storages);
|
||||
}
|
||||
remove_duplicates(processed_tensor_storages);
|
||||
std::vector<TensorStorage> dedup = remove_duplicates(processed_tensor_storages);
|
||||
processed_tensor_storages = dedup;
|
||||
|
||||
bool success = true;
|
||||
for (size_t file_index = 0; file_index < file_paths_.size(); file_index++) {
|
||||
std::string file_path = file_paths_[file_index];
|
||||
@@ -1362,7 +1431,6 @@ bool ModelLoader::load_tensors(on_new_tensor_cb_t on_new_tensor_cb, ggml_backend
|
||||
if (tensor_storage.file_index != file_index) {
|
||||
continue;
|
||||
}
|
||||
|
||||
ggml_tensor* dst_tensor = NULL;
|
||||
|
||||
success = on_new_tensor_cb(tensor_storage, &dst_tensor);
|
||||
@@ -1606,4 +1674,4 @@ bool convert(const char* input_path, const char* vae_path, const char* output_pa
|
||||
}
|
||||
bool success = model_loader.save_to_gguf_file(output_path, (ggml_type)output_type);
|
||||
return success;
|
||||
}
|
||||
}
|
||||
|
||||
1
model.h
@@ -7,6 +7,7 @@
|
||||
#include <set>
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
#include <tuple>
|
||||
#include <vector>
|
||||
|
||||
#include "ggml/ggml-backend.h"
|
||||
|
||||
295
pmid.hpp
Normal file
@@ -0,0 +1,295 @@
|
||||
#ifndef __PMI_HPP__
|
||||
#define __PMI_HPP__
|
||||
|
||||
#include "ggml_extend.hpp"
|
||||
|
||||
#include "clip.hpp"
|
||||
#include "lora.hpp"
|
||||
|
||||
struct FuseBlock : public GGMLBlock {
|
||||
// network hparams
|
||||
int in_dim;
|
||||
int out_dim;
|
||||
int hidden_dim;
|
||||
bool use_residue;
|
||||
|
||||
public:
|
||||
FuseBlock(int i_d, int o_d, int h_d, bool use_residue = true)
|
||||
: in_dim(i_d), out_dim(o_d), hidden_dim(h_d), use_residue(use_residue) {
|
||||
blocks["fc1"] = std::shared_ptr<GGMLBlock>(new Linear(in_dim, hidden_dim, true));
|
||||
blocks["fc2"] = std::shared_ptr<GGMLBlock>(new Linear(hidden_dim, out_dim, true));
|
||||
blocks["layernorm"] = std::shared_ptr<GGMLBlock>(new LayerNorm(in_dim));
|
||||
}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* x) {
|
||||
// x: [N, channels, h, w]
|
||||
|
||||
auto fc1 = std::dynamic_pointer_cast<Linear>(blocks["fc1"]);
|
||||
auto fc2 = std::dynamic_pointer_cast<Linear>(blocks["fc2"]);
|
||||
auto layer_norm = std::dynamic_pointer_cast<LayerNorm>(blocks["layernorm"]);
|
||||
|
||||
struct ggml_tensor* r = x;
|
||||
// x = ggml_nn_layer_norm(ctx, x, ln_w, ln_b);
|
||||
x = layer_norm->forward(ctx, x);
|
||||
// x = ggml_add(ctx, ggml_mul_mat(ctx, fc1_w, x), fc1_b);
|
||||
x = fc1->forward(ctx, x);
|
||||
x = ggml_gelu_inplace(ctx, x);
|
||||
x = fc2->forward(ctx, x);
|
||||
// x = ggml_add(ctx, ggml_mul_mat(ctx, fc2_w, x), fc2_b);
|
||||
if (use_residue)
|
||||
x = ggml_add(ctx, x, r);
|
||||
return x;
|
||||
}
|
||||
};
|
||||
|
||||
struct FuseModule : public GGMLBlock {
|
||||
// network hparams
|
||||
int embed_dim;
|
||||
|
||||
public:
|
||||
FuseModule(int imb_d)
|
||||
: embed_dim(imb_d) {
|
||||
blocks["mlp1"] = std::shared_ptr<GGMLBlock>(new FuseBlock(imb_d * 2, imb_d, imb_d, false));
|
||||
blocks["mlp2"] = std::shared_ptr<GGMLBlock>(new FuseBlock(imb_d, imb_d, imb_d, true));
|
||||
blocks["layer_norm"] = std::shared_ptr<GGMLBlock>(new LayerNorm(embed_dim));
|
||||
}
|
||||
|
||||
struct ggml_tensor* fuse_fn(struct ggml_context* ctx,
|
||||
struct ggml_tensor* prompt_embeds,
|
||||
struct ggml_tensor* id_embeds) {
|
||||
auto mlp1 = std::dynamic_pointer_cast<FuseBlock>(blocks["mlp1"]);
|
||||
auto mlp2 = std::dynamic_pointer_cast<FuseBlock>(blocks["mlp2"]);
|
||||
auto layer_norm = std::dynamic_pointer_cast<LayerNorm>(blocks["layer_norm"]);
|
||||
|
||||
auto prompt_embeds0 = ggml_cont(ctx, ggml_permute(ctx, prompt_embeds, 2, 0, 1, 3));
|
||||
auto id_embeds0 = ggml_cont(ctx, ggml_permute(ctx, id_embeds, 2, 0, 1, 3));
|
||||
// concat is along dim 2
|
||||
auto stacked_id_embeds = ggml_concat(ctx, prompt_embeds0, id_embeds0, 2);
|
||||
stacked_id_embeds = ggml_cont(ctx, ggml_permute(ctx, stacked_id_embeds, 1, 2, 0, 3));
|
||||
|
||||
// stacked_id_embeds = mlp1.forward(ctx, stacked_id_embeds);
|
||||
// stacked_id_embeds = ggml_add(ctx, stacked_id_embeds, prompt_embeds);
|
||||
// stacked_id_embeds = mlp2.forward(ctx, stacked_id_embeds);
|
||||
// stacked_id_embeds = ggml_nn_layer_norm(ctx, stacked_id_embeds, ln_w, ln_b);
|
||||
|
||||
stacked_id_embeds = mlp1->forward(ctx, stacked_id_embeds);
|
||||
stacked_id_embeds = ggml_add(ctx, stacked_id_embeds, prompt_embeds);
|
||||
stacked_id_embeds = mlp2->forward(ctx, stacked_id_embeds);
|
||||
stacked_id_embeds = layer_norm->forward(ctx, stacked_id_embeds);
|
||||
|
||||
return stacked_id_embeds;
|
||||
}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx,
|
||||
struct ggml_tensor* prompt_embeds,
|
||||
struct ggml_tensor* id_embeds,
|
||||
struct ggml_tensor* class_tokens_mask,
|
||||
struct ggml_tensor* class_tokens_mask_pos,
|
||||
struct ggml_tensor* left,
|
||||
struct ggml_tensor* right) {
|
||||
// x: [N, channels, h, w]
|
||||
|
||||
struct ggml_tensor* valid_id_embeds = id_embeds;
|
||||
// # slice out the image token embeddings
|
||||
// print_ggml_tensor(class_tokens_mask_pos, false);
|
||||
ggml_set_name(class_tokens_mask_pos, "class_tokens_mask_pos");
|
||||
ggml_set_name(prompt_embeds, "prompt_embeds");
|
||||
// print_ggml_tensor(valid_id_embeds, true, "valid_id_embeds");
|
||||
// print_ggml_tensor(class_tokens_mask_pos, true, "class_tokens_mask_pos");
|
||||
struct ggml_tensor* image_token_embeds = ggml_get_rows(ctx, prompt_embeds, class_tokens_mask_pos);
|
||||
ggml_set_name(image_token_embeds, "image_token_embeds");
|
||||
struct ggml_tensor* stacked_id_embeds = fuse_fn(ctx, image_token_embeds, valid_id_embeds);
|
||||
|
||||
stacked_id_embeds = ggml_cont(ctx, ggml_permute(ctx, stacked_id_embeds, 0, 2, 1, 3));
|
||||
if (left && right) {
|
||||
stacked_id_embeds = ggml_concat(ctx, left, stacked_id_embeds, 2);
|
||||
stacked_id_embeds = ggml_concat(ctx, stacked_id_embeds, right, 2);
|
||||
} else if (left) {
|
||||
stacked_id_embeds = ggml_concat(ctx, left, stacked_id_embeds, 2);
|
||||
} else if (right) {
|
||||
stacked_id_embeds = ggml_concat(ctx, stacked_id_embeds, right, 2);
|
||||
}
|
||||
stacked_id_embeds = ggml_cont(ctx, ggml_permute(ctx, stacked_id_embeds, 0, 2, 1, 3));
|
||||
class_tokens_mask = ggml_cont(ctx, ggml_transpose(ctx, class_tokens_mask));
|
||||
class_tokens_mask = ggml_repeat(ctx, class_tokens_mask, prompt_embeds);
|
||||
prompt_embeds = ggml_mul(ctx, prompt_embeds, class_tokens_mask);
|
||||
struct ggml_tensor* updated_prompt_embeds = ggml_add(ctx, prompt_embeds, stacked_id_embeds);
|
||||
ggml_set_name(updated_prompt_embeds, "updated_prompt_embeds");
|
||||
return updated_prompt_embeds;
|
||||
}
|
||||
};
|
||||
|
||||
struct PhotoMakerIDEncoderBlock : public CLIPVisionModelProjection {
|
||||
PhotoMakerIDEncoderBlock()
|
||||
: CLIPVisionModelProjection(OPENAI_CLIP_VIT_L_14) {
|
||||
blocks["visual_projection_2"] = std::shared_ptr<GGMLBlock>(new Linear(1024, 1280, false));
|
||||
blocks["fuse_module"] = std::shared_ptr<GGMLBlock>(new FuseModule(2048));
|
||||
}
|
||||
|
||||
struct ggml_tensor* forward(struct ggml_context* ctx,
|
||||
struct ggml_tensor* id_pixel_values,
|
||||
struct ggml_tensor* prompt_embeds,
|
||||
struct ggml_tensor* class_tokens_mask,
|
||||
struct ggml_tensor* class_tokens_mask_pos,
|
||||
struct ggml_tensor* left,
|
||||
struct ggml_tensor* right) {
|
||||
// x: [N, channels, h, w]
|
||||
auto vision_model = std::dynamic_pointer_cast<CLIPVisionModel>(blocks["vision_model"]);
|
||||
auto visual_projection = std::dynamic_pointer_cast<CLIPProjection>(blocks["visual_projection"]);
|
||||
auto visual_projection_2 = std::dynamic_pointer_cast<Linear>(blocks["visual_projection_2"]);
|
||||
auto fuse_module = std::dynamic_pointer_cast<FuseModule>(blocks["fuse_module"]);
|
||||
|
||||
struct ggml_tensor* shared_id_embeds = vision_model->forward(ctx, id_pixel_values); // [N, hidden_size]
|
||||
struct ggml_tensor* id_embeds = visual_projection->forward(ctx, shared_id_embeds); // [N, proj_dim(768)]
|
||||
struct ggml_tensor* id_embeds_2 = visual_projection_2->forward(ctx, shared_id_embeds); // [N, 1280]
|
||||
|
||||
id_embeds = ggml_cont(ctx, ggml_permute(ctx, id_embeds, 2, 0, 1, 3));
|
||||
id_embeds_2 = ggml_cont(ctx, ggml_permute(ctx, id_embeds_2, 2, 0, 1, 3));
|
||||
|
||||
id_embeds = ggml_concat(ctx, id_embeds, id_embeds_2, 2); // [batch_size, seq_length, 1, 2048] check whether concat at dim 2 is right
|
||||
id_embeds = ggml_cont(ctx, ggml_permute(ctx, id_embeds, 1, 2, 0, 3));
|
||||
|
||||
struct ggml_tensor* updated_prompt_embeds = fuse_module->forward(ctx,
|
||||
prompt_embeds,
|
||||
id_embeds,
|
||||
class_tokens_mask,
|
||||
class_tokens_mask_pos,
|
||||
left, right);
|
||||
return updated_prompt_embeds;
|
||||
}
|
||||
};
|
||||
|
||||
struct PhotoMakerIDEncoder : public GGMLModule {
|
||||
public:
|
||||
SDVersion version = VERSION_XL;
|
||||
PhotoMakerIDEncoderBlock id_encoder;
|
||||
float style_strength;
|
||||
|
||||
std::vector<float> ctm;
|
||||
std::vector<ggml_fp16_t> ctmf16;
|
||||
std::vector<int> ctmpos;
|
||||
|
||||
std::vector<ggml_fp16_t> zeros_left_16;
|
||||
std::vector<float> zeros_left;
|
||||
std::vector<ggml_fp16_t> zeros_right_16;
|
||||
std::vector<float> zeros_right;
|
||||
|
||||
public:
|
||||
PhotoMakerIDEncoder(ggml_backend_t backend, ggml_type wtype, SDVersion version = VERSION_XL, float sty = 20.f)
|
||||
: GGMLModule(backend, wtype),
|
||||
version(version),
|
||||
style_strength(sty) {
|
||||
id_encoder.init(params_ctx, wtype);
|
||||
}
|
||||
|
||||
std::string get_desc() {
|
||||
return "pmid";
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, struct ggml_tensor*>& tensors, const std::string prefix) {
|
||||
id_encoder.get_param_tensors(tensors, prefix);
|
||||
}
|
||||
|
||||
struct ggml_cgraph* build_graph( // struct ggml_allocr* allocr,
|
||||
struct ggml_tensor* id_pixel_values,
|
||||
struct ggml_tensor* prompt_embeds,
|
||||
std::vector<bool>& class_tokens_mask) {
|
||||
ctm.clear();
|
||||
ctmf16.clear();
|
||||
ctmpos.clear();
|
||||
zeros_left.clear();
|
||||
zeros_left_16.clear();
|
||||
zeros_right.clear();
|
||||
zeros_right_16.clear();
|
||||
|
||||
ggml_context* ctx0 = compute_ctx;
|
||||
|
||||
struct ggml_cgraph* gf = ggml_new_graph(compute_ctx);
|
||||
|
||||
int64_t hidden_size = prompt_embeds->ne[0];
|
||||
int64_t seq_length = prompt_embeds->ne[1];
|
||||
ggml_type type = GGML_TYPE_F32;
|
||||
|
||||
struct ggml_tensor* class_tokens_mask_d = ggml_new_tensor_1d(ctx0, type, class_tokens_mask.size());
|
||||
|
||||
struct ggml_tensor* id_pixel_values_d = to_backend(id_pixel_values);
|
||||
struct ggml_tensor* prompt_embeds_d = to_backend(prompt_embeds);
|
||||
|
||||
struct ggml_tensor* left = NULL;
|
||||
struct ggml_tensor* right = NULL;
|
||||
for (int i = 0; i < class_tokens_mask.size(); i++) {
|
||||
if (class_tokens_mask[i]) {
|
||||
ctm.push_back(0.f); // here use 0.f instead of 1.f to make a scale mask
|
||||
ctmf16.push_back(ggml_fp32_to_fp16(0.f)); // here use 0.f instead of 1.f to make a scale mask
|
||||
ctmpos.push_back(i);
|
||||
} else {
|
||||
ctm.push_back(1.f); // here use 1.f instead of 0.f to make a scale mask
|
||||
ctmf16.push_back(ggml_fp32_to_fp16(1.f)); // here use 0.f instead of 1.f to make a scale mask
|
||||
}
|
||||
}
|
||||
if (ctmpos[0] > 0) {
|
||||
left = ggml_new_tensor_3d(ctx0, type, hidden_size, 1, ctmpos[0]);
|
||||
}
|
||||
if (ctmpos[ctmpos.size() - 1] < seq_length - 1) {
|
||||
right = ggml_new_tensor_3d(ctx0, type,
|
||||
hidden_size, 1, seq_length - ctmpos[ctmpos.size() - 1] - 1);
|
||||
}
|
||||
struct ggml_tensor* class_tokens_mask_pos = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, ctmpos.size());
|
||||
|
||||
{
|
||||
if (type == GGML_TYPE_F16)
|
||||
set_backend_tensor_data(class_tokens_mask_d, ctmf16.data());
|
||||
else
|
||||
set_backend_tensor_data(class_tokens_mask_d, ctm.data());
|
||||
set_backend_tensor_data(class_tokens_mask_pos, ctmpos.data());
|
||||
if (left) {
|
||||
if (type == GGML_TYPE_F16) {
|
||||
for (int i = 0; i < ggml_nelements(left); ++i)
|
||||
zeros_left_16.push_back(ggml_fp32_to_fp16(0.f));
|
||||
set_backend_tensor_data(left, zeros_left_16.data());
|
||||
} else {
|
||||
for (int i = 0; i < ggml_nelements(left); ++i)
|
||||
zeros_left.push_back(0.f);
|
||||
set_backend_tensor_data(left, zeros_left.data());
|
||||
}
|
||||
}
|
||||
if (right) {
|
||||
if (type == GGML_TYPE_F16) {
|
||||
for (int i = 0; i < ggml_nelements(right); ++i)
|
||||
zeros_right_16.push_back(ggml_fp32_to_fp16(0.f));
|
||||
set_backend_tensor_data(right, zeros_right_16.data());
|
||||
} else {
|
||||
for (int i = 0; i < ggml_nelements(right); ++i)
|
||||
zeros_right.push_back(0.f);
|
||||
set_backend_tensor_data(right, zeros_right.data());
|
||||
}
|
||||
}
|
||||
}
|
||||
struct ggml_tensor* updated_prompt_embeds = id_encoder.forward(ctx0,
|
||||
id_pixel_values_d,
|
||||
prompt_embeds_d,
|
||||
class_tokens_mask_d,
|
||||
class_tokens_mask_pos,
|
||||
left, right);
|
||||
ggml_build_forward_expand(gf, updated_prompt_embeds);
|
||||
|
||||
return gf;
|
||||
}
|
||||
|
||||
void compute(const int n_threads,
|
||||
struct ggml_tensor* id_pixel_values,
|
||||
struct ggml_tensor* prompt_embeds,
|
||||
std::vector<bool>& class_tokens_mask,
|
||||
struct ggml_tensor** updated_prompt_embeds,
|
||||
ggml_context* output_ctx) {
|
||||
auto get_graph = [&]() -> struct ggml_cgraph* {
|
||||
// return build_graph(compute_allocr, id_pixel_values, prompt_embeds, class_tokens_mask);
|
||||
return build_graph(id_pixel_values, prompt_embeds, class_tokens_mask);
|
||||
};
|
||||
|
||||
// GGMLModule::compute(get_graph, n_threads, updated_prompt_embeds);
|
||||
GGMLModule::compute(get_graph, n_threads, true, updated_prompt_embeds, output_ctx);
|
||||
}
|
||||
};
|
||||
|
||||
#endif // __PMI_HPP__
|
||||
1103
stable-diffusion.cpp
@@ -49,6 +49,7 @@ enum schedule_t {
|
||||
DEFAULT,
|
||||
DISCRETE,
|
||||
KARRAS,
|
||||
AYS,
|
||||
N_SCHEDULES
|
||||
};
|
||||
|
||||
@@ -59,18 +60,17 @@ enum sd_type_t {
|
||||
SD_TYPE_Q4_0 = 2,
|
||||
SD_TYPE_Q4_1 = 3,
|
||||
// SD_TYPE_Q4_2 = 4, support has been removed
|
||||
// SD_TYPE_Q4_3 (5) support has been removed
|
||||
SD_TYPE_Q5_0 = 6,
|
||||
SD_TYPE_Q5_1 = 7,
|
||||
SD_TYPE_Q8_0 = 8,
|
||||
SD_TYPE_Q8_1 = 9,
|
||||
// k-quantizations
|
||||
SD_TYPE_Q2_K = 10,
|
||||
SD_TYPE_Q3_K = 11,
|
||||
SD_TYPE_Q4_K = 12,
|
||||
SD_TYPE_Q5_K = 13,
|
||||
SD_TYPE_Q6_K = 14,
|
||||
SD_TYPE_Q8_K = 15,
|
||||
// SD_TYPE_Q4_3 = 5, support has been removed
|
||||
SD_TYPE_Q5_0 = 6,
|
||||
SD_TYPE_Q5_1 = 7,
|
||||
SD_TYPE_Q8_0 = 8,
|
||||
SD_TYPE_Q8_1 = 9,
|
||||
SD_TYPE_Q2_K = 10,
|
||||
SD_TYPE_Q3_K = 11,
|
||||
SD_TYPE_Q4_K = 12,
|
||||
SD_TYPE_Q5_K = 13,
|
||||
SD_TYPE_Q6_K = 14,
|
||||
SD_TYPE_Q8_K = 15,
|
||||
SD_TYPE_IQ2_XXS = 16,
|
||||
SD_TYPE_IQ2_XS = 17,
|
||||
SD_TYPE_IQ3_XXS = 18,
|
||||
@@ -79,9 +79,13 @@ enum sd_type_t {
|
||||
SD_TYPE_IQ3_S = 21,
|
||||
SD_TYPE_IQ2_S = 22,
|
||||
SD_TYPE_IQ4_XS = 23,
|
||||
SD_TYPE_I8,
|
||||
SD_TYPE_I16,
|
||||
SD_TYPE_I32,
|
||||
SD_TYPE_I8 = 24,
|
||||
SD_TYPE_I16 = 25,
|
||||
SD_TYPE_I32 = 26,
|
||||
SD_TYPE_I64 = 27,
|
||||
SD_TYPE_F64 = 28,
|
||||
SD_TYPE_IQ1_M = 29,
|
||||
SD_TYPE_BF16 = 30,
|
||||
SD_TYPE_COUNT,
|
||||
};
|
||||
|
||||
@@ -95,7 +99,7 @@ enum sd_log_level_t {
|
||||
};
|
||||
|
||||
typedef void (*sd_log_cb_t)(enum sd_log_level_t level, const char* text, void* data);
|
||||
typedef void (*sd_progress_cb_t)(int step,int steps,float time, void* data);
|
||||
typedef void (*sd_progress_cb_t)(int step, int steps, float time, void* data);
|
||||
|
||||
SD_API void sd_set_log_callback(sd_log_cb_t sd_log_cb, void* data);
|
||||
SD_API void sd_set_progress_callback(sd_progress_cb_t cb, void* data);
|
||||
@@ -117,6 +121,7 @@ SD_API sd_ctx_t* new_sd_ctx(const char* model_path,
|
||||
const char* control_net_path_c_str,
|
||||
const char* lora_model_dir,
|
||||
const char* embed_dir_c_str,
|
||||
const char* stacked_id_embed_dir_c_str,
|
||||
bool vae_decode_only,
|
||||
bool vae_tiling,
|
||||
bool free_params_immediately,
|
||||
@@ -124,7 +129,9 @@ SD_API sd_ctx_t* new_sd_ctx(const char* model_path,
|
||||
enum sd_type_t wtype,
|
||||
enum rng_type_t rng_type,
|
||||
enum schedule_t s,
|
||||
bool keep_control_net_cpu);
|
||||
bool keep_clip_on_cpu,
|
||||
bool keep_control_net_cpu,
|
||||
bool keep_vae_on_cpu);
|
||||
|
||||
SD_API void free_sd_ctx(sd_ctx_t* sd_ctx);
|
||||
|
||||
@@ -140,7 +147,10 @@ SD_API sd_image_t* txt2img(sd_ctx_t* sd_ctx,
|
||||
int64_t seed,
|
||||
int batch_count,
|
||||
const sd_image_t* control_cond,
|
||||
float control_strength);
|
||||
float control_strength,
|
||||
float style_strength,
|
||||
bool normalize_input,
|
||||
const char* input_id_images_path);
|
||||
|
||||
SD_API sd_image_t* img2img(sd_ctx_t* sd_ctx,
|
||||
sd_image_t init_image,
|
||||
@@ -154,7 +164,12 @@ SD_API sd_image_t* img2img(sd_ctx_t* sd_ctx,
|
||||
int sample_steps,
|
||||
float strength,
|
||||
int64_t seed,
|
||||
int batch_count);
|
||||
int batch_count,
|
||||
const sd_image_t* control_cond,
|
||||
float control_strength,
|
||||
float style_strength,
|
||||
bool normalize_input,
|
||||
const char* input_id_images_path);
|
||||
|
||||
SD_API sd_image_t* img2vid(sd_ctx_t* sd_ctx,
|
||||
sd_image_t init_image,
|
||||
@@ -195,4 +210,4 @@ SD_API uint8_t* preprocess_canny(uint8_t* img,
|
||||
}
|
||||
#endif
|
||||
|
||||
#endif // __STABLE_DIFFUSION_H__
|
||||
#endif // __STABLE_DIFFUSION_H__
|
||||
|
||||
10
tae.hpp
@@ -200,16 +200,8 @@ struct TinyAutoEncoder : public GGMLModule {
|
||||
return "taesd";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return taesd.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return taesd.get_params_num();
|
||||
}
|
||||
|
||||
bool load_from_file(const std::string& file_path) {
|
||||
LOG_INFO("loading taesd from '%s'", file_path.c_str());
|
||||
LOG_INFO("loading taesd from '%s', decode_only = %s", file_path.c_str(), decode_only ? "true" : "false");
|
||||
alloc_params_buffer();
|
||||
std::map<std::string, ggml_tensor*> taesd_tensors;
|
||||
taesd.get_param_tensors(taesd_tensors);
|
||||
|
||||
2
thirdparty/.clang-format
vendored
Normal file
@@ -0,0 +1,2 @@
|
||||
DisableFormat: true
|
||||
SortIncludes: Never
|
||||
2585
thirdparty/stb_image_resize.h
vendored
Normal file
232
thirdparty/zip.c
vendored
@@ -36,6 +36,7 @@
|
||||
#include <unistd.h>
|
||||
#endif
|
||||
|
||||
#define USE_EXTERNAL_MZCRC
|
||||
#include "miniz.h"
|
||||
#include "zip.h"
|
||||
|
||||
@@ -1834,3 +1835,234 @@ int zip_extract(const char *zipname, const char *dir,
|
||||
|
||||
return zip_archive_extract(&zip_archive, dir, on_extract, arg);
|
||||
}
|
||||
|
||||
#if defined(__SSE4_2__) || defined(__AVX512F__)
|
||||
#include <immintrin.h>
|
||||
#endif
|
||||
|
||||
// Phil Katz 32-Bit Cyclic Redundancy Check Uber Alles
|
||||
// Goes 73 GiB/s on an AMD Ryzen Threadripper PRO 7995WX
|
||||
// "Fast CRC Computation for Generic Polynomials Using PCLMULQDQ Instruction"
|
||||
// V. Gopal, E. Ozturk, et al., 2009, http://intel.ly/2ySEwL0
|
||||
mz_ulong mz_crc32(mz_ulong init, const uint8_t *buf, size_t len) {
|
||||
uint32_t crc = ~init;
|
||||
#if defined(__AVX512F__) && defined(__VPCLMULQDQ__) && defined(__PCLMUL__)
|
||||
if (len >= 256) {
|
||||
_Alignas(__m512) static const uint64_t k1k2[] = {
|
||||
0x011542778a, 0x01322d1430, 0x011542778a, 0x01322d1430,
|
||||
0x011542778a, 0x01322d1430, 0x011542778a, 0x01322d1430,
|
||||
};
|
||||
_Alignas(__m512) static const uint64_t k3k4[] = {
|
||||
0x0154442bd4, 0x01c6e41596, 0x0154442bd4, 0x01c6e41596,
|
||||
0x0154442bd4, 0x01c6e41596, 0x0154442bd4, 0x01c6e41596,
|
||||
};
|
||||
_Alignas(__m512) static const uint64_t k5k6[] = {
|
||||
0x01751997d0,
|
||||
0x00ccaa009e,
|
||||
};
|
||||
_Alignas(__m512) static const uint64_t k7k8[] = {
|
||||
0x0163cd6124,
|
||||
0x0000000000,
|
||||
};
|
||||
_Alignas(__m512) static const uint64_t poly[] = {
|
||||
0x01db710641,
|
||||
0x01f7011641,
|
||||
};
|
||||
__m512i x0, x1, x2, x3, x4, x5, x6, x7, x8, y5, y6, y7, y8;
|
||||
__m128i a0, a1, a2, a3;
|
||||
x1 = _mm512_loadu_si512((__m512i *)(buf + 0x00));
|
||||
x2 = _mm512_loadu_si512((__m512i *)(buf + 0x40));
|
||||
x3 = _mm512_loadu_si512((__m512i *)(buf + 0x80));
|
||||
x4 = _mm512_loadu_si512((__m512i *)(buf + 0xC0));
|
||||
x1 = _mm512_xor_si512(x1, _mm512_castsi128_si512(_mm_cvtsi32_si128(crc)));
|
||||
x0 = _mm512_load_si512((__m512i *)k1k2);
|
||||
buf += 256;
|
||||
len -= 256;
|
||||
while (len >= 256) {
|
||||
x5 = _mm512_clmulepi64_epi128(x1, x0, 0x00);
|
||||
x6 = _mm512_clmulepi64_epi128(x2, x0, 0x00);
|
||||
x7 = _mm512_clmulepi64_epi128(x3, x0, 0x00);
|
||||
x8 = _mm512_clmulepi64_epi128(x4, x0, 0x00);
|
||||
x1 = _mm512_clmulepi64_epi128(x1, x0, 0x11);
|
||||
x2 = _mm512_clmulepi64_epi128(x2, x0, 0x11);
|
||||
x3 = _mm512_clmulepi64_epi128(x3, x0, 0x11);
|
||||
x4 = _mm512_clmulepi64_epi128(x4, x0, 0x11);
|
||||
y5 = _mm512_loadu_si512((__m512i *)(buf + 0x00));
|
||||
y6 = _mm512_loadu_si512((__m512i *)(buf + 0x40));
|
||||
y7 = _mm512_loadu_si512((__m512i *)(buf + 0x80));
|
||||
y8 = _mm512_loadu_si512((__m512i *)(buf + 0xC0));
|
||||
x1 = _mm512_xor_si512(x1, x5);
|
||||
x2 = _mm512_xor_si512(x2, x6);
|
||||
x3 = _mm512_xor_si512(x3, x7);
|
||||
x4 = _mm512_xor_si512(x4, x8);
|
||||
x1 = _mm512_xor_si512(x1, y5);
|
||||
x2 = _mm512_xor_si512(x2, y6);
|
||||
x3 = _mm512_xor_si512(x3, y7);
|
||||
x4 = _mm512_xor_si512(x4, y8);
|
||||
buf += 256;
|
||||
len -= 256;
|
||||
}
|
||||
x0 = _mm512_load_si512((__m512i *)k3k4);
|
||||
x5 = _mm512_clmulepi64_epi128(x1, x0, 0x00);
|
||||
x1 = _mm512_clmulepi64_epi128(x1, x0, 0x11);
|
||||
x1 = _mm512_xor_si512(x1, x2);
|
||||
x1 = _mm512_xor_si512(x1, x5);
|
||||
x5 = _mm512_clmulepi64_epi128(x1, x0, 0x00);
|
||||
x1 = _mm512_clmulepi64_epi128(x1, x0, 0x11);
|
||||
x1 = _mm512_xor_si512(x1, x3);
|
||||
x1 = _mm512_xor_si512(x1, x5);
|
||||
x5 = _mm512_clmulepi64_epi128(x1, x0, 0x00);
|
||||
x1 = _mm512_clmulepi64_epi128(x1, x0, 0x11);
|
||||
x1 = _mm512_xor_si512(x1, x4);
|
||||
x1 = _mm512_xor_si512(x1, x5);
|
||||
while (len >= 64) {
|
||||
x2 = _mm512_loadu_si512((__m512i *)buf);
|
||||
x5 = _mm512_clmulepi64_epi128(x1, x0, 0x00);
|
||||
x1 = _mm512_clmulepi64_epi128(x1, x0, 0x11);
|
||||
x1 = _mm512_xor_si512(x1, x2);
|
||||
x1 = _mm512_xor_si512(x1, x5);
|
||||
buf += 64;
|
||||
len -= 64;
|
||||
}
|
||||
a0 = _mm_load_si128((__m128i *)k5k6);
|
||||
a1 = _mm512_extracti32x4_epi32(x1, 0);
|
||||
a2 = _mm512_extracti32x4_epi32(x1, 1);
|
||||
a3 = _mm_clmulepi64_si128(a1, a0, 0x00);
|
||||
a1 = _mm_clmulepi64_si128(a1, a0, 0x11);
|
||||
a1 = _mm_xor_si128(a1, a3);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
a2 = _mm512_extracti32x4_epi32(x1, 2);
|
||||
a3 = _mm_clmulepi64_si128(a1, a0, 0x00);
|
||||
a1 = _mm_clmulepi64_si128(a1, a0, 0x11);
|
||||
a1 = _mm_xor_si128(a1, a3);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
a2 = _mm512_extracti32x4_epi32(x1, 3);
|
||||
a3 = _mm_clmulepi64_si128(a1, a0, 0x00);
|
||||
a1 = _mm_clmulepi64_si128(a1, a0, 0x11);
|
||||
a1 = _mm_xor_si128(a1, a3);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
a2 = _mm_clmulepi64_si128(a1, a0, 0x10);
|
||||
a3 = _mm_setr_epi32(~0, 0, ~0, 0);
|
||||
a1 = _mm_srli_si128(a1, 8);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
a0 = _mm_loadl_epi64((__m128i *)k7k8);
|
||||
a2 = _mm_srli_si128(a1, 4);
|
||||
a1 = _mm_and_si128(a1, a3);
|
||||
a1 = _mm_clmulepi64_si128(a1, a0, 0x00);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
a0 = _mm_load_si128((__m128i *)poly);
|
||||
a2 = _mm_and_si128(a1, a3);
|
||||
a2 = _mm_clmulepi64_si128(a2, a0, 0x10);
|
||||
a2 = _mm_and_si128(a2, a3);
|
||||
a2 = _mm_clmulepi64_si128(a2, a0, 0x00);
|
||||
a1 = _mm_xor_si128(a1, a2);
|
||||
crc = _mm_extract_epi32(a1, 1);
|
||||
}
|
||||
#endif
|
||||
#if defined(__SSE4_2__) && defined(__PCLMUL__)
|
||||
if (len >= 64) {
|
||||
_Alignas(__m128) static const uint64_t k1k2[] = {
|
||||
0x0154442bd4,
|
||||
0x01c6e41596,
|
||||
};
|
||||
_Alignas(__m128) static const uint64_t k3k4[] = {
|
||||
0x01751997d0,
|
||||
0x00ccaa009e,
|
||||
};
|
||||
_Alignas(__m128) static const uint64_t k5k0[] = {
|
||||
0x0163cd6124,
|
||||
0x0000000000,
|
||||
};
|
||||
_Alignas(__m128) static const uint64_t poly[] = {
|
||||
0x01db710641,
|
||||
0x01f7011641,
|
||||
};
|
||||
__m128i x0, x1, x2, x3, x4, x5, x6, x7, x8, y5, y6, y7, y8;
|
||||
x1 = _mm_loadu_si128((__m128i *)(buf + 0x00));
|
||||
x2 = _mm_loadu_si128((__m128i *)(buf + 0x10));
|
||||
x3 = _mm_loadu_si128((__m128i *)(buf + 0x20));
|
||||
x4 = _mm_loadu_si128((__m128i *)(buf + 0x30));
|
||||
x1 = _mm_xor_si128(x1, _mm_cvtsi32_si128(crc));
|
||||
x0 = _mm_load_si128((__m128i *)k1k2);
|
||||
buf += 64;
|
||||
len -= 64;
|
||||
while (len >= 64) {
|
||||
x5 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x6 = _mm_clmulepi64_si128(x2, x0, 0x00);
|
||||
x7 = _mm_clmulepi64_si128(x3, x0, 0x00);
|
||||
x8 = _mm_clmulepi64_si128(x4, x0, 0x00);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x11);
|
||||
x2 = _mm_clmulepi64_si128(x2, x0, 0x11);
|
||||
x3 = _mm_clmulepi64_si128(x3, x0, 0x11);
|
||||
x4 = _mm_clmulepi64_si128(x4, x0, 0x11);
|
||||
y5 = _mm_loadu_si128((__m128i *)(buf + 0x00));
|
||||
y6 = _mm_loadu_si128((__m128i *)(buf + 0x10));
|
||||
y7 = _mm_loadu_si128((__m128i *)(buf + 0x20));
|
||||
y8 = _mm_loadu_si128((__m128i *)(buf + 0x30));
|
||||
x1 = _mm_xor_si128(x1, x5);
|
||||
x2 = _mm_xor_si128(x2, x6);
|
||||
x3 = _mm_xor_si128(x3, x7);
|
||||
x4 = _mm_xor_si128(x4, x8);
|
||||
x1 = _mm_xor_si128(x1, y5);
|
||||
x2 = _mm_xor_si128(x2, y6);
|
||||
x3 = _mm_xor_si128(x3, y7);
|
||||
x4 = _mm_xor_si128(x4, y8);
|
||||
buf += 64;
|
||||
len -= 64;
|
||||
}
|
||||
x0 = _mm_load_si128((__m128i *)k3k4);
|
||||
x5 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x11);
|
||||
x1 = _mm_xor_si128(x1, x2);
|
||||
x1 = _mm_xor_si128(x1, x5);
|
||||
x5 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x11);
|
||||
x1 = _mm_xor_si128(x1, x3);
|
||||
x1 = _mm_xor_si128(x1, x5);
|
||||
x5 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x11);
|
||||
x1 = _mm_xor_si128(x1, x4);
|
||||
x1 = _mm_xor_si128(x1, x5);
|
||||
while (len >= 16) {
|
||||
x2 = _mm_loadu_si128((__m128i *)buf);
|
||||
x5 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x11);
|
||||
x1 = _mm_xor_si128(x1, x2);
|
||||
x1 = _mm_xor_si128(x1, x5);
|
||||
buf += 16;
|
||||
len -= 16;
|
||||
}
|
||||
x2 = _mm_clmulepi64_si128(x1, x0, 0x10);
|
||||
x3 = _mm_setr_epi32(~0, 0, ~0, 0);
|
||||
x1 = _mm_srli_si128(x1, 8);
|
||||
x1 = _mm_xor_si128(x1, x2);
|
||||
x0 = _mm_loadl_epi64((__m128i *)k5k0);
|
||||
x2 = _mm_srli_si128(x1, 4);
|
||||
x1 = _mm_and_si128(x1, x3);
|
||||
x1 = _mm_clmulepi64_si128(x1, x0, 0x00);
|
||||
x1 = _mm_xor_si128(x1, x2);
|
||||
x0 = _mm_load_si128((__m128i *)poly);
|
||||
x2 = _mm_and_si128(x1, x3);
|
||||
x2 = _mm_clmulepi64_si128(x2, x0, 0x10);
|
||||
x2 = _mm_and_si128(x2, x3);
|
||||
x2 = _mm_clmulepi64_si128(x2, x0, 0x00);
|
||||
x1 = _mm_xor_si128(x1, x2);
|
||||
crc = _mm_extract_epi32(x1, 1);
|
||||
}
|
||||
#endif
|
||||
static uint32_t tab[256];
|
||||
if (!tab[255]) {
|
||||
// generates table for byte-wise crc calculation on the polynomial
|
||||
// x^32+x^26+x^23+x^22+x^16+x^12+x^11+x^10+x^8+x^7+x^5+x^4+x^2+x+1
|
||||
uint32_t polynomial = 0xedb88320; // bits are reversed
|
||||
for (int d = 0; d < 256; ++d) {
|
||||
uint32_t r = d;
|
||||
for (int i = 0; i < 8; ++i)
|
||||
r = r >> 1 ^ (r & 1 ? polynomial : 0);
|
||||
tab[d] = r;
|
||||
}
|
||||
}
|
||||
for (size_t i = 0; i < len; ++i)
|
||||
crc = crc >> 8 ^ tab[(crc & 255) ^ buf[i]];
|
||||
return ~crc & 0xffffffff;
|
||||
}
|
||||
|
||||
12
unet.hpp
@@ -396,7 +396,7 @@ public:
|
||||
if (c_concat->ne[3] != x->ne[3]) {
|
||||
c_concat = ggml_repeat(ctx, c_concat, x);
|
||||
}
|
||||
x = ggml_concat(ctx, x, c_concat);
|
||||
x = ggml_concat(ctx, x, c_concat, 2);
|
||||
}
|
||||
|
||||
if (y != NULL) {
|
||||
@@ -491,7 +491,7 @@ public:
|
||||
control_offset--;
|
||||
}
|
||||
|
||||
h = ggml_concat(ctx, h, h_skip);
|
||||
h = ggml_concat(ctx, h, h_skip, 2);
|
||||
|
||||
std::string name = "output_blocks." + std::to_string(output_block_idx) + ".0";
|
||||
|
||||
@@ -543,14 +543,6 @@ struct UNetModel : public GGMLModule {
|
||||
return "unet";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return unet.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return unet.get_params_num();
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, struct ggml_tensor*>& tensors, const std::string prefix) {
|
||||
unet.get_param_tensors(tensors, prefix);
|
||||
}
|
||||
|
||||
127
util.cpp
@@ -25,6 +25,9 @@
|
||||
#include "ggml/ggml.h"
|
||||
#include "stable-diffusion.h"
|
||||
|
||||
#define STB_IMAGE_RESIZE_IMPLEMENTATION
|
||||
#include "stb_image_resize.h"
|
||||
|
||||
bool ends_with(const std::string& str, const std::string& ending) {
|
||||
if (str.length() >= ending.length()) {
|
||||
return (str.compare(str.length() - ending.length(), ending.length(), ending) == 0);
|
||||
@@ -40,6 +43,13 @@ bool starts_with(const std::string& str, const std::string& start) {
|
||||
return false;
|
||||
}
|
||||
|
||||
bool contains(const std::string& str, const std::string& substr) {
|
||||
if (str.find(substr) != std::string::npos) {
|
||||
return true;
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
void replace_all_chars(std::string& str, char target, char replacement) {
|
||||
for (size_t i = 0; i < str.length(); ++i) {
|
||||
if (str[i] == target) {
|
||||
@@ -88,6 +98,43 @@ std::string get_full_path(const std::string& dir, const std::string& filename) {
|
||||
}
|
||||
}
|
||||
|
||||
std::vector<std::string> get_files_from_dir(const std::string& dir) {
|
||||
std::vector<std::string> files;
|
||||
|
||||
WIN32_FIND_DATA findFileData;
|
||||
HANDLE hFind;
|
||||
|
||||
char currentDirectory[MAX_PATH];
|
||||
GetCurrentDirectory(MAX_PATH, currentDirectory);
|
||||
|
||||
char directoryPath[MAX_PATH]; // this is absolute path
|
||||
sprintf(directoryPath, "%s\\%s\\*", currentDirectory, dir.c_str());
|
||||
|
||||
// Find the first file in the directory
|
||||
hFind = FindFirstFile(directoryPath, &findFileData);
|
||||
|
||||
// Check if the directory was found
|
||||
if (hFind == INVALID_HANDLE_VALUE) {
|
||||
printf("Unable to find directory.\n");
|
||||
return files;
|
||||
}
|
||||
|
||||
// Loop through all files in the directory
|
||||
do {
|
||||
// Check if the found file is a regular file (not a directory)
|
||||
if (!(findFileData.dwFileAttributes & FILE_ATTRIBUTE_DIRECTORY)) {
|
||||
files.push_back(std::string(currentDirectory) + "\\" + dir + "\\" + std::string(findFileData.cFileName));
|
||||
}
|
||||
} while (FindNextFile(hFind, &findFileData) != 0);
|
||||
|
||||
// Close the handle
|
||||
FindClose(hFind);
|
||||
|
||||
sort(files.begin(), files.end());
|
||||
|
||||
return files;
|
||||
}
|
||||
|
||||
#else // Unix
|
||||
#include <dirent.h>
|
||||
#include <sys/stat.h>
|
||||
@@ -102,6 +149,7 @@ bool is_directory(const std::string& path) {
|
||||
return (stat(path.c_str(), &buffer) == 0 && S_ISDIR(buffer.st_mode));
|
||||
}
|
||||
|
||||
// TODO: add windows version
|
||||
std::string get_full_path(const std::string& dir, const std::string& filename) {
|
||||
DIR* dp = opendir(dir.c_str());
|
||||
|
||||
@@ -121,6 +169,27 @@ std::string get_full_path(const std::string& dir, const std::string& filename) {
|
||||
return "";
|
||||
}
|
||||
|
||||
std::vector<std::string> get_files_from_dir(const std::string& dir) {
|
||||
std::vector<std::string> files;
|
||||
|
||||
DIR* dp = opendir(dir.c_str());
|
||||
|
||||
if (dp != nullptr) {
|
||||
struct dirent* entry;
|
||||
|
||||
while ((entry = readdir(dp)) != nullptr) {
|
||||
std::string fname = dir + "/" + entry->d_name;
|
||||
if (!is_directory(fname))
|
||||
files.push_back(fname);
|
||||
}
|
||||
closedir(dp);
|
||||
}
|
||||
|
||||
sort(files.begin(), files.end());
|
||||
|
||||
return files;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
// get_num_physical_cores is copy from
|
||||
@@ -161,8 +230,8 @@ int32_t get_num_physical_cores() {
|
||||
return n_threads > 0 ? (n_threads <= 4 ? n_threads : n_threads / 2) : 4;
|
||||
}
|
||||
|
||||
static sd_progress_cb_t sd_progress_cb = NULL;
|
||||
void* sd_progress_cb_data = NULL;
|
||||
static sd_progress_cb_t sd_progress_cb = NULL;
|
||||
void* sd_progress_cb_data = NULL;
|
||||
|
||||
std::u32string utf8_to_utf32(const std::string& utf8_str) {
|
||||
std::wstring_convert<std::codecvt_utf8<char32_t>, char32_t> converter;
|
||||
@@ -207,9 +276,42 @@ std::string path_join(const std::string& p1, const std::string& p2) {
|
||||
return p1 + "/" + p2;
|
||||
}
|
||||
|
||||
sd_image_t* preprocess_id_image(sd_image_t* img) {
|
||||
int shortest_edge = 224;
|
||||
int size = shortest_edge;
|
||||
sd_image_t* resized = NULL;
|
||||
uint32_t w = img->width;
|
||||
uint32_t h = img->height;
|
||||
uint32_t c = img->channel;
|
||||
|
||||
// 1. do resize using stb_resize functions
|
||||
|
||||
unsigned char* buf = (unsigned char*)malloc(sizeof(unsigned char) * 3 * size * size);
|
||||
if (!stbir_resize_uint8(img->data, w, h, 0,
|
||||
buf, size, size, 0,
|
||||
c)) {
|
||||
fprintf(stderr, "%s: resize operation failed \n ", __func__);
|
||||
return resized;
|
||||
}
|
||||
|
||||
// 2. do center crop (likely unnecessary due to step 1)
|
||||
|
||||
// 3. do rescale
|
||||
|
||||
// 4. do normalize
|
||||
|
||||
// 3 and 4 will need to be done in float format.
|
||||
|
||||
resized = new sd_image_t{(uint32_t)shortest_edge,
|
||||
(uint32_t)shortest_edge,
|
||||
3,
|
||||
buf};
|
||||
return resized;
|
||||
}
|
||||
|
||||
void pretty_progress(int step, int steps, float time) {
|
||||
if (sd_progress_cb) {
|
||||
sd_progress_cb(step,steps,time, sd_progress_cb_data);
|
||||
sd_progress_cb(step, steps, time, sd_progress_cb_data);
|
||||
return;
|
||||
}
|
||||
if (step == 0) {
|
||||
@@ -255,9 +357,8 @@ std::string trim(const std::string& s) {
|
||||
return rtrim(ltrim(s));
|
||||
}
|
||||
|
||||
static sd_log_cb_t sd_log_cb = NULL;
|
||||
void* sd_log_cb_data = NULL;
|
||||
|
||||
static sd_log_cb_t sd_log_cb = NULL;
|
||||
void* sd_log_cb_data = NULL;
|
||||
|
||||
#define LOG_BUFFER_SIZE 1024
|
||||
|
||||
@@ -265,18 +366,8 @@ void log_printf(sd_log_level_t level, const char* file, int line, const char* fo
|
||||
va_list args;
|
||||
va_start(args, format);
|
||||
|
||||
const char* level_str = "DEBUG";
|
||||
if (level == SD_LOG_INFO) {
|
||||
level_str = "INFO ";
|
||||
} else if (level == SD_LOG_WARN) {
|
||||
level_str = "WARN ";
|
||||
} else if (level == SD_LOG_ERROR) {
|
||||
level_str = "ERROR";
|
||||
}
|
||||
|
||||
static char log_buffer[LOG_BUFFER_SIZE + 1];
|
||||
|
||||
int written = snprintf(log_buffer, LOG_BUFFER_SIZE, "[%s] %s:%-4d - ", level_str, sd_basename(file).c_str(), line);
|
||||
int written = snprintf(log_buffer, LOG_BUFFER_SIZE, "%s:%-4d - ", sd_basename(file).c_str(), line);
|
||||
|
||||
if (written >= 0 && written < LOG_BUFFER_SIZE) {
|
||||
vsnprintf(log_buffer + written, LOG_BUFFER_SIZE - written, format, args);
|
||||
@@ -471,4 +562,4 @@ sd_image_f32_t clip_preprocess(sd_image_f32_t image, int size) {
|
||||
}
|
||||
|
||||
return result;
|
||||
}
|
||||
}
|
||||
|
||||
8
util.h
@@ -3,11 +3,13 @@
|
||||
|
||||
#include <cstdint>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "stable-diffusion.h"
|
||||
|
||||
bool ends_with(const std::string& str, const std::string& ending);
|
||||
bool starts_with(const std::string& str, const std::string& start);
|
||||
bool contains(const std::string& str, const std::string& substr);
|
||||
|
||||
std::string format(const char* fmt, ...);
|
||||
|
||||
@@ -17,10 +19,16 @@ bool file_exists(const std::string& filename);
|
||||
bool is_directory(const std::string& path);
|
||||
std::string get_full_path(const std::string& dir, const std::string& filename);
|
||||
|
||||
std::vector<std::string> get_files_from_dir(const std::string& dir);
|
||||
|
||||
std::u32string utf8_to_utf32(const std::string& utf8_str);
|
||||
std::string utf32_to_utf8(const std::u32string& utf32_str);
|
||||
std::u32string unicode_value_to_utf32(int unicode_value);
|
||||
|
||||
sd_image_t* preprocess_id_image(sd_image_t* img);
|
||||
|
||||
// std::string sd_basename(const std::string& path);
|
||||
|
||||
typedef struct {
|
||||
uint32_t width;
|
||||
uint32_t height;
|
||||
|
||||
8
vae.hpp
@@ -526,14 +526,6 @@ struct AutoEncoderKL : public GGMLModule {
|
||||
return "vae";
|
||||
}
|
||||
|
||||
size_t get_params_mem_size() {
|
||||
return ae.get_params_mem_size();
|
||||
}
|
||||
|
||||
size_t get_params_num() {
|
||||
return ae.get_params_num();
|
||||
}
|
||||
|
||||
void get_param_tensors(std::map<std::string, struct ggml_tensor*>& tensors, const std::string prefix) {
|
||||
ae.get_param_tensors(tensors, prefix);
|
||||
}
|
||||
|
||||