stable-diffusion : ggml-alloc integration and gpu acceleration

Context
I had been waiting for a response #72 from @leejet regarding updating the GGML version so I could integrate ggml-alloc, the new memory management approach for GGML (also necessary for integrating the CUDA backend). He synchronized it with the latest changes in @ggerganov GGML repository, but it seems that in the process, the GGML library stopped working directly (at least on my computer, it's causing issues when computing) because it still has the dynamic mode integrated.
I'm using this ggml version in this PR.
I have removed the ggml submodule and directly implemented one that I already have, in which I refactored the Conv1D and Conv2D operations to use ggml_im2col and ggml_mul_mat. This will make it easier for me to create the different CUDA kernels.
Build this PR
git clone https://github.com/FSSRepo/stable-diffusion.cpp.git
cd stable-diffusion.cpp
git clone https://github.com/FSSRepo/ggml.git
mkdir build
cd build
cmake ..
cmake --build . --config Release

I've noticed that ggml-alloc consumes much more memory than leejet's current dynamic mode. I'm likely overlooking some configuration to prevent extensive compute graphs from consuming too much memory with ggml-alloc.
Here are the comparative details and results of the integration with ggml-alloc
./sd  -m kotosmix_q4_0.bin --verbose -p "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" -n "bad quality, ugly" --steps 10 -t 6 -s 48747

Note: the stable-diffusion.cpp commit afec505 that worked for me.
Memory usage:

compute graph
dynamic mode (runtime memory)
ggml-alloc (compute buffer size)

learned condition (CLIP Text Encoder)
13.37MB
25.54 MB

diffusion (UNet Model)
570.65MB
1147.33 MB

AutoEncoderKL (VAE Decode)
1674.14MB
2496.00 MB

Timings:

compute graph
dynamic mode
ggml-alloc

learned condition (CLIP Text Encoder)
0.26 s
0.23 s

diffusion (UNet Model)
~ 18.22 s per sample
~ 13.81 s per sample

AutoEncoderKL (VAE Decode)
27 s
24.22 s

Result image (the same in both cases):

Over the next few days, I will be making some improvements, trying to investigate and gather feedback with ggml-alloc to avoid the current memory usage, and once I have fixed that, I can start working on implementing the CUDA kernels.
Progress overview
Here's a summary of what I will be carrying out in the coming days to complete this PR:
ggml operations that utilize stable diffusion and their availability across backends.:

operation
cpu backend
CUDA backend

ADD
✅
✅

MUL
✅
✅

REPEAT
✅
✅

CONCAT
✅
❌ (need implement)

NORM
✅
✅

GET_ROWS
✅
✅

DIAG_MASK_INF
✅
✅

GROUP_NORM
✅
❌ (need implement)

MUL_MAT
✅
✅

SCALE
✅
✅

PERMUTE
✅
-

RESHAPE
✅
-

VIEW
✅
-

SOFT_MAX
✅
✅

IM2COL
✅
✅

UPSCALE
✅
❌ (need implement)

UNARY
✅
❌ (only GeLU and SiLU, missing GeLU Quick)

Tasks for complete this PR:

[x] Integrate ggml-alloc to simplify the use of ggml-backend.
[ ] Improve memory usage, determine if I'm using the ggml-alloc API correctly, I might be missing something that could help me reduce unnecessary memory usage during compute. can you help me? @slaren @ggerganov.
[ ] Implement the missing CUDA kernels ggml_concat, ggml_group_norm, ggml_upscale, ggml_gelu_quick.
[ ] Offload the model to the GPU and conduct performance and memory usage tests.
[ ] Implement the use of custom VAEs. I will try to implement the gguf format to simplify the loading of tensors.

I've noticed that ggml-alloc consumes much more memory than leejet's current dynamic mode. I'm likely overlooking some configuration to prevent extensive compute graphs from consuming too much memory with ggml-alloc.

Nothing comes directly to mind. Need to have a more detailed look in the code when I get some time.
One possible error is to avoid reusing the same compute buffer for new graphs - I had some troubles with that when adding ggml-alloc to whisper.cpp: https://github.com/ggerganov/whisper.cpp/pull/1270#discussion_r1320847591
It does not seem very related the increased memory that you observed though.

Great job! I'll find some time to review and merge it.

@FSSRepo thanks.
I need to comment out line 245-257 in stable-diffusion.cpp, then it works for me:
    //if(allocr != NULL) {
        //ggml_allocr_alloc(allocr, embedding);
    //}

Otherwise I get: GGML_ASSERT: ggml-alloc.c:116: tensor->data == 0
One step takes about 11s with 512x512, which is faster than before.

Otherwise I get: GGML_ASSERT: ggml-alloc.c:116: tensor->data == 0

I will try to reproduce this, because it's a strange behaivor

Can you provide more information, system operative, hardware

Of course: Windows 11 / VS / CUDA / GTX 3090 / i7-12700k

I need to comment out line 245-257 in stable-diffusion.cpp

    //if(allocr != NULL) {
        //ggml_allocr_alloc(allocr, embedding);
    //}

245 to 257 are 12 lines or you refer 245 to 247?

you refer 245 to 247?

Yes, my mistake...

With /D GGML_USE_CUBLAS one step 512x512 is between 8 and 9 seconds.

With /D GGML_USE_CUBLAS one step 512x512 is between 8 and 9 seconds.

The backend to compute is fixed to cpu_backend

The backend to compute is fixed to cpu_backend

I see. Interestingly it is already a little faster (8~ vs 10~ seconds/step).

The backend to compute is fixed to cpu_backend

I see. Interestingly it is already a little faster (8~ vs 10~ seconds/step).

Possibly matrix multiplications are being delegated to the GPU.

I need to comment out line 245-247 in stable-diffusion.cpp, then it works for me:

Are you using stable diffusion 2 when this error happen?

With the main branch I get this error: GGML_ASSERT: ggml.c: ctx->mem_buffer != NULL imartinez/privateGPT#90 So maybe it is not enough space on my HDD...

What matter that in this context?

@FSSRepo yes, it was with stable diffusion 2. I had that error GGML_ASSERT: ggml.c: ctx->mem_buffer != NULL with the main branch, so I thought maybe this one is caused by missing disc space too (but no it is not, so it does not matter in this context).

Instead of calling ggml_allocr_alloc after every tensor, you can do something like this to allocate all the tensors created in a context:
for (struct ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
        ggml_allocr_alloc(alloc, t);
    }

You can also use a similar loop calling ggml_nbytes to calculate the size of the buffer.
Caveat: the functions to enumerate the tensors in a context are only in the llama.cpp repository currently, they still need to be synced to the ggml repo.
Related: https://github.com/ggerganov/ggml/issues/578

Why did you close this?

To resolve merge conflicts, I rebased the master branch of FSSRepo/stable-diffusion.cpp onto the latest master branch of leejet/stable-diffusion.cpp. I also updated the URL for GGML to https://github.com/FSSRepo/ggml/tree/master to make it easier to review the changes and facilitate future merges.

It seems like I mistakenly closed this pull request.

It seems like I mistakenly closed this pull request.

I am unable to push my changes to https://github.com/FSSRepo/stable-diffusion.cpp/tree/master at the moment, which prevents me from reopening this pull request. Can you please pull https://github.com/leejet/stable-diffusion.cpp/tree/pr75 and push the corresponding changes to https://github.com/FSSRepo/stable-diffusion.cpp/tree/master? This way, I can reopen this pull request.

Try now

Sorry for the inconvenience

Sorry for the inconvenience

Don't worry. Anyway, I will need your help to understand how some aspects of this stable diffusion implementation work.

I still can't reopen this PR. It says 'there are now new commits.' It might require you to push the new changes to the master.

I still can't reopen this PR. It says 'there are now new commits.' It might require you to push the new changes to the master.

I will try it

I've updated the submodule 'ggml' in this branch https://github.com/leejet/stable-diffusion.cpp/tree/pr75 to the latest commit of FSSRepo/sd-cpp-fixes, and this branch also includes the changes you made to stable-diffusion.cpp earlier.

You can pull this branch to your local machine, rebase your local stable-diffusion.cpp master branch to this branch, and then push the changes to https://github.com/FSSRepo/stable-diffusion.cpp/tree/master.

@leejet can you reopen now?

reopen

@leejet I'm trying to get the kernel to work for the 'group_norm' operation, but I haven't been able to perform the sum in parallel, mainly due to the limit of concurrent threads and the size of the tensors passing through this operation (very large, and this operation in its CPU implementation seems to be not very parallelizable). I want to understand the difference between ggml_norm and ggml_group_norm and thus arrive at a clearer approach.

I want to understand the difference between ggml_norm and ggml_group_norm and thus arrive at a clearer approach.

I think this might help
https://gaoxiangluo.github.io/2021/08/01/Group-Norm-Batch-Norm-Instance-Norm-which-is-better/figure6.png

I want to understand the difference between ggml_norm and ggml_group_norm and thus arrive at a clearer approach.

I think this might help
https://gaoxiangluo.github.io/2021/08/01/Group-Norm-Batch-Norm-Instance-Norm-which-is-better/figure6.png

Seems good. While I was showering, the solution came to me, hahaha.

group norm kernel finished 😎
Tensor 3D [4, 4, 8] and n_group = 4 (left: CUDA, right: ggml cpu backend):

@slaren Is there a function to transfer GPU data to CPU without having to create this function? tensor->data in this case is a pointer of memory in gpu.
void ggml_fallback_tensor_f32(struct ggml_tensor* tensor) {
    if(ggml_backend_is_cpu(ggml_get_backend(tensor))) {
        return;
    }
    // for cuda backend bring data from device
    assert(tensor->type == GGML_TYPE_F32);
    int size_ = ggml_nbytes(tensor);
    float* data = (float*)malloc(size_);
    ggml_backend_tensor_get(tensor, data, 0, size_);
    tensor->data = data;
}

You can use ggml_backend_tensor_copy to copy a tensor between backends. Allocate a tensor on the CPU backend with the exact same shape and call ggml_backend_tensor_copy(gpu_tensor, cpu_tensor). Is that what you need? You cannot just change the data pointer of a tensor.

@leejet I've noticed that this is very sensitive to very small changes. Previously, the 'eps' parameter of 'ggml norm' was set to 1e-5. When you updated ggml, and you had to manually set that parameter, you changed it to 1e-6. This resulted in the appearance of small artifacts and variations in the final image, especially in the CLIP model.

Different epsilon (eps) values do indeed lead to different results. The default epsilon value for PyTorch's LayerNorm is 1e-5, which is also the value used in Stable-Diffusion. For GroupNorm, Stable-Diffusion uses an epsilon value of 1e-6. I made a mistake in the update earlier, and to align with Stable-Diffusion, I have just reverted the epsilon value of LayerNorm to 1e-5.
I recommend that you rebase my latest master branch to incorporate this change.
You can find more information about this in the PyTorch documentation here: PyTorch LayerNorm.

I have been conducting tests between data types to understand the outcomes across different backends and ensure that I am not doing something wrong.
I have used the same model, prompt, negative prompt, seed, and number of steps.
sd  -m $model --verbose -p "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" -n "bad quality, ugly" --steps 10 -t 6 -s 48747

Model q4_0 Full CPU Backend (clip text model, unet model, vae decoder)
When GGML_CUBLAS=ON the matrix multiplications are performed on the GPU.

Commit
GGML_CUBLAS=OFF
GGML_CUBLAS=ON

afec505

7f72088

Model f16 Full CPU Backend (clip text model, unet model, vae decoder)
When GGML_CUBLAS=ON the matrix multiplications are performed on the GPU.

Commit
GGML_CUBLAS=OFF
GGML_CUBLAS=ON

afec505

7f72088

CLIP Text Model full offloaded to GPU, UNet and VAE still with CPU Backend
The matrix multiplications are performed on the GPU.

q4_0
f16

Researching the reason or operations that diverge between the CPU and GPU, it's only when a matrix multiplication ggml_mul_mat is performed that the results change.

@leejet I'm noticing that the master version of stable-diffusion is consuming the same amount as my ggml-alloc implementation, despite using your dynamic mode, which before updating ggml, used half as much.

This doesn't look good at all; the ggml_repeat operation on the GPU takes too long and is performed multiple times during the diffusion computation.
It takes almost 2 minutes to reach an unimplemented operation, which makes me reconsider whether it's worth implementing GPU acceleration given the current limitations of ggml. Additionally, it consumes a lot of memory just for the UNet parameters and the computation in F16.

UNet with CPU Backend
sd  -m kotosmix_f16.bin --verbose -p "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" -n "bad quality, ugly" --steps 10 -t 6 -s 48747
Option:
    n_threads:       6
    mode:            txt2img
    model_path:      kotosmix_f16.bin
    output_path:     output.png
    init_img:
    prompt:          beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality
    negative_prompt: bad quality, ugly
    cfg_scale:       7.00
    width:           512
    height:          512
    sample_method:   euler_a
    schedule:        default
    sample_steps:    10
    strength:        0.75
    rng:             cuda
    seed:            48747
System Info:
    BLAS = 0
    SSE3 = 1
    AVX = 1
    AVX2 = 1
    AVX512 = 0
    AVX512_VBMI = 0
    AVX512_VNNI = 0
    FMA = 1
    NEON = 0
    ARM_FMA = 0
    F16C = 1
    FP16_VA = 0
    WASM_SIMD = 0
    VSX = 0
[INFO]  stable-diffusion.cpp:3362 - loading model from 'kotosmix_f16.bin'
[DEBUG] stable-diffusion.cpp:3372 - verifying magic
[DEBUG] stable-diffusion.cpp:3390 - loading hparams
[INFO]  stable-diffusion.cpp:3405 - model type: SD1.x
[INFO]  stable-diffusion.cpp:3417 - ftype: f16
[DEBUG] stable-diffusion.cpp:3422 - loading vocab
[DEBUG] stable-diffusion.cpp:3453 - ggml tensor size = 320 bytes
[DEBUG] stable-diffusion.cpp:789  - clip params backend buffer size =  236.18 MB
[DEBUG] stable-diffusion.cpp:792  - clip tensor count = 437
[DEBUG] stable-diffusion.cpp:1954 - Using cpu backend
[DEBUG] stable-diffusion.cpp:1961 - unet params backend buffer size =  1641.16 MB
[DEBUG] stable-diffusion.cpp:1965 - unet tensor count = 688
[DEBUG] stable-diffusion.cpp:3058 - vae params backend buffer size =  95.47 MB
[DEBUG] stable-diffusion.cpp:3067 - vae tensor count = 164
[DEBUG] stable-diffusion.cpp:3464 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:861  - clip params allocated
[DEBUG] stable-diffusion.cpp:2050 - unet params allocated
[DEBUG] stable-diffusion.cpp:3109 - vae params allocated
[DEBUG] stable-diffusion.cpp:3567 - loading weights
[DEBUG] stable-diffusion.cpp:3685 - model size = 1969.67MB
[INFO]  stable-diffusion.cpp:3696 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3698 - loading model from 'kotosmix_f16.bin' completed, taking 1.44s
[INFO]  stable-diffusion.cpp:3723 - running in eps-prediction mode
[DEBUG] stable-diffusion.cpp:3750 - finished loaded file
[DEBUG] stable-diffusion.cpp:1018 - parse 'beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality' to [['beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality', 1], ]
[DEBUG] stable-diffusion.cpp:429  - split prompt "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" to tokens ["beautiful</w>", "anime</w>", "girl</w>", ",</w>", "white</w>", "hair</w>", ",</w>", "blue</w>", "eyes</w>", ",</w>", "realistic</w>", ",</w>", "4</w>", "k</w>", ",</w>", "high</w>", "quality</w>", ]
begin: learned condition compute buffer size: 25.54 MB
start compute
[DEBUG] stable-diffusion.cpp:3804 - computing condition graph completed, taking 135 ms
[DEBUG] stable-diffusion.cpp:1018 - parse 'bad quality, ugly' to [['bad quality, ugly', 1], ]
[DEBUG] stable-diffusion.cpp:429  - split prompt "bad quality, ugly" to tokens ["bad</w>", "quality</w>", ",</w>", "ugly</w>", ]
start compute
[DEBUG] stable-diffusion.cpp:3804 - computing condition graph completed, taking 136 ms
[INFO]  stable-diffusion.cpp:4450 - get_learned_condition completed, taking 277 ms
[INFO]  stable-diffusion.cpp:4464 - start sampling
begin: diffusion compute buffer size: 1147.33 MB
[INFO]  stable-diffusion.cpp:3963 - sampling using Euler A method
start compute
time from start compute to reach group norm: 334 ms

