Block Tiling Using Shared Memory

Block Tiling Using Shared Memory
We will use three simplified examples from n-body, lavaMD and Matrix-Matrix Multiplication benchmarks. The notation of streamSeq might differ from what is curently in Futhark -- I keep around the chunk size and the current index k of streamSeq, i.e., the next chunk starts at array element k and ends at array element k+chunk-1.
I assume the analysis is run on extracted kernels, i.e., which can be seen as a perfect map nest that is going to be executed in parallel on GPU. I also assume that map, redomap, scan, filter, streamMap/Red, etc., appearing inside the kernel have been turned to StreamSeq constructs.
1. How to Recognize the Beneficial Cases for Block Tiling?
In essence block-tiling is beneficial whenever we can detect a streamSeq whose input array is invariant to one of the parallel dimensions of the kernel.
(A) n-body
fun [f32, n] main(f32 eps, [f32, n] bodies) =
map( fn f32 (f32 body) =>  -- 1-dim Kernel
            streamSeq( fn f32 (int kk, [f32, chunk] bodies_chunk, f32 acc) =>
                            loop (acc) = for k < chunk do
                                acc + (bodies_chunk[i] - body)
                            in acc
                     , bodies, 0f32)
   , bodies)

n-body is the simplest case demonstrating such an access: I assume the kernel is one-dimensional and contains only the outermost map which is fed array bodies. However, array bodies is also fed to the streamSeq that forms the kernel body. Assuming the one-dimensional block size is TILE (e.g., 128, 256, etc.), one case choose chunk = TILE, and just before the loop make the block threads cooperatively write the current TILE-chunk of bodies-chunk to fast (local/shared) memory such that inside the loop all accesses to bodies_chunk are from fast memory (rather than global memory, even if, in this case, it might be efficiently cached).
(B) lavaMD
fun [[f32, m], n] main([[f32, m], n] AA, [[int, p],n] inds_nhbs) =
map( fn [f32,m] ([f32,m] A, int i) => -- Kernel's first parallel dimension
            map( fn f32 (f32 a) =>         -- Kernel's second parallel dimension
                    streamSeq( fn f32 (int kk1, [int, chunk1] inds_nhbs, f32 acc1) =>
                            loop (acc1) = for k1 < chunk1 do
                                let ind = inds_nhbs[i, k1] in
                                let invar_AArow = AA[ind] in
                                streamSeq( fn f32 (int kk2, [int, chunk2] invar_AArow_chunk, f32 acc2) =>
                                            loop (acc2) = for k2 < chunk2 do
                                                    acc2 + invar_AArow_chunk[k2] * 2.0f32
                                            in acc2
                                         , invar_AArow, acc1)
                            in acc1
                    , iota(inds_nhbs), 0f32)
               , A)
   , zip(AA, iota(n)))

lavaMD presents a slightly more complex case that basically demonstrates that we cannot rely on simple pattern-matching, but rather we need to do a full pass of invariant-array detection.  In the example above invar_AArow is invariant to the second (map) parallel dimension of the kernel. Note that this array is created somewhere deep inside the kernel, but it is fed as input to the second streamSeq and as such can be blocked tiled, i.e., used directly from shared memory.
Note that this requires a 2-D grid, but only a 1-dim block (i.e., the second map is split in one or a number of blocks), and the TILE size can be the dimension of the 1-d block. This would allow to amortize one global-memory access over TILE accesses from fast memory (related to invar_AArow_chunk).
Note however that if the invariant map would have been the first one, block-tiling would have required to interchange the parallel dimensions, i.e., bring it innermost because fast memory is shared only within the same workgroup.
(C) MM-Mult
fun [[f32,q],m] main([[f32,n],m] A, [[f32,q],n] B0) =
  let B = transpose(B0) in
  map(fn [f32,q] ( [f32,n] a ) =>       -- Kernel's first parallel dimension
            map(fn f32 ( [f32,n] b ) =>  -- Kernel's second parallel dimension
                    streamSeq(fn f32 (int kk, [f32,chunk] a_c, [f32,chunk] b_c, int acc) =>
                                  loop (acc) = for k < chunk do
                                          acc + a_c[k]*b_c[k]
                                  in acc
                             , a, b, 0f32 )
               , B)
     , A )

The third example is the trickier one of matrix-matrix multiplication: here a is invariant to the second map, while b is invariant to the first map, and they are both fed to the streamSeq that makes the body. In order to cache both accesses in shared memory, one needs to create a 2-dimensional workgroup (so that accesses on both a and b can be shared).
How to Implement Tiling?
Once tiling opportunities have been detected, block tiling can be implemented by (i) adding a notion of fast (shared/local) memory to the Kernel, by (ii) inserting a special function call that performs the per-thread copying to fast memory just between the invariant streamSeq and its corresponding loop, and, finally, by (iii) replacing accesses to the invariant array with accesses from fast memory.
streamSeq provides a convenient notation for doing this because the indexing into array chunks starts at 0 and goes until chunk, hence exactly how we need to access fast memory! We use here a CUDA notation -- we will have to extend the kernel with some magic things such as a notion of block/workgroup-size (TILE) and indexing within a workgroup (threadIdx.x), and obviously, a notion of fast memory parameterized over the workgroup size, and a way to copy to fast memory, denoted below with copyFastMem.
(A) n-body
For n-body we assume a 1-dim block of size TILE, and as such i corresponds to blockIdx.x*blockDim.x + threadIdx.x. We choose the chunk of the invariant streamSeq to be TILE. The translated code becomes:
fun [f32, n] main(f32 eps, [f32, n] bodies) =
map( fn f32 (int i, f32 body) =>  -- 1-dim Kernel;
                                                  -- 1-dim workgroup/block of size TILE, threadIdx.x in [0,TILE-1]
                                                  -- fast memory: shMem_bodies[TILE]
            streamSeq( fn f32 (int kk, [f32, TILE] bodies_chunk, f32 acc) =>
                            let shMem_bodies = copyFastMem(bodies_chunk[threadIdx.x], threadIdx.x) in
                            loop (acc) = for k < TILE do
                                acc + (shMem_bodies[k] - body)
                            in acc
                     , bodies, 0f32)
   , zip(iota(n), bodies)

Note that with this bit of magic, when we finally translate a streamSeq to a loop we need to obviously extend the indexing to arrays such as bodies_chunk but not to fast-memory arrays, such as shMem_bodies.   Note that copyFastMem should actually check that the index of the global-memory array is in range, and if it is not it should obviously not write to shared memory.
(B) lavaMD
We treat lavaMD similarly, but note that while the grid has semantically a 2-D structure, we need only use a 1-dim structure for the block/workgroup (of size TILE), such as one or multiple blocks cover the inner (second) map. This is potentially subject to multi-versioned code since this makes sense only if the second map is large enough!  We assume for simplicity that TILE divides m, the size of the second map.
fun [[f32, m], n] main([[f32, m], n] AA, [[int, p],n] inds_nhbs) =
map( fn [f32,m] ([f32,m] A, int i) => -- Kernel's first parallel dimension
            map( fn f32 (f32 a) =>         -- Kernel's second parallel dimension
                                                        -- 1-dim Block of size TILE, threadIdx.x in [0,TILE-1]
                                                        -- fast memory: shMem_AArow[TILE]
                    streamSeq( fn f32 (int kk1, [int, chunk1] inds_nhbs, f32 acc1) =>
                            loop (acc1) = for k1 < chunk1 do
                                let ind = inds_nhbs[i, k1] in
                                let invar_AArow = AA[ind] in
                                streamSeq( fn f32 (int kk2, [int, TILE] invar_AArow_chunk, f32 acc2) =>
                                            let shMem_AArow = copyFastMem(
                                                                                invar_AArow_chunk[threadIdx.x],
                                                                                threadIdx.x ) in
                                            loop (acc2) = for k2 < TILE do
                                                    acc2 + shMem_AArow[k2] * 2.0f32
                                            in acc2
                                         , invar_AArow, acc1)
                            in acc1
                    , iota(inds_nhbs), 0f32)
               , A)
   , zip(AA, iota(n)))

(C) MM-Mult
This is the tricky case because is requires 2-dim tiling, i.e., need to create a 2-dim block/workspace such that the access to both arrays can be tiled in fast memory.   To understand how the indexing is transformed, it is instructive to review the transformation first: Semantically the two maps that give the parallel dimensions of the kernel are strip-mined with a TILE factor and then the outer strip is interchanged inward, yielding the doall structure below, which corresponds to the original map-based program:
DOALL ii = 0, m-1, TILE      -- (ii/TILE) gives the workgroup number on Y dimension
    DOALL jj = 0, q-1, TILE   -- (jj/TILE) gives the workgroup number on X dimension
            DOALL i = ii, min(ii+TILE-1, m-1), 1    -- within workgroup: (i - ii) => threadIdx.y
                DOALL j = jj, min(jj+TILE-1, q-1), 1 -- within workgroup: (j - jj) => threadIdx.x
                        -- kernel code
                        acc = 0.0;
                        DO kk = 0, n-1, TILE
                            shMemA[i-ii, j-jj] = A[i, j - jj + kk];
                            shMemB[j-jj, i-ii] = B[j, i - ii + kk];
                            __sync_barrier_sh_mem();
                           DO k = kk, min(kk+TILE-1, n-1), 1
                               acc += shMemA[i-ii, k - kk] * shMemB[j-jj, k - kk]
                        ENDDO ENDDO
                        C[i,j] = acc
ENDDOALL ENDDOALL ENDDOALL ENDDOALL

Rewriting it as a Futhark kernel this would correspond to:
fun [[f32,q],m] main([[f32,n],m] A, [[f32,q],n] B0) =
  let B = transpose(B0) in
  map(fn [f32,q] (int i, [f32,n] a ) =>        -- Kernel's first parallel dimension
            map(fn f32 ( int j, [f32,n] b ) =>  -- Kernel's second parallel dimension
                                                              -- 2-dim Block of size TILE x TILE,
                                                              -- threadIdx.x, threadIdx.y in [0,TILE-1]
                                                              -- fast memory: shMemA[TILE,TILE]
                                                              -- fast memory: shMemB[TILE,TILE+1]
                    streamSeq(fn f32 (int kk, [f32,chunk] a_c, [f32,chunk] b_c, int acc) =>
                                  let shMemA = copyFastMem( a_c[threadIdx.x+kk]
                                                                               , threadIdx.y, threadIdx.x ) in
                                  let shMemB = copyFastMem( b_c[threadIdx.y+kk]
                                                                               , threadIdx.x, threadIdx.y ) in
                                  loop (acc) = for k < chunk do
                                          acc += shMemA[threadIdx.y, k] * shMemB[threadIdx.x, k]
                                  in acc
                             , a, b, 0f32 )
               , zip(iota(q),B) )
     , zip(iota(m),A) )

Three Important Optimization Observations for Matrix-Matrix Multiplication (2-dim Tiling)
First, note that in the code above the access to global-memory array B in the imperative code and to array b_c in the Futhark-kernel code when copying to fast memory is UNCOALESCED, because its innermost index is threadIdx.y+kk, i.e., coalesced would require the innermost index to end in threadIdx.x!   As such I am counting on the transpose to be delayed and to change the indexing into the coalesced one (and if for some reason the transpose is missing then matrix B should be transposed).
Second, note the inner dimension of shMemB is TILE+1 instead of TILE. This is because the (uncoalesced) access to shMemB (both read and write) will result in frequent bank-conflicts if the inner size is a multiple of 16, i.e., TILE. So making it TILE+1 alleviates that to a good extent. Note however that on my NVIDIA card, if you also make  the inner dimension of shMemA to be TILE+1 this generates a 2x slowdown because of plummeting occupancy (I assume).
Another workaround would be to also semantically transpose shMemB, but if too difficult, the TILE+1 inner dimension should do the trick.
Third, when doing block-tiling one should aim to maximize the workgroup/block size, since one access to global memory is amortized over TILE accesses to fast memory. This might be less important for one-dimensional tiling since the gain from 256 to 1024 amortization is probably small (diminishing returns). However for two-dimensional tiling it seems to make a difference: the matrix-matrix multiplication with TILE=32 is about 1.2x times faster than with TILE=16.

This has been implemented and merged, although example 2 is not handled yet.


Feature/navbar

first stab at local state on the front end vs. app state
first stab at styles
reorg'd some of the container logic
Navbar won't match Designs because of limited V1 views; this version has been approved by jesse

closes #12

closes #52

@oconn made revisions. haven't done too much destructuring so hoping that looks right


prefer-compact should warn if filter uses default identity function

As prefer-compact considers an explicit identity function a warning:
_.filter(arr, (x) => x)

It should consider the implicit one a warning as well, shouldn't it?
_.filter(arr)

You're right, it should.

Fixed in version 2.6.1. Thanks for opening the issue!


Add initial version of Upgrade-Support announcement blog post

This PR is for the blog post of Upgrade-Support announcement, it's still a WIP.
TODO

[ ] Add screenshots for sharing a solution;
[ ] Put in correct blog post announcement date;
[ ] Remove WIP notice from Upgrade-Support README.md.

Nice :)
I guess that when this is published we'll have to also update the copy of the releases + react-native-comm/releases 👍
--
Also, this is going to be a separate blogpost from the one for 0.62? Should it come out before that one?