UNet with GPU Backend
sd  -m kotosmix_f16.bin --verbose -p "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" -n "bad quality, ugly" --steps 10 -t 6 -s 48747
Option:
    n_threads:       6
    mode:            txt2img
    model_path:      kotosmix_f16.bin
    output_path:     output.png
    init_img:
    prompt:          beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality
    negative_prompt: bad quality, ugly
    cfg_scale:       7.00
    width:           512
    height:          512
    sample_method:   euler_a
    schedule:        default
    sample_steps:    10
    strength:        0.75
    rng:             cuda
    seed:            48747
System Info:
    BLAS = 1
    SSE3 = 1
    AVX = 1
    AVX2 = 1
    AVX512 = 0
    AVX512_VBMI = 0
    AVX512_VNNI = 0
    FMA = 1
    NEON = 0
    ARM_FMA = 0
    F16C = 1
    FP16_VA = 0
    WASM_SIMD = 0
    VSX = 0
[INFO]  stable-diffusion.cpp:3362 - loading model from 'kotosmix_f16.bin'
[DEBUG] stable-diffusion.cpp:3372 - verifying magic
[DEBUG] stable-diffusion.cpp:3390 - loading hparams
[INFO]  stable-diffusion.cpp:3405 - model type: SD1.x
[INFO]  stable-diffusion.cpp:3417 - ftype: f16
[DEBUG] stable-diffusion.cpp:3422 - loading vocab
[DEBUG] stable-diffusion.cpp:3453 - ggml tensor size = 320 bytes
[DEBUG] stable-diffusion.cpp:789  - clip params backend buffer size =  236.18 MB
[DEBUG] stable-diffusion.cpp:792  - clip tensor count = 437
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3050 Laptop GPU, compute capability 8.6
[DEBUG] stable-diffusion.cpp:1951 - Using cuda backend
[DEBUG] stable-diffusion.cpp:1961 - unet params backend buffer size =  1641.16 MB
[DEBUG] stable-diffusion.cpp:1965 - unet tensor count = 688
[DEBUG] stable-diffusion.cpp:3058 - vae params backend buffer size =  95.47 MB
[DEBUG] stable-diffusion.cpp:3067 - vae tensor count = 164
[DEBUG] stable-diffusion.cpp:3464 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:861  - clip params allocated
[DEBUG] stable-diffusion.cpp:2050 - unet params allocated
[DEBUG] stable-diffusion.cpp:3109 - vae params allocated
[DEBUG] stable-diffusion.cpp:3567 - loading weights
[DEBUG] stable-diffusion.cpp:3685 - model size = 1969.67MB
[INFO]  stable-diffusion.cpp:3696 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3698 - loading model from 'kotosmix_f16.bin' completed, taking 1.40s
[INFO]  stable-diffusion.cpp:3723 - running in eps-prediction mode
[DEBUG] stable-diffusion.cpp:3750 - finished loaded file
[DEBUG] stable-diffusion.cpp:1018 - parse 'beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality' to [['beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality', 1], ]
[DEBUG] stable-diffusion.cpp:429  - split prompt "beautiful anime girl, white hair, blue eyes, realistic, 4k, high quality" to tokens ["beautiful</w>", "anime</w>", "girl</w>", ",</w>", "white</w>", "hair</w>", ",</w>", "blue</w>", "eyes</w>", ",</w>", "realistic</w>", ",</w>", "4</w>", "k</w>", ",</w>", "high</w>", "quality</w>", ]
begin: learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3804 - computing condition graph completed, taking 113 ms
[DEBUG] stable-diffusion.cpp:1018 - parse 'bad quality, ugly' to [['bad quality, ugly', 1], ]
[DEBUG] stable-diffusion.cpp:429  - split prompt "bad quality, ugly" to tokens ["bad</w>", "quality</w>", ",</w>", "ugly</w>", ]
[DEBUG] stable-diffusion.cpp:3804 - computing condition graph completed, taking 87 ms
[INFO]  stable-diffusion.cpp:4450 - get_learned_condition completed, taking 206 ms
[INFO]  stable-diffusion.cpp:4464 - start sampling
begin: diffusion compute buffer size: 1147.53 MB
[INFO]  stable-diffusion.cpp:3963 - sampling using Euler A method
start cuda compute
time from start compute to reach group norm: 284432 ms
ggml_backend_cuda_graph_compute: error: op not supported node_450 (GROUP_NORM)

@slaren Did you already know that this ggml_repeat operation in the CUDA backend was inefficient? I'll have to implement a kernel for this operation and see if I can improve these times.
static void ggml_cuda_op_repeat(
    const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
    const float * src0_d, const float * src1_d, float * dst_d, const cudaStream_t & stream) {
    // guaranteed to be an integer due to the check in ggml_can_repeat
    const int64_t ne0 = dst->ne[0];
    const int64_t ne1 = dst->ne[1];
    const int64_t ne2 = dst->ne[2];
    const int64_t ne3 = dst->ne[3];

    const int64_t ne00 = src0->ne[0];
    const int64_t ne01 = src0->ne[1];
    const int64_t ne02 = src0->ne[2];
    const int64_t ne03 = src0->ne[3];

    const size_t nb0 = dst->nb[0];
    const size_t nb1 = dst->nb[1];
    const size_t nb2 = dst->nb[2];
    const size_t nb3 = dst->nb[3];

    const size_t nb00 = src0->nb[0];
    const size_t nb01 = src0->nb[1];
    const size_t nb02 = src0->nb[2];
    const size_t nb03 = src0->nb[3];

    const int nr0 = (int)(ne0/ne00);
    const int nr1 = (int)(ne1/ne01);
    const int nr2 = (int)(ne2/ne02);
    const int nr3 = (int)(ne3/ne03);

    // TODO: support for transposed / permuted tensors
    GGML_ASSERT(nb0  == sizeof(float));
    GGML_ASSERT(nb00 == sizeof(float));

    // TODO: very inefficient, implement in a kernel, or fewer cudaMemcpyAsync calls for contiguous tensors
    for                         (int i3 = 0; i3 < nr3;  i3++) {
        for                     (int k3 = 0; k3 < ne03; k3++) {
            for                 (int i2 = 0; i2 < nr2;  i2++) {
                for             (int k2 = 0; k2 < ne02; k2++) {
                    for         (int i1 = 0; i1 < nr1;  i1++) {
                        for     (int k1 = 0; k1 < ne01; k1++) {
                            for (int i0 = 0; i0 < nr0;  i0++) {
                                CUDA_CHECK(cudaMemcpyAsync(
                                              (char *)  dst_d + (i3*ne03 + k3)*nb3  + (i2*ne02 + k2)*nb2  + (i1*ne01 + k1)*nb1  + (i0*ne00)*nb0,
                                        (const char *) src0_d + (          k3)*nb03 + (          k2)*nb02 + (          k1)*nb01,
                                        ne00*nb0, cudaMemcpyDeviceToDevice, stream));
                            }
                        }
                    }
                }
            }
        }
    }

    (void) src1;
    (void) src1_d;
}

Yes, I wrote that comment. I added this implementation to test the gpt-2 example with ggml-backend, but for any real world usage it needs to be optimized. Ideally, repeat ops should be removed and instead broadcasting should be implemented in the ops that need it, but I don't know if that's possible for stable diffusion.

instead broadcasting should be implemented in the ops that need it

I don't understand what you mean.

It means that the dimension of the tensor should be automatically extended when it is required to perform some operation. The operation should behave in a similar way as if repeat was used in these cases. This is partially support is some ggml ops already, add, mul and mul mat, I think. Check this: https://numpy.org/doc/stable/user/basics.broadcasting.html

@leejet I've already implemented the CUDA backend for the UNet model. Currently, it's encountering issues as errors in some operations accumulate during the sampling process. I'll be monitoring the results of each diffusion model layer and tracking down the functions whose outcomes diverge from the CPU backend. I'll also be testing the kernels I've implemented to rule out any potential flaws in my implementation.
The performance gained is significant (1 sample per second) ; I hope to be able to address the issues and obtain at least one image generated purely by the CUDA backend.

@slaren Translation to English: Is there a way to use the ggml_backend to forcibly redirect the computation of certain operations to the CPU for debugging purposes?

Not currently. Using multiple backends simultaneously will be supported in v2.

I'm noticing that the master version of stable-diffusion is consuming the same amount as my ggml-alloc implementation, despite using your dynamic mode, which before updating ggml, used half as much.

@FSSRepo Because there have been some changes to the computation graph of ggml, it was causing issues with memory not being released in a timely manner in dynamic mode, leading to increased memory usage. I just fixed this issue. You can pull the latest master code and give it a try.

@slaren I can't find the reason why matrix multiplications differ even though they are supposed to be the same. For instance, if I use the CPU backend with GGML_CUBLAS=ON, the ggml_mult_mat operation is performed on the GPU. However, now that I'm using the CUDA backend, I'm getting different results for the same input. I feel completely lost, and what's worse, I can't test my assumptions because I can't skip having the multmat operation performed on the CPU to obtain identical results.
Partial input (same for two backends)
matrix A [320, 1280, 1, 1] - FP16
  [0, 0, 0, 0] = 34998
  [0, 0, 0, 1] = 5801
  [0, 0, 0, 2] = 4793
  [0, 0, 0, 317] = 42215
  [0, 0, 0, 318] = 9040
  [0, 0, 0, 319] = 41546
  [0, 0, 1, 0] = 38677
  [0, 0, 1, 1] = 2040
  [0, 0, 1, 2] = 39386
  [0, 0, 1, 317] = 41281
  [0, 0, 1, 318] = 8070
  [0, 0, 1, 319] = 39651
  [0, 0, 2, 0] = 36739
  [0, 0, 2, 1] = 7243
  [0, 0, 2, 2] = 6634
  [0, 0, 2, 317] = 42108
  [0, 0, 2, 318] = 40134
  [0, 0, 2, 319] = 6881
  [0, 0, 1277, 0] = 3191
  [0, 0, 1277, 1] = 38774
  [0, 0, 1277, 2] = 38211
  [0, 0, 1277, 317] = 41590
  [0, 0, 1277, 318] = 8730
  [0, 0, 1277, 319] = 8400
  [0, 0, 1278, 0] = 39801
  [0, 0, 1278, 1] = 39710
  [0, 0, 1278, 2] = 5893
  [0, 0, 1278, 317] = 8530
  [0, 0, 1278, 318] = 5669
  [0, 0, 1278, 319] = 7057
  [0, 0, 1279, 0] = 1513
  [0, 0, 1279, 1] = 39951
  [0, 0, 1279, 2] = 39174
  [0, 0, 1279, 317] = 8392
  [0, 0, 1279, 318] = 8368
  [0, 0, 1279, 319] = 9412

matrix B [320, 1, 1, 1] - FP32
  [0, 0, 0, 0] = 0.999650
  [0, 0, 0, 1] = 0.802678
  [0, 0, 0, 2] = -0.278121
  [0, 0, 0, 317] = 0.118453
  [0, 0, 0, 318] = 0.111855
  [0, 0, 0, 319] = 0.105622

cpu backend
With GGML_CUBLAS=ON, the ggml_mult_mat operation is performed on the GPU, then I expect exactly the same output
mult emb [1280, 1, 1, 1] - MUL_MAT - FP32
  [0, 0, 0, 0] = -0.060595
  [0, 0, 0, 1] = -0.043580
  [0, 0, 0, 2] = -0.003741
  [0, 0, 0, 1277] = -0.099713
  [0, 0, 0, 1278] = 0.014790
  [0, 0, 0, 1279] = -0.005692

gpu backend
mult emb [1280, 1, 1, 1] - MUL_MAT - FP32
  [0, 0, 0, 0] = -0.060599
  [0, 0, 0, 1] = -0.043596
  [0, 0, 0, 2] = -0.003743
  [0, 0, 0, 1277] = -0.099860
  [0, 0, 0, 1278] = 0.014806
  [0, 0, 0, 1279] = -0.005704

That matrix multiplication wouldn't be performed on the GPU just by building with GGML_CUBLAS because it is too small. The results seem close enough that I would assume that the matrix multiplication is correct.

That matrix multiplication wouldn't be performed on the GPU just by building with GGML_CUBLAS because it is too small. The results seem close enough that I would assume that the matrix multiplication is correct.

The problem is that this model is too sensitive to variations.

You can try using https://github.com/ggerganov/ggml/pull/586 to run some operations on the CPU backend, but there may be issues.

You can try using ggerganov/ggml#586 to run some operations on the CPU backend, but there may be issues.

How could I make sure that only the 'ggml_mult_mat' operations are performed on the CPU with that backend version?

You can use ggml_backend_sched_set_node_backend for force an operation to run on a specific backend. You can use it to set the matrix multiplications on the CPU. But that may also cause the operations around it to also use the CPU back, so you may have to use ggml_backend_sched_set_node_backend to force them back on CUDA. But there may be all kind of issues doing this, this is only tested on the gpt-2 example, so don't try this it you are not willing to look into the internals.

You can use ggml_backend_sched_set_node_backend for force an operation to run on a specific backend. You can use it to set the matrix multiplications on the CPU. But that may also cause the operations around it to also use the CPU back, so you may have to use ggml_backend_sched_set_node_backend to force them back on CUDA. But there may be all kind of issues doing this, this is only tested on the gpt-2 example, so don't try this it you are not willing to look into the internals.

I feel a bit lazy about having to analyze how to implement the backend v2 in sd.cpp, so I'm going to create the CUDA kernels ggml_mul_mat  that produce identical results to the CPU backend.

@slaren The results of ggml_add do not match between backends for three-dimensional tensors.

Yeah I am not surprised, the CUDA backend was pretty much only tested with llama.cpp. I think that the ggml_add implementation only considers the first 2 dimensions.

For some reason the results of ggml_group_norm differs between backends:
left: cpu backend, right: cuda backend

The calculation of the normalization for each group is performed in a single thread, similar to how it works in the CPU backend. I had tried to make the summation process parallel by using shared memory to share the sum among threads, but it only worked for small and precise tensors (with a few decimal places). When dealing with very large tensors and those with many decimal places, it started to produce random values in the decimals between each execution. I always made sure to synchronize the threads and perform the summation using atomicAdd, but it consistently encountered issues.
In short, the kernel group norm is highly inefficient when it comes to utilizing GPU resources for large tensors with shapes like [64, 64, 320].

Not currently. Using multiple backends simultaneously will be supported in v2.

v2 merged 3 days ago

v2 merged 3 days ago. easier to debug specific ops now with that version.

I've been reviewing how the new backend works, and the truth is that adapting it to stable diffusion will be too cumbersome for me because I need to change the way tensors are loaded, as they must be loaded into their respective backend to perform certain operations in a specific backend. It's too time-consuming, and I'd prefer to do it properly when I adapt the gguf format to load the models. Debugging stable diffusion is proving to be quite challenging for me. To be honest, I'm feeling a bit discouraged.

@slaren I was wondering if tensors need to be definitely in the backend's buffer where I want to perform the operation. I was thinking of keeping the tensors in GPU memory, but when I use ggml_backend_sched_set_node_backend to perform a certain operation on the CPU, the tensors will be transferred to RAM and then return the result to the GPU.