I guess that when this is published we'll have to also update the copy of the releases + react-native-comm/releases 👍

Yes! probably also discuss the approach towards upgrade-related issues on the main repo.

Also, this is going to be a separate blogpost from the one for 0.62? Should it come out before that one?

Perhaps @TheSavior can answer this?

I think this could use some more context as to why we are doing this.
Here is some additional content that you can choose to use in addition to what you have.

After a ton of helpful feedback and discussion by the React Native community, we are eager to announce Upgrade-Support, a central place to request and give help when upgrading your app.
We know that when people upgrade, the occasional XCode build errors, redboxes, or other confusing problems can be frustrating. This requires searching the internet for answers, hopefully stumbling across an issue or blog post where someone had the same error. While we are working on reducing the likelihood you hit these errors with projects like autolinking (LINK NEEDED) and the upgrade helper (LINK NEEDED), we are hoping we can improve the likelihood you can find others solving the same issues.
How it Works
Upgrade Support is a GitHub repo hosting an issue tracker where people can post any issues they have while upgrading their projects. This should help increase the likelihood that people can find others with the same problems and work together to figure out a solution.
If there is a common problem associated with a part of the upgrade, the Upgrade Helper will begin linking to those issues and solutions helping more people get through upgrades smoothly.

@lucasbento, do you have a timeline in mind for this? Now that 0.62 is getting closer I'm wondering if this is a good time, or if it would be too rushed.

What's the status of this blog post draft? Is it going up with 0.62 or...?

yeah I think we can close, in the end in the 0.62 blogpost we added a section about this 👍
But pinging @lucasbento just to make sure

Yep, can be closed 🙂


Updating the SecretManager package name

Issue: https://github.com/aspnet/Docs/issues/943
Here is the manual for installing the SecretManager http://docs.asp.net/en/latest/security/app-secrets.html#installing-the-secret-manager-tool
Here is the issue about the old name and NullReferenceException: https://github.com/aspnet/UserSecrets/issues/34

Hi @sguryev, I'm your friendly neighborhood .NET Foundation Pull Request Bot (You can call me DNFBOT). Thanks for your contribution!

This seems like a small (but important) contribution, so no Contribution License Agreement is required at this point. Real humans will now evaluate your PR.

TTYL, DNFBOT;

Thanks! 195ce020865bc46e83dedca1183395885d671a32

Thank you for the great job!


Extract hook helper

create_hook already existed in hooks.bats. The essential body of create_hook was already in use in exec.bats, rehash.bats, version-name.bats, version-origin.bats and which.bats.
Extracted existing create_hook helper from hooks.bats. Modified it to write the hook to RBENV_HOOK_PATH and to accept hook body from stdin.

I can squash this before merging. Left as is for review

Nice. I was wondering how to check stdin first and use the arg as a fallback (checking number of args wasn't much nicer since we'd want third arg to default to empty string)
Fixed.

Removed string-argument support for hook body; only accepted now via STDIN

Thanks!


Use a temporary directory to store unit test headless run files

Generally it's good practice to store temporary files to the OS defined directory. This avoids IDEs or other OS level operations (ie. spotlight) indexing the files as they are created.
This also has some fixes to avoid log entries being written after a test host has already been disposed. Not sure whether this will be accepted as a solution, but it seems like about the best we're going to be able to do given the static-ness of Logger. Works well in practice.
Some further changes are required osu!-side to make use of this new class (in non-standard scenarios).

drwxr-xr-x  3 smgi    smgi       60 Aug 19 10:33 of-test-headless
Is there a reason why this doesn't append a GUID in all cases? We can't have tests re-using the same dir.

drwxr-xr-x 3 smgi smgi 60 Aug 19 10:33 of-test-headless
Is there a reason why this doesn't append a GUID in all cases? We can't have tests re-using the same dir.

The top level directory is intended so it's easy to clean up the files in one go. It's never used directly for a test - each test will use a subdirectory in that base path. GUIDs are already added to tests where they need unique directories, and the ones which don't (due to requiring constant names between test runs) tend to use a CleanRunGameHost or similar.
Basically nothing has changed from master to this in this regard. That said, if there is a way to handle a global "tests started" and "tests ended" I wouldn't be against making a unique subdirectory in the of-test-headless path for the full run, and potentially nuke it afterwards.

Ah, I thought that was the top-level directory for hosts. It's fine.


Scraper rejects XHTML documents?

Scraper seems to compare the Content-Type of the page being scraped against text/html, and reject the page if it does not match.
This will probably reject pages served as application/xhtml+xml. However, scraping XHTML should be very similar to HTML.
Or am I reading the code wrong?

XHTML is now authorized in commit https://github.com/miniflux/miniflux/commit/3b6e44c33131eb06287d070c3f4b6d19b465f4f5.
Don't hesitate to contribute if you see other issues like this one.


We need some basic information on the ActionPanes collection.

This page implies that a ThisWorkbook class must be added to a VSTO object.  The BaseWorkbook class does not contain the ActionPanes collection.  What is the base class for ThisWorkbook that actually contains the ActionPanes collection?  How would I add a ThisWorkbook class to a VSTO for Excel in Visual Studio 2019?
Also, there is no such thing as Globals.ThisWorkbook.  There is only a Globals.ThisAddIn and this does not contain a ActionPanes collection.
Where is the documentation for the ThisWorkbook class?

Document Details
⚠ Do not edit this section. It is required for docs.microsoft.com ➟ GitHub issue linking.

ID: 3433475e-7319-0e37-22a3-bf49eb58e86f
Version Independent ID: 1a85e45e-204b-46d1-6488-e3f801d83fc0
Content: Add Actions pane to Word documents or Excel workbooks - Visual Studio
Content Source: docs/vsto/how-to-add-an-actions-pane-to-word-documents-or-excel-workbooks.md
Product: visual-studio-windows
Technology: office-development
GitHub Login: @John-Hart
Microsoft Alias: johnhart

@CLIRResearch -- thank you for your feedback.  Please consider these options:

https://social.msdn.microsoft.com/Forums/vstudio
https://www.stackoverflow.com
https://www.reddit.com/r/VisualStudio/

@John-Hart -- please look into this issue.

This walkthrough only applies to document-level projects for Excel and Word which always have ThisWorkbook  or ThisDocument as the base add-in object.  See

It is not a problem of the documentation. You need to create a "Excel VSTO Workbook" or "Word VSTO Document" project to use this walkthrough.


create completion from fine-tuned model

Thanks for the PR!
Instead of adding the model property to the CompletionRequest type, I think it should be better if the model was the first parameter of completeFromModel, just like the complete method. It because the model property should not be available at complete, and it's required for completeFromModel.


Very poor output quality

I have noticed that while it massively increases the inference speed, it massively decreases the quality of the outputs, instruct models become very obstinate and give completely irrelevant responses, words become misspelled, it repeats lines over and over, and also sometimes spams Chinese letters

Tried on Chronos  13b, wizard-lm 13b, and Pygmalion 7b, I used temperatures in between 0.5 and 1, and a context length of 2048, lower temperatures do seem to wrangle it into behaving a little more, but I have to lower the temperature so much the output is too “dry” to be useful

And just to be clear, is this in ExLlama's web UI or in Ooba?

occam's fork of koboldAI that allows using exllama

Not OP, but for context, the Kobold fork is here if you want to check it, turbo.
https://github.com/0cc4m/KoboldAI/tree/4bit-plugin (KoboldAI implementation to support GTPQ and exllama)
https://github.com/0cc4m/exllama (exllama fork on transformers branch, which builds exllama to work on Kobold)
For this one, he reformatted the code to work on Kobold, but the rest and upstream are directly from main exllama.
They added Kobold samplers to that exllama fork.
So it seems these samplers are added