@slaren In the end, I went through the reverse process; I switched to a CPU backend and forced the operations I wanted to test to run on the GPU. It turned out to be much easier, and I never thought of it from the beginning. I was overcomplicating things, but now I'm achieving multiple goals with a single effort.
bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor) {
    ggml_cuda_func_t func;
    const bool any_on_device = tensor->backend == GGML_BACKEND_GPU
        || (tensor->src[0] != nullptr && (tensor->src[0]->backend == GGML_BACKEND_GPU || tensor->src[0]->backend == GGML_BACKEND_GPU_SPLIT))
        || (tensor->src[1] != nullptr && tensor->src[1]->backend == GGML_BACKEND_GPU);

    if (!any_on_device &&
        // ops to test in cuda
        tensor->op != GGML_OP_UNARY &&
        tensor->op != GGML_OP_MUL_MAT &&
        tensor->op != GGML_OP_NORM &&
        tensor->op != GGML_OP_GROUP_NORM
    ) {
        return false;
    }

    if(tensor->op == GGML_OP_UNARY && ggml_get_unary_op(tensor) != GGML_UNARY_OP_GELU) {
        return false;
    }

    switch (tensor->op) {

@Jonathhhan Could you try the latest commit on your GPU?
git clone https://github.com/FSSRepo/stable-diffusion.cpp.git --recursive
cd stable-diffusion.cpp
mkdir build
cd build
cmake .. -DSD_CUBLAS=ON
cmake --build . --config Release

And you start generating an image, all the operations are performed on the GPU, but it may seem a bit slow since the parameters will be in RAM memory since the backend used is the CPU backend, but the computation is directed to the GPU.
When enabling the CUDA backend, I am getting NaN (Not-a-Number) results in the tensor outputs. If all operations are being executed correctly on the GPU, it might not be an issue with any kernel errors. This could be due to a memory management issue in the CUDA backend.

@slaren It seems there might be a memory handling issue in the CUDA backend, as when I switch to the CPU backend and send all operations to the GPU, I get correct results. So, the operations themselves don't seem to be the problem; the only differing factor in this case is the memory management.

@leejet Can you give me permission to convert the classes into structs to add support as a dynamic library and also make it easier to use in CLI, server, and other potential uses?

I only have CUDA support with Windows at the moment, so I tried with VS. And I get this error:
ggml_allocr_alloc: not enough space in the buffer (needed 20971520, largest block available 212032)
GGML_ASSERT: ggml-alloc.c:148: !"not enough space in the buffer"

I only have CUDA support with Windows at the moment, so I tried with VS. And I get this error:
ggml_allocr_alloc: not enough space in the buffer (needed 20971520, largest block available 212032)
GGML_ASSERT: ggml-alloc.c:148: !"not enough space in the buffer"

Send me full console log

@FSSRepo I tried to models but both with the same result:
stable-diffusion-nano-2-1-ggml-model-f16.bin
v2-1_512-ema-pruned-ggml-model-f16.bin
That's the full log:
System Info:     BLAS = 1     SSE3 = 1     AVX = 1     AVX2 = 1     AVX512 = 0     AVX512_VBMI = 0     AVX512_VNNI = 0     FMA = 1     NEON = 0     ARM_FMA = 0     F16C = 1     FP16_VA = 0     WASM_SIMD = 0     VSX = 0 [INFO]  stable-diffusion.cpp:3368 - loading model from 'data/models/v2-1_512-ema-pruned-ggml-model-f16.bin' [DEBUG] stable-diffusion.cpp:3378 - verifying magic [DEBUG] stable-diffusion.cpp:3396 - loading hparams [INFO]  stable-diffusion.cpp:3411 - model type: SD2.x [INFO]  stable-diffusion.cpp:3423 - ftype: f16 [DEBUG] stable-diffusion.cpp:3428 - loading vocab [DEBUG] stable-diffusion.cpp:3459 - ggml tensor size = 320 bytes [DEBUG] stable-diffusion.cpp:797  - Using CPU backend - clip [DEBUG] stable-diffusion.cpp:804  - clip params backend buffer size =  675.18 MB [DEBUG] stable-diffusion.cpp:807  - clip tensor count = 869 ggml_init_cublas: found 1 CUDA devices:   Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6 [DEBUG] stable-diffusion.cpp:1969 - Using CPU backend - unet [DEBUG] stable-diffusion.cpp:1977 - unet params backend buffer size =  1653.34 MB [DEBUG] stable-diffusion.cpp:1981 - unet tensor count = 688 [DEBUG] stable-diffusion.cpp:3070 - vae params backend buffer size =  179.03 MB [DEBUG] stable-diffusion.cpp:3079 - vae tensor count = 328 [DEBUG] stable-diffusion.cpp:3470 - preparing memory for the weights [DEBUG] stable-diffusion.cpp:876  - clip params allocated [DEBUG] stable-diffusion.cpp:2066 - unet params allocated [DEBUG] stable-diffusion.cpp:3121 - vae params allocated [DEBUG] stable-diffusion.cpp:3570 - loading weights [DEBUG] stable-diffusion.cpp:3685 - model size = 2486.00MB [INFO]  stable-diffusion.cpp:3696 - total memory buffer size = 2507.55MB (clip 675.18MB, unet 1653.34MB, vae 179.03MB) [INFO]  stable-diffusion.cpp:3698 - loading model from 'data/models/v2-1_512-ema-pruned-ggml-model-f16.bin' completed, taking 1.51s begin: diffusion compute buffer size: 10.20 MB [DEBUG] stable-diffusion.cpp:3790 - check is_using_v_parameterization_for_sd2, taking 0.91s [INFO]  stable-diffusion.cpp:3723 - running in eps-prediction mode [DEBUG] stable-diffusion.cpp:3750 - finished loaded file [DEBUG] stable-diffusion.cpp:1029 - parse 'a lovely cat' to [['a lovely cat', 1], ] [DEBUG] stable-diffusion.cpp:443  - split prompt "a lovely cat" to tokens ["a</w>", "lovely</w>", "cat</w>", ] begin: learned condition compute buffer size: 63.83 MB [DEBUG] stable-diffusion.cpp:3809 - computing condition graph completed, taking 463 ms [DEBUG] stable-diffusion.cpp:1029 - parse '' to [['', 1], ] [DEBUG] stable-diffusion.cpp:443  - split prompt "" to tokens [] [DEBUG] stable-diffusion.cpp:3809 - computing condition graph completed, taking 295 ms [INFO]  stable-diffusion.cpp:4459 - get_learned_condition completed, taking 767 ms [INFO]  stable-diffusion.cpp:4473 - start sampling [INFO]  stable-diffusion.cpp:3969 - sampling using Euler A method ggml_allocr_alloc: not enough space in the buffer (needed 20971520, largest block available 212032) GGML_ASSERT: C:\Users\Jonat\Desktop\of_v20231019_vs_release\examples\eigen\ofStableDiffusionExample\src\stableDiffusion\ggml\ggml-alloc.c:148: !"not enough space in the buffer"

It does work with v1-5-pruned-emaonly-ggml-model-f32.bin.

@FSSRepo I tried to models but both with the same result:
stable-diffusion-nano-2-1-ggml-model-f16.bin
v2-1_512-ema-pruned-ggml-model-f16.bin
That's the full log:

System Info:
BLAS = 1
SSE3 = 1
AVX = 1
AVX2 = 1
AVX512 = 0
AVX512_VBMI = 0
AVX512_VNNI = 0
FMA = 1
NEON = 0
ARM_FMA = 0
F16C = 1
FP16_VA = 0
WASM_SIMD = 0
VSX = 0
[INFO]  stable-diffusion.cpp:3368 - loading model from 'data/models/v2-1_512-ema-pruned-ggml-model-f16.bin'
[DEBUG] stable-diffusion.cpp:3378 - verifying magic
[DEBUG] stable-diffusion.cpp:3396 - loading hparams
[INFO]  stable-diffusion.cpp:3411 - model type: SD2.x
[INFO]  stable-diffusion.cpp:3423 - ftype: f16
[DEBUG] stable-diffusion.cpp:3428 - loading vocab
[DEBUG] stable-diffusion.cpp:3459 - ggml tensor size = 320 bytes
[DEBUG] stable-diffusion.cpp:797  - Using CPU backend - clip
[DEBUG] stable-diffusion.cpp:804  - clip params backend buffer size =  675.18 MB
[DEBUG] stable-diffusion.cpp:807  - clip tensor count = 869
ggml_init_cublas: found 1 CUDA devices:
Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6
[DEBUG] stable-diffusion.cpp:1969 - Using CPU backend - unet
[DEBUG] stable-diffusion.cpp:1977 - unet params backend buffer size =  1653.34 MB
[DEBUG] stable-diffusion.cpp:1981 - unet tensor count = 688
[DEBUG] stable-diffusion.cpp:3070 - vae params backend buffer size =  179.03 MB
[DEBUG] stable-diffusion.cpp:3079 - vae tensor count = 328
[DEBUG] stable-diffusion.cpp:3470 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:876  - clip params allocated
[DEBUG] stable-diffusion.cpp:2066 - unet params allocated
[DEBUG] stable-diffusion.cpp:3121 - vae params allocated
[DEBUG] stable-diffusion.cpp:3570 - loading weights
[DEBUG] stable-diffusion.cpp:3685 - model size = 2486.00MB
[INFO]  stable-diffusion.cpp:3696 - total memory buffer size = 2507.55MB (clip 675.18MB, unet 1653.34MB, vae 179.03MB)
[INFO]  stable-diffusion.cpp:3698 - loading model from 'data/models/v2-1_512-ema-pruned-ggml-model-f16.bin' completed, taking 1.51s
begin: diffusion compute buffer size: 10.20 MB
[DEBUG] stable-diffusion.cpp:3790 - check is_using_v_parameterization_for_sd2, taking 0.91s
[INFO]  stable-diffusion.cpp:3723 - running in eps-prediction mode
[DEBUG] stable-diffusion.cpp:3750 - finished loaded file
[DEBUG] stable-diffusion.cpp:1029 - parse 'a lovely cat' to [['a lovely cat', 1], ]
[DEBUG] stable-diffusion.cpp:443  - split prompt "a lovely cat" to tokens ["a", "lovely", "cat", ]
begin: learned condition compute buffer size: 63.83 MB
[DEBUG] stable-diffusion.cpp:3809 - computing condition graph completed, taking 463 ms
[DEBUG] stable-diffusion.cpp:1029 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:443  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:3809 - computing condition graph completed, taking 295 ms
[INFO]  stable-diffusion.cpp:4459 - get_learned_condition completed, taking 767 ms
[INFO]  stable-diffusion.cpp:4473 - start sampling
[INFO]  stable-diffusion.cpp:3969 - sampling using Euler A method
ggml_allocr_alloc: not enough space in the buffer (needed 20971520, largest block available 212032)
GGML_ASSERT: C:\Users\Jonat\Desktop\of_v20231019_vs_release\examples\eigen\ofStableDiffusionExample\src\stableDiffusion\ggml\ggml-alloc.c:148: !"not enough space in the buffer"

I will try a model you suggested. Thanks.

It seems to be a problem that only occurs with stable diffusion version 2. I can't verify it because my computer only has 16 GB of RAM.

Try disabling CUDA with -DSD_CUBLAS=OFF.
Same result. But the 1.5 model works fine.

Try disabling CUDA with -DSD_CUBLAS=OFF.

Same result. But the 1.5 model works fine.

Can you share me the model stable-diffusion 2 converted, please?

Can you share me the model stable-diffusion 2 converted, please?

https://huggingface.co/stabilityai/stable-diffusion-2-1-base/tree/main
https://huggingface.co/NikolayKozloff/stable-diffusion-nano-2-1-ckpt/tree/main

@Jonathhhan test now!

@FSSRepo yes, it is working with v2 now. Thanks.
This is a result (512x512):

[INFO]  stable-diffusion.cpp:3966 - step 1 sampling completed, taking 8.23s
[INFO]  stable-diffusion.cpp:3966 - step 2 sampling completed, taking 8.13s
[INFO]  stable-diffusion.cpp:3966 - step 3 sampling completed, taking 8.11s
[INFO]  stable-diffusion.cpp:3966 - step 4 sampling completed, taking 8.13s
[INFO]  stable-diffusion.cpp:3966 - step 5 sampling completed, taking 8.17s
[INFO]  stable-diffusion.cpp:3966 - step 6 sampling completed, taking 8.20s
[INFO]  stable-diffusion.cpp:3966 - step 7 sampling completed, taking 8.16s
[INFO]  stable-diffusion.cpp:3966 - step 8 sampling completed, taking 8.24s
[INFO]  stable-diffusion.cpp:4482 - sampling completed, taking 65.44s
compute_first_stage: vae compute buffer size: 2496.00 MB
[DEBUG] stable-diffusion.cpp:4406 - computing vae [mode: DECODE] graph completed, taking 13.82s
[INFO]  stable-diffusion.cpp:4493 - decode_first_stage completed, taking 13.82s
[INFO]  stable-diffusion.cpp:4501 - txt2img completed in 80.15s

It is not much faster (but still around one second less) than it was just with /D GGML_USE_CUBLAS, but that is what you mentioned earlier:

And you start generating an image, all the operations are performed on the GPU, but it may seem a bit slow since the parameters will be in RAM memory since the backend used is the CPU backend, but the computation is directed to the GPU.

@leejet Create a converter to the gguf format in C++ that accepts .ckpt and safetensors models. I tested it with sd 1.5 and sd 2.1 models, although it probably works with the majority. It uses the quantization function included in ggml, making this project not require the installation of PyTorch and the entire Python environment. There is still an issue with excessive memory usage because internally, gguf_write_to_file creates a copy of all the data. This could be improved by making it write through a stream.
Now, stable diffusion supports the gguf format, which will make it easier for me to implement the ggml backend v2 and see if the memory management issue in CUDA is resolved.

@Jonathhhan Could you try out the latest changes on your GPU?

@FFSRepo not sure how to compile convert.cpp for converting the model...

@Jonathhhan You must clone the repository and update the submodule, then compile with CMake. The converter is already compiled as well, without the need for anything else.

@FSSRepo thanks. I have convert.exe now. Need to figure out how to install CUDA (which seems to be needed by convert.exe) on wsl, because I have it only installed on Windows atm...

@FSSRepo thanks. I have convert.exe now. Need to figure out how to install CUDA (which seems to be needed by convert.exe) on wsl, because I have it only installed on Windows atm...
.cpp/bin/convert /mnt/c/Users/Jonat/Downloads/v1-5-pruned-emaonly.ckpt -tp f16
model type: checkpoint
reading tensors metadata
Tensor name very large
Num Tensors readed: 1131
Stable Diffusion 1.x - v1-5-pruned-emaonly.ckpt
preprocessing 1131 tensors
using embedded vocab
writing vocab: 49408 tokens
converting 1131 tensors

CUDA error 100 at /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml-cuda.cu:6111: no CUDA-capable device is detected
current device: 21850

It seems to be a bug when creating the gguf file related to ggml, as the conversion is done on the CPU. Try building the convert with DSD_CUBLAS=OFF and convert the file, then rebuild with DSD_CUBLAS=ON to use the converted model with sd.

@FSSRepo now I get:
.cpp/bin/convert /mnt/c/Users/Jonat/Downloads/v1-5-pruned-emaonly.ckpt -tp f32
model type: checkpoint
reading tensors metadata
Tensor name very large
Num Tensors readed: 1131
Stable Diffusion 1.x - v1-5-pruned-emaonly.ckpt
preprocessing 1131 tensors
using embedded vocab
writing vocab: 49408 tokens
converting 1131 tensors
Segmentation fault

It's very strange because I have the same model, and I don't encounter any errors (using Windows):
C:\proyects\stable-diffusion.cpp\build\bin\Release>convert.exe v1-5-pruned-emaonly.ckpt -tp f16
model type: checkpoint
reading tensors metadata
Tensor name very large
Num Tensors readed: 1131
Stable Diffusion 1.x - v1-5-pruned-emaonly.ckpt
preprocessing 1131 tensors
using embedded vocab
writing vocab: 49408 tokens
converting 1131 tensors
converting: alphas_cumprod | f32 => f32
converting: clip.txt_mdl.embeddings.position_embedding.weight | f32 => f16
converting: clip.txt_mdl.embeddings.token_embedding.weight | f32 => f16
converting: clip.txt_mdl.enc.layers.0.layer_norm1.bias | f32 => f32
converting: clip.txt_mdl.enc.layers.0.layer_norm1.weight | f32 => f32
converting: clip.txt_mdl.enc.layers.0.layer_norm2.bias | f32 => f32
converting: clip.txt_mdl.enc.layers.0.layer_norm2.weight | f32 => f32
converting: clip.txt_mdl.enc.layers.0.mlp.fc1.bias | f32 => f32
converting: clip.txt_mdl.enc.layers.0.mlp.fc1.weight | f32 => f16
...
converting: unet.time_embed.2.bias | f32 => f32
converting: unet.time_embed.2.weight | f32 => f16

CLIP Model Tensor count: 196
UNET Model Tensor count: 686
VAE Model Tensor count: 248

saving gguf file
model saved 'v1-5-pruned-emaonly-f16.gguf' correctly.
C:\proyects\stable-diffusion.cpp\build\bin\Release>

Could you try using the gdb?

@Jonathhhan You can try now; on my Linux, it seems to be fixed already.

@FSSRepo yes, it works now. Here is the log:
System Info:
    BLAS = 1
    SSE3 = 1
    AVX = 1
    AVX2 = 1
    AVX512 = 0
    AVX512_VBMI = 0
    AVX512_VNNI = 0
    FMA = 1
    NEON = 0
    ARM_FMA = 0
    F16C = 1
    FP16_VA = 0
    WASM_SIMD = 0
    VSX = 0
[INFO]  stable-diffusion.cpp:3385 - loading model from 'data/models/v1-5-pruned-emaonly-f16.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   0:                              sd.model.name str
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   1:                             sd.model.dtype i32
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   2:                           sd.model.version i8
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   3:                            sd.vocab.tokens arr
[INFO]  stable-diffusion.cpp:3413 - Stable Diffusion 1.x | v1-5-pruned-emaonly.ckpt
[INFO]  stable-diffusion.cpp:3422 - model data type: f16
[DEBUG] stable-diffusion.cpp:3426 - loading vocab
[DEBUG] stable-diffusion.cpp:3445 - ggml tensor size = 352 bytes
[DEBUG] stable-diffusion.cpp:795  - Using CPU backend - clip
[DEBUG] stable-diffusion.cpp:802  - clip params backend buffer size =  236.18 MB
[DEBUG] stable-diffusion.cpp:805  - clip tensor count = 437
[DEBUG] stable-diffusion.cpp:1968 - Using CPU backend - unet
[DEBUG] stable-diffusion.cpp:1976 - unet params backend buffer size =  1641.16 MB
[DEBUG] stable-diffusion.cpp:1980 - unet tensor count = 688
[DEBUG] stable-diffusion.cpp:3071 - vae params backend buffer size =  179.03 MB
[DEBUG] stable-diffusion.cpp:3080 - vae tensor count = 328
[DEBUG] stable-diffusion.cpp:3456 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:874  - clip params allocated
[DEBUG] stable-diffusion.cpp:2065 - unet params allocated
[DEBUG] stable-diffusion.cpp:3122 - vae params allocated
[DEBUG] stable-diffusion.cpp:3474 - loading weights
[DEBUG] stable-diffusion.cpp:3578 - model size = 2034.89MB
[INFO]  stable-diffusion.cpp:3588 - total memory buffer size = 2056.36MB (clip 236.18MB, unet 1641.16MB, vae 179.03MB)
[INFO]  stable-diffusion.cpp:3590 - loading model from 'data/models/v1-5-pruned-emaonly-f16.gguf' completed, taking 0.55s
[INFO]  stable-diffusion.cpp:3614 - running in eps-prediction mode
[DEBUG] stable-diffusion.cpp:3641 - finished loaded file
[DEBUG] stable-diffusion.cpp:1028 - parse 'a lovely cat' to [['a lovely cat', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "a lovely cat" to tokens ["a</w>", "lovely</w>", "cat</w>", ]
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 130 ms
[DEBUG] stable-diffusion.cpp:1028 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 93 ms
[INFO]  stable-diffusion.cpp:4351 - get_learned_condition completed, taking 227 ms
[INFO]  stable-diffusion.cpp:4365 - start sampling
[DEBUG] stable-diffusion.cpp:2313 - diffusion compute buffer size: 1147.33 MB
[INFO]  stable-diffusion.cpp:3861 - sampling using Euler A method
[INFO]  stable-diffusion.cpp:3854 - step 1 sampling completed, taking 6.06s
[INFO]  stable-diffusion.cpp:3854 - step 2 sampling completed, taking 5.79s
[INFO]  stable-diffusion.cpp:3854 - step 3 sampling completed, taking 5.80s
[INFO]  stable-diffusion.cpp:3854 - step 4 sampling completed, taking 5.82s
[INFO]  stable-diffusion.cpp:3854 - step 5 sampling completed, taking 5.89s
[INFO]  stable-diffusion.cpp:3854 - step 6 sampling completed, taking 5.92s
[INFO]  stable-diffusion.cpp:3854 - step 7 sampling completed, taking 5.94s
[INFO]  stable-diffusion.cpp:3854 - step 8 sampling completed, taking 5.95s
[INFO]  stable-diffusion.cpp:4370 - sampling completed, taking 47.21s
[DEBUG] stable-diffusion.cpp:4284 - vae compute buffer size: 2496.00 MB
CUDA Fallback disabled
[DEBUG] stable-diffusion.cpp:4294 - computing vae [mode: DECODE] graph completed, taking 20.96s
[INFO]  stable-diffusion.cpp:4381 - decode_first_stage completed, taking 20.97s
[INFO]  stable-diffusion.cpp:4389 - txt2img completed in 68.40s

@FSSRepo with v1-5-pruned-emaonly-f32.gguf it works to generate a second image, but it seems that CUDA is only used for the first image:
System Info:
    BLAS = 1
    SSE3 = 1
    AVX = 1
    AVX2 = 1
    AVX512 = 0
    AVX512_VBMI = 0
    AVX512_VNNI = 0
    FMA = 1
    NEON = 0
    ARM_FMA = 0
    F16C = 1
    FP16_VA = 0
    WASM_SIMD = 0
    VSX = 0
[INFO]  stable-diffusion.cpp:3385 - loading model from 'data/models/v1-5-pruned-emaonly-f32.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   0:                              sd.model.name str
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   1:                             sd.model.dtype i32
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   2:                           sd.model.version i8
[DEBUG] stable-diffusion.cpp:3403 - load_from_file: - kv   3:                            sd.vocab.tokens arr
[INFO]  stable-diffusion.cpp:3413 - Stable Diffusion 1.x | v1-5-pruned-emaonly.ckpt
[INFO]  stable-diffusion.cpp:3422 - model data type: f32
[DEBUG] stable-diffusion.cpp:3426 - loading vocab
[DEBUG] stable-diffusion.cpp:3445 - ggml tensor size = 352 bytes
[DEBUG] stable-diffusion.cpp:795  - Using CPU backend - clip
[DEBUG] stable-diffusion.cpp:802  - clip params backend buffer size =  470.66 MB
[DEBUG] stable-diffusion.cpp:805  - clip tensor count = 437
[DEBUG] stable-diffusion.cpp:1968 - Using CPU backend - unet
[DEBUG] stable-diffusion.cpp:1976 - unet params backend buffer size =  2156.24 MB
[DEBUG] stable-diffusion.cpp:1980 - unet tensor count = 688
[DEBUG] stable-diffusion.cpp:3071 - vae params backend buffer size =  179.03 MB
[DEBUG] stable-diffusion.cpp:3080 - vae tensor count = 328
[DEBUG] stable-diffusion.cpp:3456 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:874  - clip params allocated
[DEBUG] stable-diffusion.cpp:2065 - unet params allocated
[DEBUG] stable-diffusion.cpp:3122 - vae params allocated
[DEBUG] stable-diffusion.cpp:3474 - loading weights
[DEBUG] stable-diffusion.cpp:3578 - model size = 2784.45MB
[INFO]  stable-diffusion.cpp:3588 - total memory buffer size = 2805.93MB (clip 470.66MB, unet 2156.24MB, vae 179.03MB)
[INFO]  stable-diffusion.cpp:3590 - loading model from 'data/models/v1-5-pruned-emaonly-f32.gguf' completed, taking 0.85s
[INFO]  stable-diffusion.cpp:3614 - running in eps-prediction mode
[INFO]  stable-diffusion.cpp:3624 - running with Karras schedule
[DEBUG] stable-diffusion.cpp:3641 - finished loaded file
[DEBUG] stable-diffusion.cpp:1028 - parse 'a lovely cat' to [['a lovely cat', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "a lovely cat" to tokens ["a</w>", "lovely</w>", "cat</w>", ]
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 166 ms
[DEBUG] stable-diffusion.cpp:1028 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 120 ms
[INFO]  stable-diffusion.cpp:4351 - get_learned_condition completed, taking 290 ms
[INFO]  stable-diffusion.cpp:4365 - start sampling
[DEBUG] stable-diffusion.cpp:2313 - diffusion compute buffer size: 1147.33 MB
[INFO]  stable-diffusion.cpp:3861 - sampling using Euler A method
[INFO]  stable-diffusion.cpp:3854 - step 1 sampling completed, taking 6.15s
[INFO]  stable-diffusion.cpp:3854 - step 2 sampling completed, taking 6.15s
[INFO]  stable-diffusion.cpp:3854 - step 3 sampling completed, taking 6.18s
[INFO]  stable-diffusion.cpp:3854 - step 4 sampling completed, taking 6.23s
[INFO]  stable-diffusion.cpp:3854 - step 5 sampling completed, taking 6.21s
[INFO]  stable-diffusion.cpp:3854 - step 6 sampling completed, taking 6.20s
[INFO]  stable-diffusion.cpp:3854 - step 7 sampling completed, taking 6.18s
[INFO]  stable-diffusion.cpp:3854 - step 8 sampling completed, taking 6.12s
[INFO]  stable-diffusion.cpp:4370 - sampling completed, taking 49.53s
[DEBUG] stable-diffusion.cpp:4284 - vae compute buffer size: 2496.00 MB
CUDA Fallback disabled
[DEBUG] stable-diffusion.cpp:4294 - computing vae [mode: DECODE] graph completed, taking 73.43s
[INFO]  stable-diffusion.cpp:4381 - decode_first_stage completed, taking 73.43s
[INFO]  stable-diffusion.cpp:4389 - txt2img completed in 123.25s
[DEBUG] stable-diffusion.cpp:1028 - parse 'a lovely cat' to [['a lovely cat', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "a lovely cat" to tokens ["a</w>", "lovely</w>", "cat</w>", ]
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 290 ms
[DEBUG] stable-diffusion.cpp:1028 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:444  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:962  - learned condition compute buffer size: 25.54 MB
[DEBUG] stable-diffusion.cpp:3701 - computing condition graph completed, taking 269 ms
[INFO]  stable-diffusion.cpp:4351 - get_learned_condition completed, taking 563 ms
[INFO]  stable-diffusion.cpp:4365 - start sampling
[DEBUG] stable-diffusion.cpp:2313 - diffusion compute buffer size: 1147.33 MB
[INFO]  stable-diffusion.cpp:3861 - sampling using Euler A method
[INFO]  stable-diffusion.cpp:3854 - step 1 sampling completed, taking 44.84s
[INFO]  stable-diffusion.cpp:3854 - step 2 sampling completed, taking 44.73s
[INFO]  stable-diffusion.cpp:3854 - step 3 sampling completed, taking 44.75s
[INFO]  stable-diffusion.cpp:3854 - step 4 sampling completed, taking 44.87s
[INFO]  stable-diffusion.cpp:3854 - step 5 sampling completed, taking 44.21s
[INFO]  stable-diffusion.cpp:3854 - step 6 sampling completed, taking 44.69s
[INFO]  stable-diffusion.cpp:3854 - step 7 sampling completed, taking 44.61s
[INFO]  stable-diffusion.cpp:3854 - step 8 sampling completed, taking 44.39s
[INFO]  stable-diffusion.cpp:4370 - sampling completed, taking 357.23s
[DEBUG] stable-diffusion.cpp:4284 - vae compute buffer size: 2496.00 MB
[DEBUG] stable-diffusion.cpp:4294 - computing vae [mode: DECODE] graph completed, taking 73.25s
[INFO]  stable-diffusion.cpp:4381 - decode_first_stage completed, taking 73.25s
[INFO]  stable-diffusion.cpp:4389 - txt2img completed in 431.04s

I guess my mistake was, that I was compiling with /D SD_CUBLAS=ON /D GGML_USE_CUBLAS. With /D SD_CUBLAS=ON alone I do not have that issue. Interestingly image generation with /D SD_CUBLAS=ON takes twice as long as with /D GGML_USE_CUBLAS (but still much faster than with cpu alone).

Maybe my mistake was, that I was compiling with /D SD_USE_CUBLAS and /D GGML_USE_CUBLAS together?
It was because I thought /D GGML_USE_CUBLAS is also needed for using CUDA. And with /D GGML_USE_CUBLAS its twice as fast, while /D GGML_USE_CUBLAS does not make a big difference.

This difference shouldn't happen since the same tensor values are being used. The only distinction is that one falls back to GPU, and the other I am using the CUDA backend entirely. In the end, matrix multiplication is being performed on the GPU in both cases, so it shouldn't differ for any reason. I have already tried making the two source tensors contiguous in the CUDA backend case, and the same difference persists.
Fallback to CUDA - Backend CUDA
========= UNET - compute embedding-src0:            NONE ========
[0] | 34998.000000 - 34998.000000
[1] | 5801.000000 - 5801.000000
[2] | 4793.000000 - 4793.000000
[3] | 6385.000000 - 6385.000000
[4] | 5797.000000 - 5797.000000
[409596] | 8951.000000 - 8951.000000
[409597] | 8392.000000 - 8392.000000
[409598] | 8368.000000 - 8368.000000
[409599] | 9412.000000 - 9412.000000
sum: 0.0000 | diff: 0.0000 | avg: 0.0000
========= UNET - compute embedding-src1:            CONT ========
[0] | 0.999650 - 0.999650
[1] | 0.802678 - 0.802678
[2] | -0.278121 - -0.278121
[3] | 0.176803 - 0.176803
[4] | -0.278017 - -0.278017
[316] | 0.125435 - 0.125435
[317] | 0.118453 - 0.118453
[318] | 0.111855 - 0.111855
[319] | 0.105622 - 0.105622
sum: 0.0000 | diff: 0.0000 | avg: 0.0000
========= UNET - compute embedding:         MUL_MAT ========
[0] | 0.064568 - 0.064636
[1] | -0.027617 - -0.027771
[2] | 0.060006 - 0.059998
[3] | -0.004918 - -0.004974
[4] | -0.007953 - -0.007965
[1276] | 0.071102 - 0.071045
[1277] | 0.078870 - 0.078918
[1278] | -0.072765 - -0.072876
[1279] | 0.001293 - 0.001343
sum: 0.2171 | diff: 0.0073 | avg: 0.0002

I don't know what to do at this point.

 I was using /D GGML_USE_CUBLAS instead of /D GGML_CUBLAS.
/D SD_USE_CUBLAS /D GGML_CUBLAS it works as expected.

@slaren @ggerganov  I have already identified the error that produces NaN, and the truth is that it seems very strange to me that it occurs in a ggml_reshape_4d operation that does not modify the data of the tensor but only changes its shape. However, at a certain point, with the same range of elements, this happens for some reason only in the CUDA backend.
How I arrived at this result: what I did was directly extract the nodes from the graph after they were calculated in CUDA, considering that they could come from a different backend, so it's not a data transfer error or anything. I did this for both the CPU backend but directed all calculations to CUDA, using the CUDA backend in its entirety.
// DEBUG STABLE DIFFUSION
#if 1
#include <stdarg.h>

bool track_unet = false;
int clip = 0;
int view_count = 0;

void log_file(const char* format, ...) {
    FILE *file = fopen("C:\\proyects\\nodes.txt", "a");
    if (file == NULL) {
        perror("Error opening file");
        return;
    }

    va_list args;
    va_start(args, format);
    vfprintf(file, format, args);
    va_end(args);
    fclose(file);
}

void save_tensor(ggml_tensor* tensor,const char* target, const char* name, int num, int source, const char* source_name) {
    char file_name[150];
    sprintf(file_name,
        source != -1 ? "C:\\proyects\\output%s\\%s.%i-src%i_%s.tensor" :
        "C:\\proyects\\output%s\\%s.%i.tensor", target, name, num, source, source_name);
    FILE* fp = fopen(file_name, "wb");
    int ttype = (int)tensor->type;
    int op = (int)tensor->op;
    fwrite(&tensor->n_dims, 1, sizeof(tensor->n_dims), fp);
    fwrite(&ttype, 1, sizeof(ttype), fp);
    fwrite(&op, 1, sizeof(op), fp);
    for (int i = 0; i < tensor->n_dims; ++i) {
        int ne_ = (int) tensor->ne[i];
        fwrite(&ne_, 1, sizeof(ne_), fp);
    }
    if(tensor->backend == GGML_BACKEND_GPU) {
        void* cpu_block = malloc(ggml_nbytes(tensor));
        ggml_backend_tensor_get(tensor, cpu_block, 0, ggml_nbytes(tensor));
        fwrite(cpu_block, 1, ggml_nbytes(tensor), fp);
        free(cpu_block);
    } else {
        fwrite(tensor->data, 1, ggml_nbytes(tensor), fp);
    }
    fclose(fp);
}
#endif

bool ggml_cuda_compute_forward(struct ggml_compute_params * params, struct ggml_tensor * tensor) {
...
...
    func(tensor->src[0], tensor->src[1], tensor);
#if 1
    // DEBUG STABLE DIFFUSION
    if(strcmp(ggml_get_name(tensor), "node_1") == 0) { // skip CLIP Model
        track_unet = clip == 2;
        clip++;
    }
    if(track_unet && tensor->op != GGML_OP_CONT && tensor->op != GGML_OP_REPEAT && ggml_nbytes(tensor) < 10 * 1024 * 1024) {
        const char* t_name = ggml_get_name(tensor);
        if(strcmp(t_name, "node_1313") == 0) { // this node differs of correct calculation, then extract the source tensors
            printf("saving %s -> %s %s\n", t_name, ggml_type_name(tensor->src[0]->type), ggml_type_name(tensor->src[1]->type));
            if(tensor->op != GGML_OP_ADD) {
                printf("Error");
                return false;
            }
            save_tensor(tensor->src[0], !any_on_device ? "" : "-cuda", t_name, view_count, 0, ggml_get_name(tensor->src[0]));
            log_file("%s.%i-src0_%s\n", t_name, view_count, ggml_get_name(tensor->src[0]));
            if(tensor->src[1] != nullptr) {
                save_tensor(tensor->src[1], !any_on_device ? "" : "-cuda", t_name, view_count, 1, ggml_get_name(tensor->src[1]));
                log_file("%s.%i-src1_%s\n", t_name, view_count, ggml_get_name(tensor->src[1]));
            }
        }
        save_tensor(tensor, !any_on_device ? "" : "-cuda", t_name, view_count, -1, "");
        log_file("%s.%i\n", t_name, view_count);
        view_count++;
        track_unet = strcmp(t_name, "UNET Finish") != 0;
    }
#endif

The tensors saved in files are compared by a program that I created, between output/*.tensors and output-cuda/*.tensors. And this was my result diff.txt.
Expected - Result got
====== 907 nodes are the same ====
DownSample.ggml_conv_2d
=========   node_1306.908:          IM2COL ========
=========  (reshaped).909:         RESHAPE ========
=========   node_1309.910:         MUL_MAT ========
=========  (reshaped).911:         RESHAPE ========
========= leaf_495 (reshaped).912:         RESHAPE ========
========= node_1313.913-src0_ (reshaped):         RESHAPE ========
[64] | -0.064575 - -0.083374
[65] | 0.822998 - 0.804199
[66] | 0.315186 - 0.296387
[67] | 0.213379 - 0.194580
[68] | -0.182861 - -0.201660
[69] | 0.581787 - 0.562988
[70] | 0.611572 - 0.592773
[71] | -12.768799 - -12.787598
[72] | -0.514893 - -0.533691
[73] | -2.653564 - -2.672363
[74] | -3.124268 - -3.143066
[75] | -4.483643 - -4.502441
[76] | -3.282471 - -3.301270
[77] | -2.776611 - -2.795410
[78] | -1.917236 - -1.936035
[79] | -0.438232 - -0.457031
[80] | 0.397461 - 0.378662
[81] | -1.384033 - -1.402832
[82] | -1.179932 - -1.198730
[83] | -2.741455 - -2.760254
[84] | -2.175049 - -2.193848
[85] | -1.257080 - -1.275879
[86] | -1.279541 - -1.298340
[87] | -0.096436 - -0.115234
[88] | -0.425537 - -0.444336
[89] | -1.247314 - -1.266113
[90] | -0.580322 - -0.599121
[91] | -1.952393 - -1.971191
[92] | -2.442627 - -2.461426
[93] | -1.206299 - -1.225098
[94] | -1.929932 - -1.948730
[95] | -0.755615 - -0.774414
[96] | -0.780518 - -0.799316
[97] | -1.012939 - -1.031738
[98] | -0.980225 - -0.999023
[99] | -1.137939 - -1.156738
[100] | -2.335205 - -2.354004
[101] | -1.342041 - -1.360840
[102] | -1.116455 - -1.135254
[103] | -2.581299 - -2.600098
[104] | -0.876221 - -0.895020
[105] | -0.557373 - -0.576172
[106] | -1.130127 - -1.148926
[107] | -1.634033 - -1.652832
[108] | -0.835693 - -0.854492
[109] | -0.975342 - -0.994141
[110] | -1.453369 - -1.472168
[111] | -2.552002 - -2.570801
[112] | -0.099182 - -0.117981
[113] | -0.276123 - -0.294922
[114] | -0.053467 - -0.072266
[115] | -1.451416 - -1.470215
[116] | -1.898682 - -1.917480
[117] | -0.780518 - -0.799316
[118] | -0.740967 - -0.759766
[119] | -1.302979 - -1.321777
[120] | -1.677002 - -1.695801
[121] | -1.841064 - -1.859863
[122] | -1.365479 - -1.384277
[123] | -1.784424 - -1.803223
[124] | -0.824463 - -0.843262
[125] | -0.946533 - -0.965332
[126] | -0.863037 - -0.881836
[127] | -1.076416 - -1.095215
[128] | 2.249146 - 2.200562
[129] | 0.255981 - 0.207397
[130] | 0.192993 - 0.144409
[131] | -0.188599 - -0.237183
[132] | -0.955444 - -1.004028
[133] | -1.004761 - -1.053345
[134] | 0.160645 - 0.112061
[135] | -3.637573 - -3.686157
[136] | 0.205811 - 0.157227
[137] | -1.526245 - -1.574829
[138] | -1.247925 - -1.296509
[139] | -1.507690 - -1.556274
[140] | -2.215698 - -2.264282
[141] | -1.834839 - -1.883423
[142] | -0.789917 - -0.838501
[143] | -0.168823 - -0.217407
[144] | 0.971802 - 0.923218
[145] | -0.194702 - -0.243286
[146] | -0.614624 - -0.663208
[147] | -1.289917 - -1.338501
[148] | -1.118042 - -1.166626
[149] | -1.609253 - -1.657837
[150] | -0.751831 - -0.800415
[151] | 0.182983 - 0.134399
[152] | 0.303589 - 0.255005
[153] | 0.499146 - 0.450562
[154] | -0.518433 - -0.567017
[155] | 1.111450 - 1.062866
[156] | -0.216431 - -0.265015
[157] | 0.066406 - 0.017822
[158] | 0.525513 - 0.476929
[159] | 0.707153 - 0.658569
[160] | -0.302368 - -0.350952
[161] | -0.870972 - -0.919556
[162] | 0.210083 - 0.161499
[163] | -0.141235 - -0.189819
[164] | -0.018433 - -0.067017
[165] | -0.115112 - -0.163696
[166] | 0.590942 - 0.542358
[167] | 0.681763 - 0.633179
[168] | -0.803589 - -0.852173
[169] | -1.966675 - -2.015259
[170] | -0.286499 - -0.335083
[171] | -1.099487 - -1.148071
[172] | -1.865112 - -1.913696
[173] | -1.215698 - -1.264282
[174] | -1.346558 - -1.395142
[175] | -0.598999 - -0.647583
[176] | -0.268433 - -0.317017
[177] | 0.197144 - 0.148560
[178] | -1.863159 - -1.911743
[179] | -1.562378 - -1.610962
[180] | -2.313354 - -2.361938
[181] | -2.319214 - -2.367798
[182] | -0.727417 - -0.776001
[183] | 0.591919 - 0.543335
[184] | 0.748657 - 0.700073
[185] | -1.558472 - -1.607056
[186] | -1.594604 - -1.643188
[187] | -1.845581 - -1.894165
[188] | -2.233276 - -2.281860
[189] | -2.473511 - -2.522095
[190] | -1.244995 - -1.293579
[191] | 0.263550 - 0.214966
[192] | -2.021240 - -2.065674
[193] | -0.520752 - -0.565186
[194] | -0.142700 - -0.187134
[195] | -0.318359 - -0.362793
[196] | 0.147217 - 0.102783
[197] | 0.096985 - 0.052551
[198] | -0.368164 - -0.412598
[199] | -0.280029 - -0.324463
[200] | -0.028076 - -0.072510
[201] | 1.625244 - 1.580811
[202] | 0.323975 - 0.279541
[203] | 0.774658 - 0.730225
[204] | 0.989502 - 0.945068
[205] | 0.791260 - 0.746826
[206] | 1.777588 - 1.733154
[207] | 0.321777 - 0.277344
[208] | 1.158447 - 1.114014
[209] | 1.747314 - 1.702881
[210] | -0.736084 - -0.780518
[211] | 0.483643 - 0.439209
[212] | 1.638916 - 1.594482
[213] | 0.303711 - 0.259277
[214] | 2.407471 - 2.363037
[215] | 1.509033 - 1.464600
[216] | -0.015381 - -0.059814
[217] | 0.063110 - 0.018677
[218] | -1.278076 - -1.322510
[219] | 1.849854 - 1.805420
[220] | 1.762939 - 1.718506
[221] | 0.506104 - 0.461670
[222] | 1.128174 - 1.083740
[223] | 0.418701 - 0.374268
[224] | -1.208740 - -1.253174
[225] | -1.436279 - -1.480713
[226] | -0.382568 - -0.427002
[227] | 1.721924 - 1.677490
[228] | 0.406006 - 0.361572
[229] | 0.161865 - 0.117432
[230] | -0.120728 - -0.165161
[231] | -1.197998 - -1.242432
[232] | -0.981201 - -1.025635
[233] | 0.056519 - 0.012085
[234] | -1.854248 - -1.898682
[235] | -0.348877 - -0.393311
[236] | -2.252686 - -2.297119
[237] | -0.079834 - -0.124268
[238] | -0.020264 - -0.064697
[239] | -0.811279 - -0.855713
[240] | -1.385498 - -1.429932
[241] | -0.901123 - -0.945557
[242] | -1.254639 - -1.299072
[243] | -0.377197 - -0.421631
[244] | -1.026123 - -1.070557
[245] | 0.031494 - -0.012939
[246] | 0.604736 - 0.560303
[247] | 0.985596 - 0.941162
[248] | -2.082764 - -2.127197
[249] | -0.417236 - -0.461670
[250] | -1.188232 - -1.232666
[251] | 0.198486 - 0.154053
[252] | 0.641846 - 0.597412
[253] | 0.352295 - 0.307861
[254] | 1.655518 - 1.611084
[255] | 1.689697 - 1.645264
[8, 8, 1280, 1]
sum: 7.1562 | diff: 0.0486 | avg: 0.0001
========= node_1313.913-src1_down_sample_end.8:          REPEAT ========
=========   node_1313.913:             ADD ========
[64] | -0.064575 - -0.083374
[65] | 0.822998 - 0.804199
[66] | 0.315186 - 0.296387
[67] | 0.213379 - 0.194580
[68] | -0.182861 - -0.201660
[69] | 0.581787 - 0.562988
[70] | 0.611572 - 0.592773
... All next tensors are different. The im2col operation, which is highly sensitive to changes because it degrades float32 to float16, is starting to show significant differences.

struct ggml_tensor* forward(struct ggml_context* ctx,
                                struct ggml_tensor* x,
                                struct ggml_tensor* timesteps,
                                struct ggml_tensor* context,
                                struct ggml_tensor* t_emb = NULL) {
...
...
for (int i = 0; i < len_mults; i++) {
            int mult = channel_mult[i];
            for (int j = 0; j < num_res_blocks; j++) {
                h = input_res_blocks[i][j].forward(ctx, h, emb);  // [N, mult*model_channels, h, w]
                if (ds == attention_resolutions[0] || ds == attention_resolutions[1] || ds == attention_resolutions[2]) {
                    h = input_transformers[i][j].forward(ctx, h, context);  // [N, mult*model_channels, h, w]
                }
                hs.push_back(h);
            }
            if (i != len_mults - 1) {
                ds *= 2;
                h = input_down_samples[i].forward(ctx, h);  // In the last iteration, it begins to differ.
                hs.push_back(h);
            }
        }

Source code of DownSample.forward:
struct ggml_tensor* forward(struct ggml_context* ctx, struct ggml_tensor* x) {
        // x: [N, channels, h, w]
        struct ggml_tensor* c = nullptr;
        if (vae_downsample) {
            auto pad_x = ggml_new_tensor_4d(ctx, x->type, x->ne[0] + 1, x->ne[1] + 1, x->ne[2], x->ne[3]);

            c = ggml_map_custom2_inplace(ctx, pad_x, x, asymmetric_pad, 1, NULL);
            c = ggml_conv_2d(ctx, op_w, c, 2, 2, 0, 0, 1, 1);
        } else {
            ggml_format_name(x, "down_sample.%i", index);
            c = ggml_conv_2d(ctx, op_w, x, 2, 2, 1, 1, 1, 1);
            // im2col pass
           // mul_mat pass
           // reshape 4d fail (The data was unexpectedly changed when it should have been a view of mul_mat), And interestingly, the data changes within the index range [64-255] every time I've tried. If it were an operational error, either im2col or mul_mat should have failed.
        }
        auto r = ggml_repeat(ctx,
                                 ggml_reshape_4d(ctx, op_b, 1, 1, op_b->ne[0], 1),
                                 c);
        ggml_format_name(r, "down_sample_end.%i", index);
        c = ggml_add(ctx,
                     c,
                     r);  // [N, out_channels, h/2, w/2]
        return c;
    }

I have no idea what's going on, as DownSample.forward is executed twice successfully before failing.

@FSSRepo if you can create a small self-contained program that can reproduce the issue, it would be the best way for me to able to look into this.

@FSSRepo if you can create a small self-contained program that can reproduce the issue, it would be the best way for me to able to look into this.

ok

@slaren Create a new branch in my ggml repository that reproduces this error. The input part of the UNet model triggers the error, and since the error occurs almost at the end of this stage, it was inevitable to include some complex blocks like ResBlock, SpatialTransformer. Execute the following commands; I'm not sure if you have Linux or Windows, but I can help you with either.
git clone https://github.com/FSSRepo/ggml.git
cd ggml
git checkout ri-sd-cpp
curl -O https://huggingface.co/Steward/unet-partial/resolve/main/context.tensor
curl -O https://huggingface.co/Steward/unet-partial/resolve/main/expected.tensor
curl -O https://huggingface.co/Steward/unet-partial/resolve/main/t_emb.tensor
curl -O https://huggingface.co/Steward/unet-partial/resolve/main/noise.tensor
curl -O https://huggingface.co/Steward/unet-partial/resolve/main/unet-model.gguf
mkdir build
cmake -S . -B build -DGGML_CUBLAS=ON
cmake --build build --config Release

# run CPU Backend but fallback compute to CUDA
build/bin/stable fb

# run full CUDA Backend
build/bin/stable

# run comparison between output tensors
build/bin/stable -c

@FSSRepo if you build in debug, you should see an assert similar to this one:
stable: /home/diego/code/sd/ggml/src/ggml-cuda.cu:8862: void ggml_backend_cuda_graph_compute(ggml_backend_t, ggml_cgraph*): Assertion `node->src[j]->backend == GGML_BACKEND_GPU' failed.

There are no tensor names so it is a bit hard to debug this, but it is probably caused by this:
 q = ggml_scale_inplace(ctx, q, ggml_new_f32(ctx, 1.0f / sqrt((float)d_head)));

You cannot use ggml_new_f32 with ggml-alloc or with ggml-backend. If you are doing this correctly, the context should be no_alloc and the ggml_new_f32 should fail immediately because it would try to write to a NULL pointer, but since that is not happening, there may be deeper issues. You need to allocate the tensor with ggml-alloc and then set the value with ggml-backend, the gpt-2 example shows how to do this. The repro is too big for me to debug this, a minimal repro should only contain the bare minimum to reproduce the issue.

@slaren I already applied the suggestion you made, and it still has issues at the same point. I can't make the playback program any smaller, as isolating the inputs from DownSample and the expected outputs could work as if nothing happened.
                           node_1015.754 (       ADD): PASS
                          (permuted).755 (   PERMUTE): PASS
                           node_1018.756 (    IM2COL): PASS
                          (reshaped).757 (   RESHAPE): PASS
                   leaf_9 (reshaped).758 (   RESHAPE): PASS
                           node_1021.759 (   MUL_MAT): PASS
                          (reshaped).760 (   RESHAPE): PASS
                 leaf_247 (reshaped).761 (   RESHAPE): PASS
                           node_1025.762 (       ADD): PASS
                       down_sample.9.763 (       ADD): PASS
                           node_1027.764 (    IM2COL): PASS
                          (reshaped).765 (   RESHAPE): PASS
                           node_1030.766 (   MUL_MAT): PASS
                down_sample.9.conv2d.767 (   RESHAPE): PASS
                 leaf_248 (reshaped).768 (   RESHAPE): PASS
node_1034.769.src0 (down_sample.9.conv2d) (   RESHAPE): FAILED <-- DownSample last iteration ggml_conv2d
[64] | -0.064575 - -0.083374
[65] | 0.822998 - 0.804199
[66] | 0.315186 - 0.296387
[67] | 0.213379 - 0.194580
[68] | -0.182861 - -0.201660
[69] | 0.581787 - 0.562988
[70] | 0.611572 - 0.592773
[71] | -12.768799 - -12.787598

If you make ggml_op_can_inplace in ggml-alloc.c always return false to disable auto-inplace, the differences look within the normal margin for different backends. Does that fix the issue in the full program?

If you make ggml_op_can_inplace in ggml-alloc.c always return false to disable auto-inplace, the differences look within the normal margin for different backends. Does that fix the issue in the full program?

The headache is gone, thank you very much for your help. It works perfectly fine.

Great that it works now. Not sure if this is helpful, but it is enough to remove case GGML_OP_ADD: and case GGML_OP_MUL:.

UNet Full CUDA Backend

compute graph
CPU Backend
CUDA Backend
Speed up

CLIP Model
240 ms
26 ms
9.2

UNet per step
17.1 s
2.1 s
8.02

Results
CLI params:
sd -m $model -p "beautiful anime girl, white hair, blue eyes, realistic, masterpiece, azur lane, 4k, high quality" -n "bad quality, ugly, face malformed, bad anatomy" --sampling-method dpm++2m --steps 20 -t 6 -s 424354

CPU Backend
CUDA Backend

With CPU Backend it takes 360.87 segs, now with CUDA it takes 65 segs.

Nice speedup!
The issue is probably in the add and mul kernels. This implementation produces the same results with inplace:
template<typename dst_t, typename src0_t, typename src1_t>
static __global__ void k_add(const src0_t * x, const src1_t * y, dst_t * dst, int n_elements) {
    const int idx = blockDim.x * blockIdx.x + threadIdx.x;
    if(idx >= n_elements) {
        return;
    }
    dst[idx] = (dst_t)((float)x[idx] + (float)y[idx]);
}

static __global__ void mul_f32(const float * x, const float * y, float * dst, int n_elements) {
    const int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if(idx >= n_elements) {
        return;
    }
    dst[idx] = x[idx] * y[idx];
}

template<typename dst_t, typename src0_t, typename src1_t>
static void add_cuda(const src0_t * x, const src1_t * y, dst_t * dst, const int ne0, const int ne1, const int ne2, cudaStream_t stream) {
    int n_els = ne0 * ne1 * ne2;
    int num_blocks = (n_els + CUDA_ADD_BLOCK_SIZE - 1) / CUDA_ADD_BLOCK_SIZE;
    k_add<<<num_blocks, CUDA_ADD_BLOCK_SIZE, 0, stream>>>(x, y, dst, n_els);
}

static void mul_f32_cuda(const float * x, const float * y, float * dst, const int ne0, const int ne1, const int ne2, cudaStream_t stream) {
    int n_els = ne0 * ne1 * ne2;
    int num_blocks = (n_els + CUDA_MUL_BLOCK_SIZE - 1) / CUDA_MUL_BLOCK_SIZE;
    mul_f32<<<num_blocks, CUDA_MUL_BLOCK_SIZE, 0, stream>>>(x, y, dst, n_els);
}

inline void ggml_cuda_op_mul(
    const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
    const float * src0_dd, const float * src1_dd, float * dst_dd, const cudaStream_t & main_stream) {

    GGML_ASSERT(src0->type == GGML_TYPE_F32);
    GGML_ASSERT(src1->type == GGML_TYPE_F32);
    GGML_ASSERT( dst->type == GGML_TYPE_F32);

    mul_f32_cuda(src0_dd, src1_dd, dst_dd, dst->ne[0], dst->ne[1], dst->ne[2], main_stream);

    (void) dst;
}

inline void ggml_cuda_op_add(
    const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst,
    const float * src0_dd, const float * src1_dd, float * dst_dd, const cudaStream_t & main_stream) {

    GGML_ASSERT(src1->type == GGML_TYPE_F32);

    if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) {
        add_cuda(src0_dd, src1_dd, dst_dd, dst->ne[0], dst->ne[1], dst->ne[2], main_stream);
    } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) {
        add_cuda((const half *) src0_dd, src1_dd, (half *) dst_dd, dst->ne[0], dst->ne[1], dst->ne[2], main_stream);
    } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) {
        add_cuda((const half *) src0_dd, src1_dd, dst_dd, dst->ne[0], dst->ne[1], dst->ne[2], main_stream);
    } else {
        fprintf(stderr, "src0->type: %d  dst->type: %d\n", src0->type, dst->type);
        GGML_ASSERT(false);
    }

    (void) src1;
    (void) dst;
}

However this implementation is not good either, since it removes support for broadcasting. All of the repeats in adds and muls need to be removed and implemented with broadcasting.

I will try it

The issue is probably in the add and mul kernels. This implementation produces the same results with or without automatic inplace:

That was definitely the problem, thank you very much for helping me. Now I'll be able to focus on improving memory usage and optimizing some kernels.

Adding broadcasting to the add and mul kernels eliminates the need for repeat operations. I will also attempt to port the flash attention kernel to CUDA for cross-attention, and optimize the sampling phase to address the avoidable 500 ms delay.
UNet forward
========== CUDA Timings =========
[       ADD] - 92.428001 ms - 419 - 0.220592 ms
[       MUL] - 15.241000 ms - 125 - 0.121928 ms
[    REPEAT] - 76.947998 ms - 269 - 0.286052 ms
[    CONCAT] - 0.046000 ms - 12 - 0.003833 ms
[      NORM] - 0.171000 ms - 48 - 0.003563 ms
[GROUP_NORM] - 15.509000 ms - 61 - 0.254246 ms
[   MUL_MAT] - 189.089996 ms - 329 - 0.574742 ms
[     SCALE] - 1.059000 ms - 32 - 0.033094 ms
[      CONT] - 0.502000 ms - 129 - 0.003891 ms
[  SOFT_MAX] - 0.114000 ms - 32 - 0.003563 ms
[    IM2COL] - 16.172001 ms - 97 - 0.166722 ms
[   UPSCALE] - 0.040000 ms - 3 - 0.013333 ms
[     UNARY] - 0.411000 ms - 84 - 0.004893 ms
Total Time: 407.731049 ms

@leejet Can you give me permission to convert the classes into structs to add support as a dynamic library and also make it easier to use in CLI, server, and other potential uses?

I believe this can be achieved by adding an additional stable-diffusion-c.h/stable-diffusion-c.cpp to export C APIs, much like what was done in this PR: https://github.com/leejet/stable-diffusion.cpp/pull/65/files. This is a common approach for many libraries written in C++. I plan to take a look at merging this PR soon. I prefer not to directly change the class in stable-diffusion.cpp to struct.

Now, stable diffusion supports the gguf format, which will make it easier for me to implement the ggml backend v2 and see if the memory management issue in CUDA is resolved.

Great job! Glad to hear this news.

I briefly reviewed the GGUF section and noticed that you've made some changes to the names of certain tensors. Unless necessary, I recommend keeping the original names of tensors, specifically the ones from the original stable diffusion. The inconsistency between the names in diffusers and the original stable diffusion tensors has been quite confusing. To keep things simple, when I initially implemented stable-diffusion.cpp, I chose to retain the original tensor names from stable diffusion. If there's no compelling reason to rename them, I suggest avoiding these changes.

Sorry for the delayed response, I've been busy with work recently.

I chose to retain the original tensor names from stable diffusion. If there's no compelling reason to rename them, I suggest avoiding these changes.

The thing is, by default, with the gguf format, when saving tensors, there is a maximum character limit in the name, which is 64. Could you change GGML_MAX_NAME and still support the original tensor names? Convert already handles the conversion of HF model names to the original ones and shortens the tensor names, so there's no need to worry about that.

I downloaded the LCM_Dreamshaper_v7_4k.safetensors model and it seems to be just the UNet model (I thought it was the complete model; they don't even bother to deliver things properly.)

In fact, for many diffusers formatted model weights, the clip/unet/vae weights are stored separately.
You can download the following three model files:

https://huggingface.co/SimianLuo/LCM_Dreamshaper_v7/blob/main/text_encoder/model.safetensors
https://huggingface.co/SimianLuo/LCM_Dreamshaper_v7/blob/main/unet/diffusion_pytorch_model.safetensors
https://huggingface.co/SimianLuo/LCM_Dreamshaper_v7/blob/main/vae/diffusion_pytorch_model.safetensors

Then, place them in the following directory structure:
- model_path
    - text_encoder/model.safetensors
    - unet/diffusion_pytorch_model.safetensors
    - vae/diffusion_pytorch_model.safetensors

After that, use this script to convert the safetensors from the diffusers format to the original stable diffusion weights.

@leejet

The thing is, by default, with the gguf format, when saving tensors, there is a maximum character limit in the name, which is 64. I could change GGML_MAX_NAME and still support the original tensor names. convert program already handles the conversion of HuggingFace model names to the original ones and shortens the tensor names, so there's no need to worry about that.

What do you think?

@leejet LCM working in this PR without any python dependency:
C:\proyects\stable-diffusion.cpp>build_cuda\bin\Release\sd -m lcm_model-f16.gguf -p "Self-portrait oil painting, a beautiful cyborg with golden hair, 8k" --steps 4 --cfg-scale 1.0 --sampling-method lcm -t 6
[INFO]  stable-diffusion.cpp:3447 - loading model from 'lcm_model-f16.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3050 Laptop GPU, compute capability 8.6
[INFO]  stable-diffusion.cpp:3475 - Stable Diffusion 1.x | lcm_model
[INFO]  stable-diffusion.cpp:3484 - model data type: f16
[INFO]  stable-diffusion.cpp:3650 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3652 - loading model from 'lcm_model-f16.gguf' completed, taking 0.70s
[INFO]  stable-diffusion.cpp:3676 - running in eps-prediction mode
[INFO]  stable-diffusion.cpp:4440 - get_learned_condition completed, taking 15 ms
[INFO]  stable-diffusion.cpp:4454 - start sampling
[INFO]  stable-diffusion.cpp:4265 - sampling using LCM method
[INFO]  stable-diffusion.cpp:3907 - step 1 sampling completed taking 1.09s
[INFO]  stable-diffusion.cpp:3907 - step 2 sampling completed taking 0.98s
[INFO]  stable-diffusion.cpp:3907 - step 3 sampling completed taking 0.97s
[INFO]  stable-diffusion.cpp:3907 - step 4 sampling completed taking 0.98s
[INFO]  stable-diffusion.cpp:4459 - sampling completed, taking 4.09s
[INFO]  stable-diffusion.cpp:4470 - decode_first_stage completed, taking 25.44s
[INFO]  stable-diffusion.cpp:4478 - txt2img completed in 29.55s
save result image to 'output.png'

Result:

@leejet LCM working in this PR without any python dependency:
C:\proyects\stable-diffusion.cpp>build_cuda\bin\Release\sd -m lcm_model-f16.gguf -p "Self-portrait oil painting, a beautiful cyborg with golden hair, 8k" --steps 4 --cfg-scale 1.0 --sampling-method lcm -t 6
[INFO]  stable-diffusion.cpp:3447 - loading model from 'lcm_model-f16.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3050 Laptop GPU, compute capability 8.6
[INFO]  stable-diffusion.cpp:3475 - Stable Diffusion 1.x | lcm_model
[INFO]  stable-diffusion.cpp:3484 - model data type: f16
[INFO]  stable-diffusion.cpp:3650 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3652 - loading model from 'lcm_model-f16.gguf' completed, taking 0.70s
[INFO]  stable-diffusion.cpp:3676 - running in eps-prediction mode
[INFO]  stable-diffusion.cpp:4440 - get_learned_condition completed, taking 15 ms
[INFO]  stable-diffusion.cpp:4454 - start sampling
[INFO]  stable-diffusion.cpp:4265 - sampling using LCM method
[INFO]  stable-diffusion.cpp:3907 - step 1 sampling completed taking 1.09s
[INFO]  stable-diffusion.cpp:3907 - step 2 sampling completed taking 0.98s
[INFO]  stable-diffusion.cpp:3907 - step 3 sampling completed taking 0.97s
[INFO]  stable-diffusion.cpp:3907 - step 4 sampling completed taking 0.98s
[INFO]  stable-diffusion.cpp:4459 - sampling completed, taking 4.09s
[INFO]  stable-diffusion.cpp:4470 - decode_first_stage completed, taking 25.44s
[INFO]  stable-diffusion.cpp:4478 - txt2img completed in 29.55s
save result image to 'output.png'

Result:

that's amazing!

If someone want to test this:

sadly:
[INFO]  stable-diffusion.cpp:3447 - loading model from '4721097975058205c4edcdece2cc574b7dd7bc04-f16.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3050 Ti Laptop GPU, compute capability 8.6
[INFO]  stable-diffusion.cpp:3475 - Stable Diffusion 1.x |<PHONE_NUMBER>058205c4edcdece2cc574b7dd7bc04
[INFO]  stable-diffusion.cpp:3484 - model data type: f16
[INFO]  stable-diffusion.cpp:3650 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3652 - loading model from '4721097975058205c4edcdece2cc574b7dd7bc04-f16.gguf' completed, taking 0.56s
[INFO]  stable-diffusion.cpp:3676 - running in eps-prediction mode
Segmentation fault

sadly:

Linux, Windows? Could you provide more details, if it worked before or with another model, only when you used the LCM sampling?
Linux, Windows? Could you provide more details, if it worked before or with another model, only when you used the LCM sampling?

@leejet

The thing is, by default, with the gguf format, when saving tensors, there is a maximum character limit in the name, which is 64. I could change GGML_MAX_NAME and still support the original tensor names. convert program already handles the conversion of HuggingFace model names to the original ones and shortens the tensor names, so there's no need to worry about that.

What do you think?

I prefer to keep the original tensor name by modifying GGML_MAX_NAME.

I prefer to keep the original tensor name by modifying GGML_MAX_NAME.

Ok, I will do it now

just it's the name isn't necessary add a '/' in the model path
bin/convert ../models/LCM_Dreamshaper_v7 --type f16

#Custom VAE support when mergin model

Original VAE
Custom VAE

output

@FSSRepo thanks, very nice. Not sure if this is a common use case, but I would like to generate several images after initializing sd once.
For that it would be nice, if speed_up does not depend on free_params_immediately (because the parameters are needed for the next image). And if fallback = false; is set with each StableDiffusion::txt2img(); (it crashes, if still set to true from the last image).
And ggml_backend_buffer_free(buf_compute); added at the end of get_first_stage_encoding(ggml_context* draft_ctx, ggml_tensor* moments) releases memory (line 4408 in stable-diffusion.cpp).

@FSSRepo thanks, very nice. Not sure if this is a common use case, but I would like to generate several images after initializing sd once. For that it would be nice, if speed_up does not depend on free_params_immediately (because the parameters are needed for the next image). And if fallback = false; is set with each StableDiffusion::txt2img(); / StableDiffusion::img2img(); (it crashes, if still set to true from the last image). And ggml_backend_buffer_free(buf_compute); added at the end of get_first_stage_encoding(ggml_context* draft_ctx, ggml_tensor* moments) releases memory (line 4408 in stable-diffusion.cpp):
        ggml_allocr_free(allocr);
        int64_t t1 = ggml_time_ms();
        LOG_DEBUG("computing vae [mode: %s] graph completed, taking %.2fs", decode ? "DECODE" : "ENCODE", (t1 - t0) * 1.0f / 1000);
        result = ggml_dup_tensor(draft_ctx, res);
        copy_ggml_tensor(result, res);
        ggml_backend_buffer_free(buf_compute);
        return result;

Don't worry, I'm taking that into account. I'm going to add an option to generate multiple images once the AutoEncoder is fully on the GPU.

The group norm kernel is really poorly optimized. I've tried to make it as parallel as possible, but it produces invalid values even though I'm careful about thread synchronization.
========== CUDA Timings =========
[       ADD] - 21.414000 ms - 419 - 0.051107 ms
[       MUL] - 8.814000 ms - 125 - 0.070512 ms
[    CONCAT] - 1.504000 ms - 12 - 0.125333 ms
[      NORM] - 2.605000 ms - 48 - 0.054271 ms
[GROUP_NORM] - 434.403015 ms - 61 - 7.121361 ms
[   MUL_MAT] - 116.302002 ms - 361 - 0.322166 ms
[     SCALE] - 2.485000 ms - 32 - 0.077656 ms
[      CONT] - 10.784000 ms - 161 - 0.066981 ms
[  SOFT_MAX] - 99.375000 ms - 32 - 3.105469 ms
[    IM2COL] - 231.380005 ms - 97 - 2.385361 ms
[   UPSCALE] - 0.286000 ms - 3 - 0.095333 ms
[     UNARY] - 4.634000 ms - 84 - 0.055167 ms
Total Time: 933.985962 ms
========== CUDA Timings =========
[       ADD] - 21.429001 ms - 419 - 0.051143 ms
[       MUL] - 8.772000 ms - 125 - 0.070176 ms
[    CONCAT] - 1.473000 ms - 12 - 0.122750 ms
[      NORM] - 2.509000 ms - 48 - 0.052271 ms
[GROUP_NORM] - 434.075012 ms - 61 - 7.115984 ms
[   MUL_MAT] - 116.774002 ms - 361 - 0.323474 ms
[     SCALE] - 2.521000 ms - 32 - 0.078781 ms
[      CONT] - 10.904000 ms - 161 - 0.067727 ms
[  SOFT_MAX] - 99.410004 ms - 32 - 3.106563 ms
[    IM2COL] - 231.143997 ms - 97 - 2.382928 ms
[   UPSCALE] - 0.286000 ms - 3 - 0.095333 ms
[     UNARY] - 4.606000 ms - 84 - 0.054833 ms
Total Time: 933.902954 ms
[INFO]  stable-diffusion.cpp:3892 - step 2 sampling completed taking 1.87s

I randomly remembered xaedes talking about optimizing the graph order using genetic algorithm https://github.com/ggerganov/llama.cpp/pull/2632#issuecomment-1714366066

@leejet Added LoRA support in this PR
C:\proyects\stable-diffusion.cpp>build_cuda\bin\Release\sd -m v1-5-pruned-emaonly-f16.gguf -v -p "<lora:lcm_lora:1>a lovely cat" --steps 8 --cfg-scale 1
Option:
    n_threads:       6
    mode:            txt2img
    model_path:      v1-5-pruned-emaonly-f16.gguf
    output_path:     output.png
    init_img:
    prompt:          <lora:lcm_lora:1>a lovely cat
    negative_prompt:
    cfg_scale:       1.00
    width:           512
    height:          512
    sample_method:   euler_a
    schedule:        default
    sample_steps:    8
    strength:        0.75
    rng:             cuda
    seed:            42
System Info:
    BLAS = 1
    SSE3 = 1
    AVX = 1
    AVX2 = 1
    AVX512 = 0
    AVX512_VBMI = 0
    AVX512_VNNI = 0
    FMA = 1
    NEON = 0
    ARM_FMA = 0
    F16C = 1
    FP16_VA = 0
    WASM_SIMD = 0
    VSX = 0
[INFO]  stable-diffusion.cpp:3671 - loading model from 'v1-5-pruned-emaonly-f16.gguf'
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3050 Laptop GPU, compute capability 8.6
[DEBUG] stable-diffusion.cpp:3689 - load_from_file: - kv   0:                              sd.model.name str
[DEBUG] stable-diffusion.cpp:3689 - load_from_file: - kv   1:                             sd.model.dtype i32
[DEBUG] stable-diffusion.cpp:3689 - load_from_file: - kv   2:                           sd.model.version i8
[DEBUG] stable-diffusion.cpp:3689 - load_from_file: - kv   3:                            sd.vocab.tokens arr
[INFO]  stable-diffusion.cpp:3699 - Stable Diffusion 1.x | v1-5-pruned-emaonly.ckpt
[INFO]  stable-diffusion.cpp:3708 - model data type: f16
[DEBUG] stable-diffusion.cpp:3712 - loading vocab
[DEBUG] stable-diffusion.cpp:3728 - ggml tensor size = 416 bytes
[DEBUG] stable-diffusion.cpp:812  - Using CUDA backend - clip
[DEBUG] stable-diffusion.cpp:823  - clip params backend buffer size =  236.18 MB
[DEBUG] stable-diffusion.cpp:826  - clip tensor count = 449
[DEBUG] stable-diffusion.cpp:1956 - Using CUDA backend - unet
[DEBUG] stable-diffusion.cpp:1968 - unet params backend buffer size =  1641.16 MB
[DEBUG] stable-diffusion.cpp:1972 - unet tensor count = 706
[DEBUG] stable-diffusion.cpp:3063 - Using CUDA backend - vae
[DEBUG] stable-diffusion.cpp:3074 - vae params backend buffer size =  95.47 MB
[DEBUG] stable-diffusion.cpp:3083 - vae tensor count = 164
[DEBUG] stable-diffusion.cpp:3737 - preparing memory for the weights
[DEBUG] stable-diffusion.cpp:907  - clip params allocated
[DEBUG] stable-diffusion.cpp:2075 - unet params allocated
[DEBUG] stable-diffusion.cpp:3127 - vae params allocated
[DEBUG] stable-diffusion.cpp:3755 - loading weights
[DEBUG] stable-diffusion.cpp:3859 - model size = 1969.67MB
[INFO]  stable-diffusion.cpp:3869 - total memory buffer size = 1972.80MB (clip 236.18MB, unet 1641.16MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3871 - loading model from 'v1-5-pruned-emaonly-f16.gguf' completed, taking 0.76s
[INFO]  stable-diffusion.cpp:3895 - running in eps-prediction mode
[DEBUG] stable-diffusion.cpp:3922 - finished loaded file
[DEBUG] stable-diffusion.cpp:3930 - applying LoRA - lcm_lora: 1.00
[INFO]  stable-diffusion.cpp:3262 - loading LoRA from 'lcm_lora.gguf'
[DEBUG] stable-diffusion.cpp:3281 - load: - kv   0:                               sd.lora.name str
[DEBUG] stable-diffusion.cpp:3281 - load: - kv   1:                              sd.lora.dtype i32
[DEBUG] stable-diffusion.cpp:3281 - load: - kv   2:                               sd.lora.type i32
[DEBUG] stable-diffusion.cpp:3281 - load: - kv   3:                           sd.lora.alphas_k arr
[DEBUG] stable-diffusion.cpp:3281 - load: - kv   4:                           sd.lora.alphas_v arr
[INFO]  stable-diffusion.cpp:3288 - LoRA Type: 2 | lcm_lora.safetensors
[DEBUG] stable-diffusion.cpp:3293 - Using CUDA backend - lora
[INFO]  stable-diffusion.cpp:3317 - LoRA data type: f16
[DEBUG] stable-diffusion.cpp:3321 - calculating buffer size
[DEBUG] stable-diffusion.cpp:3329 - lora params backend buffer size =  128.27 MB
[DEBUG] stable-diffusion.cpp:3333 - loading alphas
[DEBUG] stable-diffusion.cpp:3388 - finished loaded lora
[DEBUG] stable-diffusion.cpp:3468 - apply lora buffer size: 227.85 MB
[DEBUG] stable-diffusion.cpp:3482 - computing lora
[DEBUG] stable-diffusion.cpp:3485 - free lora
[DEBUG] stable-diffusion.cpp:1059 - parse 'a lovely cat' to [['a lovely cat', 1], ]
[DEBUG] stable-diffusion.cpp:451  - split prompt "a lovely cat" to tokens ["a</w>", "lovely</w>", "cat</w>", ]
[DEBUG] stable-diffusion.cpp:993  - learned condition compute buffer size: 1.58 MB
[DEBUG] stable-diffusion.cpp:3995 - computing condition graph completed, taking 7 ms
[INFO]  stable-diffusion.cpp:4655 - get_learned_condition completed, taking 9 ms
[INFO]  stable-diffusion.cpp:4669 - start sampling
[DEBUG] stable-diffusion.cpp:2334 - diffusion compute buffer size: 552.57 MB
[INFO]  stable-diffusion.cpp:4147 - sampling using Euler A method
[INFO]  stable-diffusion.cpp:4140 - step 1 sampling completed taking 0.94s
[INFO]  stable-diffusion.cpp:4140 - step 2 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 3 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 4 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 5 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 6 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 7 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4140 - step 8 sampling completed taking 0.93s
[INFO]  stable-diffusion.cpp:4674 - sampling completed, taking 7.55s
[DEBUG] stable-diffusion.cpp:3218 - vae compute buffer size: 1664.00 MB
[DEBUG] stable-diffusion.cpp:4594 - computing vae [mode: DECODE] graph completed, taking 0.05s
[INFO]  stable-diffusion.cpp:4685 - decode_first_stage completed, taking 6.32s
[INFO]  stable-diffusion.cpp:4692 - txt2img completed in 13.88s
save result image to 'output.png'

Result:

I'm in the process of rebasing this branch onto my latest code; it's almost done. If you don't mind, once I'm done, I'll force push to your master branch.

@leejet Before merging this pull request, I need to add documentation and remove some things that I had used for debugging.Try not to overwrite my ggml submodule.

I won't change ggml when I rebase your branch

@FSSRepo something is off with your lcm results, its too sharp. + looks like you did not use the lcm sampler (minor difference)

@FSSRepo something is off with your lcm results, its too sharp. + looks like you did not use the lcm sampler (minor difference)

No, I already did the test on master, and I get the same result, so there are no differences.
This is the result with master:

@FSSRepo I've finished the rebase and I'm going to force push to your branch.

@leejet 💀🤣🤣🤣

Don't worry, I backed it up in another branch

Done. I tested both CPU and GPU and it worked fine. I've tested it on both linux and windows, and will probably test it on macos again sometime.

CPU backend
CUDA backend

Thank you very much. Could you please refrain from making any further changes? I'm going to start cleaning up some things now. I'll let you know when everything is ready

I'll let you know if I have any breaking changes, and I'll try to avoid making new commits until this pr is merged.

You can try to create a new branch and pull the remote branch. For example:
git fetch origin
git checkout -b dev origin/master
git submodule update

Remove unused code, update the documentation, and delete the 'models' folder as it will not be necessary from now on since Python is not required. Also, finally, add a memory optimization in U-Net Flash Attention, reducing memory usage by 4 times without degrading performance. Currently, this optimization applies only when using the full CPU backend; it is not applied when using CUDA

I think the Flash Attention memory optimization could be opened as a separate PR. This PR has already accomplished so much excellent work, and I can't wait to merge it!

I think the Flash Attention memory optimization could be opened as a separate PR. This PR has already accomplished so much excellent work, and I can't wait to merge it!

It's not much change, just replace some lines with ggml_flash_attention.

I think CUDA kernel for flash attention will take some time.

I think CUDA kernel for flash attention will take some time.

Yes, I will add it in a pull request after this one to avoid prolonging the wait. For now, "flash attention" will be a default OFF option in CMake. I will display a warning if it's activated using CUDA.

@FSSRepo the convert tool needs lora autodetection. because rn it converts "successfully" as a model file, but asserts during runtime, when not explicitly converted as lora.
Also, the convert tool appends the type to the file, but the lora is looked up by full name + .gguf -> does not find with the -f16. It's not the end of the world, but it's annoying to manually rename every file or add the filetype to the loraname in-prompt.
Also missing, is source value type -> if the src file is in f16 it should autoconvert to f16 etc...

@FSSRepo the convert tool needs lora autodetection. because rn it converts "successfully" as a model file, but asserts during runtime, when not explicitly converted as lora.
Also, the convert tool appends the type to the file, but the lora is looked up by full name + .gguf -> does not find with the -f16. It's not the end of the world, but it's annoying to manually rename every file or add the filetype to the loraname in-prompt.
Also missing, is source value type -> if the src file is in f16 it should autoconvert to f16 etc...
edit: Another point is missing BF16 (conversion) support.

I will fix that, I have researched extensively trying to find a way to convert bfloat16 types to float16 in C++, and I have not found any information.

I will fix that, I have researched extensively trying to find a way to convert bfloat16 types to float16 in C++, and I have not found any information. ChatGPT says me that just cast bfloat16 as float16 w/o conversion.

bf16 is just a cut down f32 (see wikipedia lmao)

$ bin/sd -m ../models/epicphotogasm_lastUnicorn-f16.gguf --cfg-scale 1.0 --sampling-method lcm --lora-model-dir ../models/loras/sd1-gguf/ --steps 8 -p "<lora:lcm-lora-sdv1-5:1><lora:Round_animal-10:1>a lovely cat"

@leejet I haven't tested it, but this pull request (PR) could potentially break compatibility with CLBlast. This is because the implementation of ggml_mul in OpenCL doesn't support broadcasting, and this PR assumes that all backends support broadcasting in ggml_add and ggml_mul.

For now, I'm going to comment out those lines to avoid issues with users.

Ok, please point this out in the readme.

@FSSRepo @leejet when loading a lora with 0.f, it now gets stuck in an endless loop. only tested this pr.

I also noticed: if a lora mismatches the base model in datatype (f16 vs f32), you get a cryptic backtrace about allocations in this branch. That might benefit from a friendlier error message.

@FSSRepo @leejet when loading a lora with 0.f, it now gets stuck in an endless loop. only tested this pr.

It should be fixed now.

I also noticed: if a lora mismatches the base model in datatype (f16 vs f32), you get a cryptic backtrace about allocations in this branch. That might benefit from a friendlier error message.

I definitely need to limit LoRA conversion to only float16 or float32. I've tested LoRAs with quantized models, and degradation becomes noticeable only with q4_0 and q4_1, q5_0 q5_1 and q8_0 is almost the same result of f16.

Trying quantizing a LoRA
C:\proyects\finish-pr\stable-diffusion.cpp\build\bin\Release>convert.exe glasses.safetensors -tp q4_0
loading model 'glasses.safetensors'
model type: safetensors
LoRA detected
Error: The LoRa conversion only supports f32 and f16.

Using f16 lora with model quantized f16, q5_1, q4_0:
C:\proyects\finish-pr\stable-diffusion.cpp>build\bin\Release\sd -m AnythingV5_v5PrtRE-q4_0.gguf -p "<lora:Kana_Arima-10:0.95>beautiful anime girl, red hair, short hair, (red eyes), cute, casual dress, masterpiece, 4k, high quality, smile" -n "bad quality, ugly, face malformed, bad anatomy" --sampling-method dpm++2m --steps 20 -t 6 -s 55543534
[INFO]  stable-diffusion.cpp:3695 - loading model from 'AnythingV5_v5PrtRE-q4_0.gguf'
[INFO]  stable-diffusion.cpp:3723 - Stable Diffusion 1.x | AnythingV5_v5PrtRE.safetensors
[INFO]  stable-diffusion.cpp:3731 - model data type: q4_0
[INFO]  stable-diffusion.cpp:3892 - total memory buffer size = 1434.05MB (clip 67.64MB, unet 1270.95MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3894 - loading model from 'AnythingV5_v5PrtRE-q4_0.gguf' completed, taking 1.05s
[INFO]  stable-diffusion.cpp:3918 - running in eps-prediction mode
[WARN]  stable-diffusion.cpp:4004 - In quantized models when applying LoRA, the images have poor quality.
[INFO]  stable-diffusion.cpp:3277 - loading LoRA from 'Kana_Arima-10.gguf'
[INFO]  stable-diffusion.cpp:3303 - LoRA Type: 2 | Kana_Arima-10.safetensors
[INFO]  stable-diffusion.cpp:3323 - LoRA data type: f16
[INFO]  stable-diffusion.cpp:3999 - lora 'Kana_Arima-10' applied, taking 1.99s
[INFO]  stable-diffusion.cpp:4653 - apply_loras completed, taking 1.99s
[INFO]  stable-diffusion.cpp:4662 - get_learned_condition completed, taking 226 ms
[INFO]  stable-diffusion.cpp:4676 - start sampling
[INFO]  stable-diffusion.cpp:4417 - sampling using DPM++ (2M) method
|==================================================| 20/20 - 19.05s/it
[INFO]  stable-diffusion.cpp:4681 - sampling completed, taking 381.45s
[INFO]  stable-diffusion.cpp:4692 - decode_first_stage completed, taking 27.96s
[INFO]  stable-diffusion.cpp:4699 - txt2img completed in 409.63s
save result image to 'output.png'

Result:

f16
q5_1
q4_0

if a lora mismatches the base model in datatype (f16 vs f32), you get a cryptic backtrace about allocations in this branch.

@ursg  Please explain more

@FSSRepo progress pretty printing is not working properly yet. It is only displayed, once completed, so I suspect missing flushes.

@FSSRepo progress pretty printing is not working properly yet. It is only displayed, once completed, so I suspect missing flushes.

Linux?

@FSSRepo progress pretty printing is not working properly yet. It is only displayed, once completed, so I suspect missing flushes.

Linux? If you can fix it, Let me know it
// add this after printf(time > 1.0f ? "| %i/%i - ...
fflush(stdout);

I came to the same solution.

I came to the same solution. and it works.

Great, i will fix it now

Using f16 lora with model quantized f16, q5_1, q4_0:
C:\proyects\finish-pr\stable-diffusion.cpp>build\bin\Release\sd -m AnythingV5_v5PrtRE-q4_0.gguf -p "<lora:Kana_Arima-10:0.95>beautiful anime girl, red hair, short hair, (red eyes), cute, casual dress, masterpiece, 4k, high quality, smile" -n "bad quality, ugly, face malformed, bad anatomy" --sampling-method dpm++2m --steps 20 -t 6 -s 55543534
[INFO]  stable-diffusion.cpp:3695 - loading model from 'AnythingV5_v5PrtRE-q4_0.gguf'
[INFO]  stable-diffusion.cpp:3723 - Stable Diffusion 1.x | AnythingV5_v5PrtRE.safetensors
[INFO]  stable-diffusion.cpp:3731 - model data type: q4_0
[INFO]  stable-diffusion.cpp:3892 - total memory buffer size = 1434.05MB (clip 67.64MB, unet 1270.95MB, vae 95.47MB)
[INFO]  stable-diffusion.cpp:3894 - loading model from 'AnythingV5_v5PrtRE-q4_0.gguf' completed, taking 1.05s
[INFO]  stable-diffusion.cpp:3918 - running in eps-prediction mode
[WARN]  stable-diffusion.cpp:4004 - In quantized models when applying LoRA, the images have poor quality.
[INFO]  stable-diffusion.cpp:3277 - loading LoRA from 'Kana_Arima-10.gguf'
[INFO]  stable-diffusion.cpp:3303 - LoRA Type: 2 | Kana_Arima-10.safetensors
[INFO]  stable-diffusion.cpp:3323 - LoRA data type: f16
[INFO]  stable-diffusion.cpp:3999 - lora 'Kana_Arima-10' applied, taking 1.99s
[INFO]  stable-diffusion.cpp:4653 - apply_loras completed, taking 1.99s
[INFO]  stable-diffusion.cpp:4662 - get_learned_condition completed, taking 226 ms
[INFO]  stable-diffusion.cpp:4676 - start sampling
[INFO]  stable-diffusion.cpp:4417 - sampling using DPM++ (2M) method
|==================================================| 20/20 - 19.05s/it
[INFO]  stable-diffusion.cpp:4681 - sampling completed, taking 381.45s
[INFO]  stable-diffusion.cpp:4692 - decode_first_stage completed, taking 27.96s
[INFO]  stable-diffusion.cpp:4699 - txt2img completed in 409.63s
save result image to 'output.png'

Result:
f16	q5_1	q4_0

Where you get "AnythingV5_v5PrtRE-q4_0.gguf"? Can you give me the link to i use here? Thanks.

If I compile on windows the `stable-diffusion.lib` does not seem to be created:
  ggml.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\ggml.dll
  stable-diffusion.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\stable-diffusion.dll
  common.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\common\common.dir\Release\common.lib
  zip.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\common\zip.dir\Release\zip.lib
  convert.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\convert.exe
LINK : fatal error LNK1181: file "..\..\Release\stable-diffusion.lib" can not opened. [C:\Users\Jon
at\Desktop\cp\stable-diffusion.cpp\build\examples\cli\sd.vcxproj]

@paulocoutinhox

Where you get "AnythingV5_v5PrtRE-q4_0.gguf"? Can you give me the link to i use here? Thanks.

download the original weights (probably safetensors) and convert using the convert executable.

Ok. Another thing:
Im trying use llama and stable-diffusion in the same EXE, but ggml.h from llama don't have "ggml_get_dynamic, ggml_init_params.dynamic, ggml_set_dynamic, ggml_used_mem_of_data, etc".
Can you update to latest llama? Thanks.

Can you update to latest llama/ggml? Thanks.

A lot of ops that sd.cpp uses doesn't available in the ggml original project.

@Jonathhhan I don't get some error when I build sd.cpp as shared libraries.

@paulocoutinhox

Where you get "AnythingV5_v5PrtRE-q4_0.gguf"? Can you give me the link to i use here? Thanks.

Can you short your comment, delete the images of quote. Here is the link of the model you can convert it to gguf:
convert AnythingV5_v5PrtRE.safetensors -tp f16

It seems this PR is ready for merging. I'll review and merge it within the next couple of days.

@FSSRepo This would be great if you could submit a PR for the optimizations made to ggml to ggerganov/ggml. This way, more projects based on ggml could benefit from these optimizations.

Has anyone tried this version on Metal? Are there any missing op kernels that would need to be implemented?

@ggerganov You can refer to the operations table that I have in the description of my pull request; I have described all the operations that require a kernel in Metal there.

If I compile on windows the stable-diffusion.lib does not seem to be created:
  ggml.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\ggml.dll
  stable-diffusion.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\stable-diffusion.dll
  common.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\common\common.dir\Release\common.lib
  zip.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\common\zip.dir\Release\zip.lib
  convert.vcxproj -> C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\build\bin\Release\convert.exe
LINK : fatal error LNK1181: file "..\..\Release\stable-diffusion.lib" can not opened. [C:\Users\Jon
at\Desktop\cp\stable-diffusion.cpp\build\examples\cli\sd.vcxproj]

Edit: stable-diffusion.lib gets created , if I build with -DBUILD_SHARED_LIBS=OFF.

@Jonathhhan I don't get some error when I build sd.cpp as shared libraries.

Apologies for any confusion. Obviously I do something wrong.
What I want to do, is to wrap stable-diffusion.cpp as a library for an (my first) Open Frameworks addon (still in progress, but it already works): https://github.com/Jonathhhan/ofxStableDiffusion

@leejet  I've been trying to add the batch count option to generate one after another, but only the first image is generated correctly. For some reason, on the next one, the parameter tensor data is erased or has invalid references when they were not released or deleted from memory at any point. This happens on both CPU and CUDA, although with different errors each time I try.

@Jonathhhan Nothing related to that commit, in the example the CLI option to release the parameters immediately was activated; I have already managed to solve it.

@FSSRepo glad that it is working now. And sorry for my wrong guesses.

@FSSRepo glad that it is working now. And sorry for my wrong guesses.

Translate to English: Don't worry, it gave me a clue when I took a moment to think.

@FSSRepo the timer for sampling is not reset between images, so sampling takes longer and longer. not sure how to properly display this.
[INFO]  stable-diffusion.cpp:4685 - generating image: 1/10
[INFO]  stable-diffusion.cpp:4693 - start sampling
[DEBUG] stable-diffusion.cpp:2366 - diffusion compute buffer size: 115.28 MB
[INFO]  stable-diffusion.cpp:4511 - sampling using LCM method
  |==================================================| 26/26 - 4.16s/it
[INFO]  stable-diffusion.cpp:4698 - sampling completed, taking 107.90s
[INFO]  stable-diffusion.cpp:4685 - generating image: 2/10
[INFO]  stable-diffusion.cpp:4693 - start sampling
[DEBUG] stable-diffusion.cpp:2366 - diffusion compute buffer size: 115.28 MB
[INFO]  stable-diffusion.cpp:4511 - sampling using LCM method
  |==================================================| 26/26 - 4.33s/it
[INFO]  stable-diffusion.cpp:4698 - sampling completed, taking 218.30s
[INFO]  stable-diffusion.cpp:4685 - generating image: 3/10
[INFO]  stable-diffusion.cpp:4693 - start sampling
[DEBUG] stable-diffusion.cpp:2366 - diffusion compute buffer size: 115.28 MB
[INFO]  stable-diffusion.cpp:4511 - sampling using LCM method
  |==================================================| 26/26 - 4.19s/it
[INFO]  stable-diffusion.cpp:4698 - sampling completed, taking 329.84s
[INFO]  stable-diffusion.cpp:4685 - generating image: 4/10
[INFO]  stable-diffusion.cpp:4693 - start sampling
[DEBUG] stable-diffusion.cpp:2366 - diffusion compute buffer size: 115.28 MB
[INFO]  stable-diffusion.cpp:4511 - sampling using LCM method
  |==================================================| 26/26 - 4.11s/it
[INFO]  stable-diffusion.cpp:4698 - sampling completed, taking 438.95s
[INFO]  stable-diffusion.cpp:4685 - generating image: 5/10
[INFO]  stable-diffusion.cpp:4693 - start sampling

One batch generation is working well for me, but if I start a second batch (with size 1, too) the app crashes and I get this error message:
[DEBUG] stable-diffusion.cpp:1101 - parse 'a lovely cat, van gogh' to [['a lovely cat, van gogh', 1], ]
[DEBUG] stable-diffusion.cpp:503  - split prompt "a lovely cat, van gogh" to tokens ["a</w>", "lovely</w>", "cat</w>", ",</w>", "van</w>", "gogh</w>", ]
[DEBUG] stable-diffusion.cpp:1034 - learned condition compute buffer size: 1.58 MB
[DEBUG] stable-diffusion.cpp:4049 - computing condition graph completed, taking 41 ms
[DEBUG] stable-diffusion.cpp:1101 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:503  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:1034 - learned condition compute buffer size: 1.58 MB
[DEBUG] stable-diffusion.cpp:4049 - computing condition graph completed, taking 38 ms
[INFO]  stable-diffusion.cpp:4673 - get_learned_condition completed, taking 81 ms
[INFO]  stable-diffusion.cpp:4685 - generating image: 1/4
[INFO]  stable-diffusion.cpp:4693 - start sampling
GGML_ASSERT: /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml.c:4065: ggml_can_mul_mat(a, b)

CUDA error 4 at /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml-cuda.cu:8550: driver shutting down
current device: 0
Segmentation fault
make: *** [../../../libs/openFrameworksCompiled/project/makefileCommon/compile.project.mk:191: RunRelease] Error 139

Anyway, it only happens in my Open Frameworks implementation (but with Linux and Windows)...

libstable-diffusion.so points to the path where libggml.so was created.
Its possible to change that to the relative path with set_target_properties in this file https://github.com/FSSRepo/stable-diffusion.cpp/blob/master/CMakeLists.txt:
add_library(${SD_LIB} stable-diffusion.h stable-diffusion.cpp)
set_target_properties(${SD_LIB} PROPERTIES
        BUILD_WITH_INSTALL_RPATH FALSE
        LINK_FLAGS "-Wl,-rpath,$ORIGIN/")
target_link_libraries(${SD_LIB} PUBLIC ggml)
target_include_directories(${SD_LIB} PUBLIC .)
target_compile_features(${SD_LIB} PUBLIC cxx_std_11)

@FSSRepo yes, it is still generated in the same directory. The difference is, that if I run the program stable-diffusion.so links to ggml.so in the same folder (wherever it is). While before stable-diffusion.so was linked to the ggml.so file in the folder where it was created.

@FSSRepo yes, it is still generated in the same directory. The difference is, if I run the program with the change stable-diffusion.so links to ggml.so in the same folder (wherever it is). While before stable-diffusion.so was linked to the ggml.so file in the folder where it was created.

this fixes your problem?
GGML_ASSERT: /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml.c:4065: ggml_can_mul_mat(a, b)

CUDA error 4 at /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml-cuda.cu:8550: driver shutting down
current device: 0
Segmentation fault
make: *** [../../../libs/openFrameworksCompiled/project/makefileCommon/compile.project.mk:191: RunRelease] Error 139

@FSSRepo no, its a different one. Maybe that is (really no idea), because the buffers/contexts are freed after each batch and they only should be freed if the program is closed or recreated if a new model is loaded?

StableDiffusion::txt2img works very well now. With StableDiffusion::img2img I get this error:
[DEBUG] stable-diffusion.cpp:1112 - parse 'a lovely cat, van gogh' to [['a lovely cat, van gogh', 1], ]
[DEBUG] stable-diffusion.cpp:514  - split prompt "a lovely cat, van gogh" to tokens ["a</w>", "lovely</w>", "cat</w>", ",</w>", "van</w>", "gogh</w>", ]
[DEBUG] stable-diffusion.cpp:1045 - learned condition compute buffer size: 1.58 MB
[DEBUG] stable-diffusion.cpp:4060 - computing condition graph completed, taking 6 ms
[DEBUG] stable-diffusion.cpp:1112 - parse '' to [['', 1], ]
[DEBUG] stable-diffusion.cpp:514  - split prompt "" to tokens []
[DEBUG] stable-diffusion.cpp:1045 - learned condition compute buffer size: 1.58 MB
[DEBUG] stable-diffusion.cpp:4060 - computing condition graph completed, taking 5 ms
[INFO]  stable-diffusion.cpp:4804 - get_learned_condition completed, taking 12 ms
[INFO]  stable-diffusion.cpp:4813 - sampling using modified DPM++ (2M) method
GGML_ASSERT: C:\Users\Jonat\Desktop\cp\stable-diffusion.cpp\ggml\src\ggml.c:5253: a->ne[2] == b->ne[2]

It was working some commits before, but with some issue: https://github.com/leejet/stable-diffusion.cpp/issues/79

With StableDiffusion::img2img I get:
[INFO]  stable-diffusion.cpp:4813 - sampling using modified DPM++ (2M) method
GGML_ASSERT: /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml.c:5253: a->ne[2] == b->ne[2]

CUDA error 4 at /mnt/c/Users/Jonat/Desktop/cp/stable-diffusion.cpp/ggml/src/ggml-cuda.cu:8550: driver shutting down
current device: 0
Segmentation fault

@Jonathhhan check now

@FSSRepo it works. Thanks. I only have this issue: https://github.com/leejet/stable-diffusion.cpp/issues/79
But maybe its because of the model. I will try this one: https://huggingface.co/stabilityai/stable-diffusion-2-1-unclip

Hi, one question. It works on iOS?

@paulocoutinhox Yes, it can be done, but it requires creating the missing kernels for the Metal backend based on the ones I have created for cuBLAS. I don't have a Mac or iPhone, and I don't know much about Metal programming, so I am the least qualified to contribute anything.

Nice. But i run i run today in one iPhone, what happen? It will run on CPU?

Nice. But i run i run today in one iPhone, what happen? It will run on CPU?

yes

I've tested it on Windows/Linux/macOS, and now it's working perfectly. The PR is ready for merging. Thank you for your excellent work!