(I'm not sure about rep penaly or rep pen slope though)

Okay. I really have enough work cut out for me with this, but I guess I should try installing Kobold at some point to see how they're using it. I would assume they're just taking the logits and passing them to the same samplers they use for other models, and that should just work. But there are some peculiarities to keep in mind, specifically regarding the cache, and that "context tokens" slider looks a little suspect. But idk.

you think maybe the code wasn't hooked up to the context correctly and it's actually running on incredibly low context size?

I'm not sure what that slider does, but if it truncates the cache that would definitely lead to degenerate output since the position embeddings for cached entries would be wrong. But, looking at the Transformers wrapper they added I think it's just an issue with how the cache is being passed around. It has to stay in sync with the sequence for every forward pass.
E.g. if the model generates an EOS token, and their generator doesn't add that to the running sequence, it has to be removed from the cache. Or something similar along those lines. The cache being out of sync is the kind of thing which might leave it working poorly without crashing. But I'd have to install it and run it in a debugger to make sure. Which I will. After doing some other stuff first.

Alright then, thank you for taking a look at it

@turboderp Apologies for this, this should have gone to me directly.
I do use the KoboldAI samplers, here's the code if you're interested. It seems to work the first time or times you generate, but breaks afterwards. I'm not yet sure why. I do call generator.gen_begin(gen_in), which resets the cache as far as I know.

Yes, using the KoboldAI samplers is the obvious choice for integrating into Kobold, so that's great. There's nothing special about the logits, after all. In fact you should just be able to bypass ExLlamaGenerator altogether and call the forward pass directly.
I'm going to install the 4-bit branch and have a play with it later today. But I don't see anything immediately wrong with how you're using it. gen_begin() should indeed reset the cache (gen_begin_reuse() should work as well, but much faster in some cases), and you're appending every token produced by the forward pass so the cache should stay in sync with the sequence.
I'll have a look though. It shouldn't be too hard to spot if the cache and the sequence go out of sync somehow.

It's not yet that user-friendly to install, you need to clone the branch, run install_requirements.sh and then install the exllama package into the conda env with ./commandline.sh, pip install git+https://github.com/0cc4m/exllama. Then you can run it with ./play.sh

I can confiirm my issue is no longer present after Occam's latest commit to his koboldAI fork, thank you very much for your help.

But... I didn't fix anything yet.

@turboderp Let me know if you need help.

I'm seeing a similar degradation in output quality.  It used to match autogptq output quite closely but latest releases seem to be producing different results.  I can get back previous quality results by setting
ExLlamaConfig.fused_attn = False
hope this can help chase things down.

Well, it's up and running. I was just using a model that didn't have any gptq_bits key in its config and I got stuck on why it wasn't being recognized. Kind of a lot going on in aiserver.py. Maybe you should refactor to less than 10k lines? ;) But it's fine now.
I had to skip the call to tpool.execute() in generate(), just calling model.core_generate() directly in order to debug in PyCharm, but I don't see that having any side effects in this case.
I'm just not seeing anything amiss. It's correctly resetting the cache on each pass, then generating one token at a time and the cache grows as it should, staying exactly one token behind the sequence, and there really isn't much else happening.
The output also looks reasonable. Just trying with 7B Llama, but with the storywriter preset it is telling me a very cute little story that doesn't seem to be degenerating with multiple passes. It does the thing that small models like to do where it starts repeating itself, but you can throw it off by adding in "Until suddenly..." or some such, and all that behaves as I'd expect.
If I swap out gen_begin with gen_begin_reuse, it even seems to be correctly reusing the cache and only re-evaluating the prompt from the first changed token, to further show that it's working. I'm not sure how useful that feature is in Kobold since you're not truncating the sequence in larger steps, so it would only accelerate things until the context is filled up. And prompt eval is really fast already, so idk.
But all in all... I can't find anything wrong at the moment.

The fused attention step is mathematically equivalent to the regular attention, but there might be slight differences related to numerical precision. Maybe if some of the sampling methods are extremely sensitive?

Here's a adjusted snippet of the code - nothing too complicated.  llama is a python class which executes a prompt.  I've had the same issue / tried it with multiple different models from thebloke.  It might just be a user issue with how I'm using the exllama code.  I've setup my code to run with either exllama, autogptq, gptq-for-llama, and llama.cpp so have been comparing them and noticed this difference / issue.
llama.model_path = "models/Nous-Hermes-13B-GPTQ"
llama.tokenizer_model_path = llama.model_path + "/tokenizer.model"
llama.model_config_path = llama.model_path + "/config.json"
llama.model_safetensors_path = llama.model_path + "/" + [x for x in os.listdir(llama.model_path) if x.endswith('.safetensors')][0]
llama.config = ExLlamaConfig(llama.model_config_path)
llama.config.model_path = llama.model_safetensors_path
# llama.config.fused_attn = False
llama.config.max_seq_len = 2048
llama.model = ExLlama(llama.config)
llama.cache = ExLlamaCache(llama.model)
llama.tokenizer = ExLlamaTokenizer(llama.tokenizer_model_path)
llama.generator = ExLlamaGenerator(llama.model, llama.tokenizer, llama.cache)
llama.generator.settings.token_repetition_penalty_max = 1.2

with torch.no_grad():
    # torch.manual_seed(42)
    llama.generator.end_beam_search()

    ids = llama.generator.tokenizer.encode(prompt)
    #llama.generator.gen_begin(ids)
    llama.generator.gen_begin_reuse(ids)

    for i in range(request.max_tokens):
        token = llama.generator.gen_single_token()
        llama.generator.gen_prune_left
        if token.item() == llama.generator.tokenizer.eos_token_id: break
        for eos_token in stopping_criteria_list :
            if llama.generator.sequence_ends_with(eos_token) :
                break
    generated_ids = llama.generator.sequence[0][len(ids[0]):]
    generated_text = llama.generator.tokenizer.decode(generated_ids)

I'll have to try and see if I can reproduce it. One thing that stands out is the call to gen_prune_left() which I haven't looked at in ages. I think it's buggy when called during a beam search. Otherwise, calling it in the generation loop would continually reset the cache so performance would suffer a lot. Maybe it's just a copy/paste error?
Hermes is a model I haven't tested, though. I have found some finetunes to be strangely sensitive to rounding errors. I'll have to check that one out I guess.

Thanks, I can run the sample code you provided have and it works cleanly.  So must be an issue in the code I'm using / how exllama is being called.  The code is trying to be general between all the various GPTQ implementations so might have some cruft causing issues.  Will do more testing and see if i can find out why.

Did some further digging.  Seems to be related to creating the generate and tokenizer objects inside the "llama" class.  When created at the top level it works, but when the exllama objects are created in a class it has the fused attention difference.  Could it be some scoping issues?  For the sample below it generates different creative texts.  But when used for instruction following it produces very  bad results when fused attention is on.
output :
Injected compiler path: C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Tools\MSVC\14.36.32532\bin\Hostx64\x64
------------------- With Fused Attention == True version --------------------
Once upon a time, the only way to get your hands on new music was by waiting for it to come out or finding an underground tape trading scene. Nowadays you can stream and download songs instantly from anywhere in world with just few clicks of mouse button!
The internet has also made sharing information about bands much easier than before – through social media sites like Facebook & Twitter as well blogs that cater specifically towards independent musicians (like this one). This makes discoverability so important because now anyone who wants access to their favorite band’s latest single without having any connection within industry gatekeepers such us record labels A&R people
Injected compiler path: C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Tools\MSVC\14.36.32532\bin\Hostx64\x64
------------------- With Fused Attention == False version --------------------
Once upon a time, the only way to get your hands on new music was by waiting for it to be released and then going out to buy […]
Filed Under: Entertainment Tagged With: Apple Music, Beats 1 Radio Station
code :
from model import ExLlama, ExLlamaCache, ExLlamaConfig
from tokenizer import ExLlamaTokenizer
from generator import ExLlamaGenerator
import torch
import os, glob
torch.set_grad_enabled(False)
torch.cuda._lazy_init()
SET THIS TO CHANGE FROM / TO FUSED ATTENTION
_use_fused_attention = False
print ("")
print ("------------------- With Fused Attention == " + str(_use_fused_attention) + " version --------------------")
print ("")
model_directory =  "../../LLaMa/Nous-Hermes-13B-GPTQ"
class LLaMa_ex :
model_directory :str | None = None
tokenizer_model_path :str | None = None
model_config_path :str | None = None
model_safetensor_path :str | None = None
n_ctx : int = 2048
config : ExLlamaConfig  | None = None
model : ExLlama | None = None
cache : ExLlamaCache | None = None
tokenizer : ExLlamaTokenizer | None = None
generator : ExLlamaGenerator  | None = None
def __init__(self, *args, **kwargs):
    for this_param in list(set(dir(self)) & set(kwargs.keys())) :
        setattr(self, this_param, kwargs[this_param])
    self.model_tokenizer_path = os.path.join(self.model_directory, "tokenizer.model")
    self.model_config_path = os.path.join(self.model_directory, "config.json")
    self.model_safetensors_path = os.path.join(self.model_directory, [x for x in os.listdir(self.model_directory) if x.endswith('.safetensors')][0])
    self.config = ExLlamaConfig(self.model_config_path)
    self.config.model_path = self.model_safetensors_path
    self.config.fused_attn = _use_fused_attention
    self.model = ExLlama(self.config)
    self.cache = ExLlamaCache(self.model)
    self.tokenizer = ExLlamaTokenizer(self.model_tokenizer_path)
    self.generator = ExLlamaGenerator(self.model, self.tokenizer, self.cache)

llama = LLaMa_ex(model_directory = model_directory)
prompt = "Once upon a time,"
llama.generator.settings.token_repetition_penalty_max = 1.5
llama.generator.settings.temperature = 0.5
llama.generator.settings.top_p = 0.1
llama.generator.settings.top_k = 40
gen_tokens = 128
torch.manual_seed(69420)
generated_text = llama.generator.generate_simple(prompt, max_new_tokens = gen_tokens)
print(generated_text)

so does this mean the fix has been rolled into the code, and if so, what files do I replace?

I will try that when I get a chance to, thank you

For what it's worth, I've noticed output quality issues as well in Kobold, which I assumed was related to the sampling swap. However, I noticed similar issues with ooba's very recent exllama support, which doesn't touch exllama's native sampling.
One revealing thing. I was using Wizard-Vicuna-30b, which uses </s> as part of its prompt format. I noticed that I got "</s>" (as in the literal string, not the EOS token) creeping into the output, which never happened with normal transformers. This suggests that exllama is not interpreting </s> as a special token. If it doesn't check special_tokens_map.json, that would explain some things.
In addition, I had issues with very early/jarring EOS, and contraction fumbling (emitting words like can'm, don've, etc.) which is normally only an issue with GPTQ models that don't use desc_act. Neither happened with regular GPTQ. The early stopping may be a symptom of incorrect interpretation of </s> in the prompt, but I'm not sure if that's plausible for contraction fumbling.

Kobold doesn't use ExLlama's sampling, only logits from the model. Ooba does use the native sampling, though, as well as ExLlama's tokenizer which is just a straight SentencePiece instance reading the model file directly. I'll have to dig into the Transformers tokenizer to see if it does something special.
The special tokens map shouldn't lead to you seeing "" in the output, especially when the file that defines that string isn't being read. I can look into some ways to take special_tokens_map.json into account, but it's going to be a little tricky when you have models on HF where that file looks like this:
{
  "bos_token": "</s>",
  "eos_token": "</s>",
  "pad_token": "[PAD]",
  "unk_token": "</s>"
}

The contractions are interesting, at least. Seems too oddly specific to not be a tokenizer issue, but I'm not sure what to make of it. I'll try to see if I can reproduce it. Have you seen it in Kobold too or just in Ooba?

Seen it in both, but it's happening constantly in ooba, every other reply. It's very weird. It manifests in a few ways: just forgetting to finish (doesn') finishing with a weird token (can'the) or cutting off the whole generation at a contraction (doesn'<EOS>). Again, the only other time I was this was with 128g cuda models without act-order, but it was rarer, and I assume it was just quantization error. Fascinating that issues can manifest like this.
For the </s> issue, I suspect this is happening because it's interpreted as a normal string (in the prompt from the chat history), causing the model to assume it should end generations with it. GPTQ parses it as EOS. This might be causing some issues, but I tried removing the EOS tokens from the prompt entirely, and the contraction glitches are still there. Weird.
For what it's worth, the initial exllama branch in Kobold (which was very early, before most of your optimizations, or even support for non-groupsize models) didn't have any generation bugs at all that I could detect.

Sorry for the question here, but the only samplers missing now are tfs (tail free sampling) and top_a, right?

@QM60 : Are you seeing it with different models or just a particular one? Is there one that does it more than others?

Okay, so I did a lot of in-depth testing and I did discover a bug in the handling of the cache during fused attention. It was a little extra sneaky because it doesn't manifest on 7B, and I've been testing a bunch on 7B because it's faster, and just validating on 13B and 33B, but not thoroughly enough to notice the difference.
Anyway, I don't know if this fixes all the issues, but it definitely improves the output on paper.

Interesting, I was just investigating the "contraction disease" issue last night, because I could swear it only started showing up after I pulled a few days ago, so I was trying different commits to narrow down exactly when it started. I was going to spend more time on it today, but it looks like 1ef63db7861ad87e99241a929a8c04e16457b7b3 did indeed fix it.

Anyway, I don't know if this fixes all the issues, but it definitely improves the output on paper.

Massive improvement for me, the contraction issues are gone. The only remaining issue is the </s> one.
Which is interesting, because sentencepiece is deliberately designed NOT to read any control tokens (including EOS) from a normal text stream - which makes complete sense! But the normal llama tokenizer does parse </s> as one token, and vicuna-1.1 and derived models rely on seeing EOS in the prompt. Empirically, for Wizard-Vicuna-30b, omitting it from the prompt is not too bad, but makes it prone to emit very short or even empty replies fairly frequently.
Honestly, I could see the argument that llama's behavior is a bug, but that's what it is. No API lets you inject token IDs into a prompt, so people are used to embedding control tokens into text. I'm not sure if using plain sentencepiece instead of llamatokenizer will cause other issues though. Transformers seemed to have multiple tokenization bugs for llama which suggests something might be tricky about it.

Alright, I celebrated too fast. Contractions are okay but there is still something off about the output from exllama. I suspect it's a sampling issue; perhaps related to sampling order? To identify the differences, I recommend using the sphinx-moth preset from ooba, which has very high temperature mediated by strict top-p and top-k. I think this is more susceptible to whatever the difference is.
Here is a simple comparison using ooba with a fixed seed, llama 30B (plain) quantized to 4 bit with act-order and no groupsize. Settings included. The difference in sanity speaks for itself although I'm kinda fond of the exllama version's energy
exllama comparison.txt

Well, the model is still FP16, regardless, whereas GPTQ-for-LLaMa can also be FP32 depending on how you use it. So it is more susceptible to numerical instability, even if it has no meaningful impact on perplexity.
I would question if sampling with a high temperature is really a good way to get more "varied" output anyway, compared to typical or tail-free sampling, or some such. It makes a lot of assumptions about how the model is supposed to act way outside of the conditions it was trained in, with some inherent assumptions about numerical accuracy. So getting it to work just right by tweaking magic numbers is a balancing act. You could try running without fused attention and MLP, or maybe try lowering the top-p threshold to see if that produces more equivalent output by cutting out some more of the noise on the tail-end of the probability distribution.
I'll have a look anyway a little later. There might be issues with the order of operations. We'll see. But generally speaking, to produce exactly the same output as other implementations in the extremes, I'd have to emulate those other implementations much more carefully, maybe even down to timings in the kernels to get CUDA threads to launch in the same order.
And of course FP32 would eat up a lot more VRAM. I have other ideas for improving accuracy, though, building on the LoRA support I'm currently working on.

I was skeptical that FP32 would explain that kind of a difference in output, so I looked into it, and it's just a sampling difference. For exllama, top_p sampling inherits the base probabilities, from before top_k was applied. That means the same value of top_p is potentially a lot more strict in transformers, where probabilities are normalized at each step.
To confirm, I was able to get the same subjective level of output quality simply by adding one line after if top_p > 0.0: in generator.py to rescale probabilities.
top_probs = top_probs / torch.sum(top_probs, dim = -1)

This could explain some output differences reported here.

I was skeptical that FP32 would explain that kind of a difference in output, so I looked into it, and it's just a sampling difference. For exllama, top_p sampling inherits the base probabilities, from before top_k was applied. That means the same value of top_p is potentially a lot more strict in transformers, where probabilities are normalized at each step.
To confirm, I was able to get the same subjective level of output quality simply by adding one line after if top_p > 0.0: in generator.py to rescale probabilities.
top_probs = top_probs / torch.sum(top_probs, dim = -1)

This could explain some output differences reported here.

Just tried this and indeed gets the same level of output quality. Nice catch!

Well... I'm not fond of making the sampling parameters interdependent in this way. But I guess it doesn't matter all that much since it's kind of all trial-and-error anyway, getting those parameters just right for subjectively satisfying output. So I've pushed an update to normalize the distribution after each sampler.

I wonder whether it is possible to utilize the original decoding pipeline from huggingface transformers to get a more aligned output?

Man... it's better, but I swear something's still off. I switched to Kobold just to rule out the tokenizer and samplers (which are the same there) but exllama still has seizures ever so often. A few highlights:

She grins mischievouslyifizziq
every word and action soDD far
Wouldn't that be fun?"ований текст надо убрать
isn' readonly bad after all…

This doesn't seem likely to be a sampler issue, and I never saw things like that from GPTQ with the same model/settings. It's much better than before, but I'd still view the current forward pass with a bit of suspicion atm...

Are these still with extremely high temperature?

I can relate to that Russian output.  As far as I'm aware it's not running with the temperature cranked, but it's super possible that the temperature is cranked thanks to some bug somewhere in the stack of spaghetti software that exists on top of these LLMs.
Something happens where the first paragraph or so of the output appears normal, and then the "deviation" happens and you get this totally random output after that, with a (apparent) preference towards russian and other non-english output.
To my knowledge, this behavior has been consistent since the beginning - I've not yet seen a change in behavior with any of the updates/fixes for coherence in this repo or on occam's kobold branch.
I'm going to set a reminder to test after work to see if I can get a definitive answer for where this problem actually shows up.  If there are any specific tweaks/changes you want me to trial, give me a heads up.
Get the context my UI frontend sends to kobold that generates iffy responses.
Drop that into the kobold UI, make sure temperature is low, try again a few times.
Drop that into the kobold exllama branch running natively with your script.
Drop that into your exllama repo doing the same.

Alright, this went pretty quick and it looks like the problem lives somewhere in the KoboldAi API or in "SillyTavern" which I am using as the frontend.  Here's an example of a car salesman chat with the two systems:

Compared to:

Kobold itself seems to be rock solid, but my very first response in sillytavern caused it to go haywire.

Suppressing the EOS token might be an instance where the model is extra sensitive to noise, or rounding errors from the FP16 math. If the model is 99.9% sure it wants to emit an EOS token, picking from the remaining 0.1% instead might be asking for trouble.
I'm thinking maybe it's more correct to lower the temperature by some amount proportional to the likelihood of the EOS token, or any other token that gets masked out. Maybe the distribution just ends up being really flat otherwise.
It's interesting that I've never really seen this myself, though.

For what it's worth, I've never seen Kobold suppress EOS when called via its api. It doesn't do that for me; I have generations stop before max tokens have been reached quite often. And my results didn't come from "super high temperature", either, I believe it was about 0.7 with tfs 0.9. But still, sphinx-moth is a normal, popular preset, despite high temp, and as strict as it is with top-k and top-p, tiny FP16 rounding errors aren't likely to cause these kinds of issues.
Banning EOS can cause gibberish, although usually it's a different kind. But the thing is, 90% of gens feel completely normal, and the seizures are sudden and occur at random places in the middle of the text, not at natural stopping points. (The contraction is a clear example of that)
I have no idea how to debug this, though, if the results are nondeterministic due to concurrency. Except maybe fuzzing and comparing logits with transformers output within a tolerance. But that would only detect the bugs, not find them.

Fuzzing is probably the way to go. Perplexity tests aren't revealing anything out of the ordinary, but then they wouldn't if it's an intermittent thing that just corrupts a few numbers under some special conditions. That wouldn't affect the average over thousands of tokens.
I think I'll set up a script to run the same inference in ExLlama and GPTQ-for-LLaMa and look for anywhere the logits deviate considerably between the two. That should at least let me reproduce the error.

For what it's worth, so far, ooba's exllama_hf adapter is giving me flawless output, or at least, indistinguishable from transformers. Not terribly surprising since I think it should have identical tokenizer and sampler behavior.
In summary, what I noticed:

Kobold's exllama = random seizures/outbursts, as mentioned
native exllama samplers = weird repetitiveness (even with sustain == -1), issues parsing special tokens in prompt
ooba's exllama HF adapter = perfect

The forward pass might be perfectly fine after all.

For what it's worth, so far, ooba's exllama_hf adapter is giving me flawless output, or at least, indistinguishable from transformers. Not terribly surprising since I think it should have identical tokenizer and sampler behavior.
In summary, what I noticed:

Kobold's exllama = random seizures/outbursts, as mentioned
native exllama samplers = weird repetitiveness (even with sustain == -1), issues parsing special tokens in prompt
ooba's exllama HF adapter = perfect

The forward pass might be perfectly fine after all.

The only issue that I see is that some models don't load for some reason (https://github.com/oobabooga/text-generation-webui/pull/2777#issuecomment-1602150909)
I'm not sure if @oobabooga or @Larryvrh may know why (Sorry for pinging you guys here, but maybe you can be interested in this)

Well, the forward pass is mutating all the time, and there definitely were some issues with it that have been resolved.
The comparison between the HF samplers and the ones in ExLlama's generator is kind of apples-to-oranges. Without some concrete examples of generations that go "wrong" in ExLlama when they shouldn't, it's really hard to do anything about it. All I get out of that is basically "I tried strumming my bass the same way I strum my guitar, and it doesn't sound the same, so I guess my bass is broken." There could absolutely be bugs in the implementation, but it's impossible to find them based on a subjective sense of something being off.
I could just keep making it more and more identical to Transformers, but the whole point of the project wasn't to create a Transformers-compatible plugin for Ooba and Kobold, it was to build an alternative platform for experimenting with techniques that don't fit well in Transformers.
The tokenizer still shouldn't behave differently, though. With regards to BOS tokens, yes, but that's a separate issue of figuring out what to do about all the incorrect tokenizer configs floating around on HF.

Was not complaining about samplers, just noting my experiences. I don't expect exllama's built-in samplers to be identical at all! If it needs new presets, that's fine. Swapping in known samplers is a very useful way to control for differences, though. And that's important, since it suggests that any complaints you're getting probably aren't due to bugs in the forward pass; at least, not anymore. That's good!
The tokenizer definitely does behave differently though, because it's just wrapping sentencepiece, which explicitly says it doesn't parse special tokens like </s> (by design). HF's tokenizer does parse them. That's not wrong per se. It might even be desirable for someone writing their own backend, since they could insert separator tokens manually, and know that users can't inject them. But it's different!

I'm not taking it as complaints, don't get me wrong. But I also don't want to doubt people when they say the output is bad, like in this case there was a bug causing the forward pass to run slightly incorrectly, and it's good that I found that. I also don't blame people for not having any other way to determine if the output is off than by comparing it to other implementations, because I don't either. It's a black box after all. Just sometimes I do wish it came more in the form of: "here's the output I got, here's (exactly) how I got it, and here's why I think it's wrong."
But yeah, it's not to be confrontational or anything, I just wish there was a better way of communicating that I don't view Transformers as the standard. Maybe I just need an FAQ. :)
As for the tokenizer, I do plan to add some support special tokens. For models that rely on them it gets messy otherwise, if they have to be inserted after encoding and disappear when decoding.

So... how do things stand with the special tokens?
Having studied it a bit, it seems the authors of SentencePiece go out of their way to explain that control symbols are categorically invalid inputs to the encoder. Meanwhile, Transformers implements a very elaborate workaround so they can be encoded anyway. I imagine those two teams don't like each other very much.
Anyway, it wouldn't be too difficult to emulate what Transformers does, but it would be kind of messy so I'm wondering how many models actually include control symbols in their prompt format. Is it unique to Wizard-Vicuna, or is it more common than that?

OpenAssistant also has special tokens like <|endoftext|> that were manually added to the tokenizer. See here: https://huggingface.co/OpenAssistant/oasst-rlhf-2-llama-30b-7k-steps-xor/blob/main/oasst-rlhf-2-llama-30b-7k-steps-xor/added_tokens.json

I see. That's an XOR release, but I assume some the full releases I'm looking at are true enough to the original. It doesn't look like the tokenizer model is changed at all, so it's really all just spread across four different config files, with contradictions and all. Transformers is quite the framework..
Anyway, this makes me wonder if textgen-web-ui makes any attempt at sanitizing user input, or if that's maybe just me overthinking things.

I believe I have a lead on this part of the issue:

Kobold's exllama = random seizures/outbursts, as mentioned

I managed to reproduce the problem with logging enabled and observed the following generation:
   id       gen     RepetitionP     TopP      softmax
 29892    24.8125     22.5628     22.5628        1       [,]
  1183    18.3281     16.7737       -inf         0       [ she]
  310     15.3672     14.1284       -inf         0       [ of]
  322     15.0312     13.7184       -inf         0       [ and]
 20265    1.68555     1.68555       -inf         0       [ Bened]
Selected: 20265 [ Bened]

Despite the comma being the only token with non-zero probability, KAI selected the  Bened token instead. The cause is a bug in torch.multinomial that has been claim fixed but not yet released (i.e. it's still active in PyTorch 2.0.1). This bug sometimes causes the multinomial function to select items with zero weight.
I worked around this bug in the KoboldAI exllama backend by checking the selected tokens and resampling when any zero probability tokens are chosen. I've verified with logging that this avoids the issue and in testing it seems to have solved all the problems with poor output quality.


ebusd 24.1

Hello,
Can you update your repo to use latest ebusd 24.1?

This issue needs to be addressed first https://github.com/john30/ebusd/issues/1382


[test-triage] chip_sw_sleep_pwm_pulses Pass patterns missing

Hierarchy of regression failure
Chip Level
Failure Description
Some pass patterns missing: ['^TEST PASSED (UVM_)?CHECKS$'] has 17 failures:

Test chip_sw_sleep_pwm_pulses has 1 failures.

0.chip_sw_sleep_pwm_pulses.3860101830
Log /container/opentitan-public/scratch/os_regression/chip_earlgrey_asic-sim-vcs/0.chip_sw_sleep_pwm_pulses/latest/run.log

  [617 / 991] Compiling Rust rlib regex_syntax v0.6.27 (32 files); 67s processwrapper-sandbox ... (60 actions, 59 running)
  INFO: From Running Cargo build script libftdi1_sys:
  Build Script Warning: Config for libftdi1 not found, falling back to default search path: `"pkg-config" "--libs" "--cflags" "libftdi1" "libftdi1 >= 1.4"` did not exit successfully: exit status: 1
  [647 / 991] Linking sw/device/lib/arch/libsim_dv.a; 126s processwrapper-sandbox ... (58 actions running)
  [668 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 195s processwrapper-sandbox ... (56 actions running)
  [678 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 246s processwrapper-sandbox ... (55 actions running)
  [688 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 324s processwrapper-sandbox ... (57 actions running)
  [693 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 410s processwrapper-sandbox ... (54 actions, 53 running)
  [694 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 503s processwrapper-sandbox ... (54 actions, 53 running)
  [710 / 991] Action sw/device/lib/testing/test_rom/chip_info.h; 618s processwrapper-sandbox ... (60 actions, 59 running)

Steps to Reproduce

Commit hash where failure was observed: c2a8c64cc
dvsim invocation command to reproduce the failure, inclusive of build and run seeds:
./util/dvsim/dvsim.py hw/top_earlgrey/dv/chip_sim_cfg.hjson -i chip_sw_sleep_pwm_pulses --fixed-seed<PHONE_NUMBER> --build-seed<PHONE_NUMBER> --waves -v h

Tests with similar or related failures

[ ] Just chip_sw_sleep_pwm_pulses

it looks like this test has been passing for multiple nightlies.
The error seen here is most likely a regression infrastructure issue (there's been a few) that has since been fixed. I will close this one down for now.