Ptxas slow

[Curefab]

Do several threads or even warps share the same instructions/instruction streams? How large is the working memory space for data the threads have to access exclusively or shared? Does it fit into registers or max larger? Can you say anything about the kind of instructions you want to combine and execute: Data movement and arithmetic instructions? Conditionals? If arithmetic instructions, can some of those be expressed as matrix multiplications (for using the Tensor Cores with their higher TFLOPS)?

The SASS instructions are 128 bits/instruction (16 bytes) with the same instruction loaded for a warp of 32 threads. All the 100.000.000 instructions (1.6 GB) need memory bandwidth. If all threads/warp have the same instructions it has to be streamed only once through L2 and distributed to the SMs as long as the warps stay roughly synchronized (which can be achieved). The high rate of 80 TFLOPs you can only achieve, if several threads/warps have the same instruction stream or you can repeat instructions in a loop. Otherwise you would need 1.28 Petabyte of generated SASS per second (80 T * 16 bytes).

[havlikmarcell]

@Curefab

The idea with AMD is (as far as I understand) to speed up to compiler by a small factor (perhaps x2?). Not sure, if that helps you, if you want 100x or 1000x.

With custom compiler I would go with an extremly simple that is O(1) and literally just “mapping” the task to the SASS, so the sequence of calculations and allocation be calculated pretty well. That’s why I was believeing in the 1000x actually. We were talking about how many instruction can we do per SASS with @njuffa in #2 post. So theoretically even on single thread it could be possible.

@Curefab

You were talking about AI in your original post.

Yeah AI library is really inaccurate, I meant an ML library but of course we wanted to go further so eventually add some custom algorithm that already sort of and artificial intelligence, that calculates “how to learn” something automatically. We have many tries on this but this feature is essential to speed it up and be useful.

@Curefab

The SASS instructions are 128 bits/instruction (16 bytes) with the same instruction loaded for a warp of 32 threads. All the 100.000.000 instructions (1.6 GB) need memory bandwidth. If all threads/warp have the same instructions it has to be streamed only once through L2 and distributed to the SMs as long as the warps stay roughly synchronized (which can be achieved). The high rate of 80 TFLOPs you can only achieve, if several threads/warps have the same instruction stream or you can repeat instructions in a loop. Otherwise you would need 1.28 Petabyte of generated SASS per second (80 T * 16 bytes).

Yes, I wanted to ask a very standard case, not calculating with any additional thing, that are obviously there. Thank you nothing 1 warp need 1 instruction only. So yes, 128 bits/instruction. So 1 instruction / 32 threads. On 1008 GB/s bandwith (4090) it means 63 billion instruction/sec. So all in all this allow us to get 63 billion * 32 FLOPS/sec = 2016 GFLOPS ~ 2TFLOPS which is already pretty nice. Only 40th of the theoretical speed. If I am right in this case.

But this is really a very far upper limit of course. But processing every single instruction from global memory eventually gives us this hard limit. If we cannot optimize it a little with compression, then this should stay there for upper bound for this ML library in this case which isn’t too nice. (But to be honest with compression it isn’t trivial how to achieve high enough compression to be worth to work on it.)

[havlikmarcell]

Ok, I am starting to believe there are these bounds in each case and a lot of people already thinked about it weather it worth to create larger kernels as I guess the ICache also will suffer due to this.

One thing I got into my mind that it can still worth to have this way to work if we can do filter out calculations by using this graph computation instead of dense calculations. So eventually we could realize that only 100th or even 1000th of the instructions are meaningful, so basically with that in mind we could reach “theroical 4000Tflops” .

But if someone else has any additional idea, don’t hesitate to share. :D
I really hope someone will figure out how to do this “graph” like neural networks in an effective manner.

[Curefab]

@havlikmarcell

Only 40th of the theoretical speed. If I am right in this case.

It is good that the threads within warps have common instructions.

When using the full memory bandwidth for the instructions and ptxas or an AMD variant as assembler, you still have to transfer the instructions to the GPU over PCIe. I am not sure if it is possible (probably not) to implement an assembler on Cuda and run instructions from there (although it would be logical and theoretically possible).

When you create custom 2 byte instructions, and use a jump table in device code for executing. Then the kernel could reach a maximum of 1/5th of the theoretical speed. The custom instruction stream could be created inside another kernel. But please be aware that registers are not dynamically addressable (= not with index). So you would have to program each register access for each instruction or use shared memory, which can be accessed with an index. Or a combination of both: A few hard-coded registers and shared memory with index.

Another way: Somehow you generate the long program code from some input. Couldn’t you instead of generating the program code instead directly execute it from the initial input?

[havlikmarcell]

@Curefab

Another way: Somehow you generate the long program code from some input. Couldn’t you instead of generating the program code instead directly execute it from the initial input?

We could and actually as there is no other way we will try. We just have to do it in spite of the fact that this won’t get 100% utilization on the computation side.
So all in all, any program that is too big to be loaded will have the hard limit 2 TFLOPS due to “loading the instructions” alone takes so much bandwidth per flops (on a 4090).

@Curefab AMD side should be similar indeed.

By the way at least this compiler is actually a very simple project compared to any rewrite of the compiler, and it could fit on the ICache which is nice. I guess my problem will be with the uncoalesced read that will happen due to the pointer to different inputs many times. Also, I am not quite sure if there are any chance to use the registers or shared memory in these cases (so also global memory) because that also a pretty big drawback (even 3-10x I believe from the 2TFLOPS). I would need more time to figure out the possibilities here. (I guess it could be used with modularity and have advanced “operation” set.)

@njuffa by the way if the whole code doesn’t fit into the ICache, it will be swapped to global memory right? So bandwidth limit applies there too? Just want to make sure we calculate correctly about these limits.

[Curefab]

There perhaps are ways to solve the coalesced read topic.

For example (not sure, if the preconditions are met in the case of your problem) to go from random loads to as coalesced ways as possible:

Instead of reading once, you use two passes. In the first pass, the read requests are generated and stored in a coalesced way. In a second pass in a loop over memory regions suitable read requests are bundled, executed and the result distributed.

Otherwise, you should try to use sectors of 32 bytes (256 bits) as L2 is arranged with this granularity.

Fully coalesced 128 bytes (1024 bits) are only needed for optimal performance in the L1 cache and for less requests there.

To increase the data size, e.g. instead of two random accesses to two arrays arr1[idx], arr2[idx], you can have a structure stored in a common array: arr[idx].data1 and arr[idx].data2. This doubles the amount of data, which is read with a single transaction.

Writing is typically more performant, as you do not have to wait for a returned value, and you could just write in a coalesced way serially and sort the data later.

[njuffa]

@havlikmarcell

by the way if the whole code doesn’t fit into the ICache, it will be swapped to global memory right? So bandwidth limit applies there too? Just want to make sure we calculate correctly about these limits.

I would suggest relying on the CUDA profiler for assessing the impact of ICache misses on kernel performance, as it may well differ by GPU architecture. As far as I recall, there is straight-line prefetching of instructions, so ICache misses are primarily an issue for loops (with a backward edge in control flow) that are too large to be fully contained in the ICache.

[njuffa]

@Curefab

is it possible to create SASS with the GPU directly in global GPU memory and afterwards run it there? Or does each kernel always have to be uploaded from the CPU?

I am not aware of the possibility of self-modifying code on the GPU. In modern security-conscious environments, writeable and executable memory attributes are usually handled in exclusive-or fashion. In order to prevent malware intrusion into a system via GPU, I would assume the same holds for GPUs. It would not surprise me if someone figured out how to circumvent such mechanisms.

[Curefab]

@njuffa

I am not aware of the possibility of self-modifying code on the GPU.

Self-modifying code would be an extreme case. But would it be possible to create SASS code on the GPU and then memcpy it into an executable region without going over the host?

Probably the hardware could do it, but the driver does not offer it?

[njuffa]

It depends how one defines self-modifying. Lets assume we have a kernel foo() running on the GPU and it calls a device function bar() via a function pointer. foo now overwrites the machine code for bar at run-time. That is what I assumed you envisioned. I would call that self modification, other might not, since foo modifies bar, not foo itself. I would say it does not matter qualitatively, as bar() could be inlined into foo(); keeping it separately as a callable functions just makes it easier to know which memory locations to modify.

It should be easy enough to try. Lets assume bar() is a function that either increments or decrements its input, and returns that as the result. Compile both variants off-line with -rdc=true (so it uses the internal ABI), and extract their respective binary code blobs. This ensures we have valid machine code to inject. Store the two binary code blobs in foo() and try overwriting bar() with either of the stored versions. Does it work?

[Curefab]

I would think that the MMU on the GPU prevents it. At least for compute shaders. Shaders with elevated privileges for initialization, error handling, debug, etc. could perhaps circumvent the MMU, but those are probably even less disclosed or accessible. (I only know that there is a user and trap mode for shaders with trap mode being able to access more special registers, probably to retrieve the location of the error or for resuming.)

But perhaps one can use special host calls to use data created on the GPU as executable code or only do a device to device copy instead of going over PCIe. The open GPU Kernel Modules released in 2022 could be a useful resource GitHub - NVIDIA/open-gpu-kernel-modules: NVIDIA Linux open GPU kernel module source to find out, what happens, when a cuda function is loaded (e.g. when calling cuLibraryLoadData()).

[Robert_Crovella]

A simple test, attempting to “read” code, doesn’t seem to work for me:

# cat t138.cu
#include <cstdio>

__device__ void foo(int &a){a += 13;}

__device__ void (*myfp)(int &a) = foo;

__global__ void k(int val) {
#ifdef TEST
  unsigned char *mycode = (unsigned char *)myfp;
  for (int i = 0; i < 1; i++) printf("%u ", mycode[i]);
#endif
  printf("\n");
  myfp(val);
  printf("val = %d\n", val);
}

int main(){

  k<<<1,1>>>(4);
  cudaDeviceSynchronize();
}
# nvcc -o t138 t138.cu -rdc=true
# compute-sanitizer ./t138
========= COMPUTE-SANITIZER

val = 17
========= ERROR SUMMARY: 0 errors
# nvcc -o t138 t138.cu -rdc=true -DTEST
# compute-sanitizer ./t138
========= COMPUTE-SANITIZER
========= Invalid __global__ read of size 1 bytes
=========     at 0xa0 in k(int)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7fcc0fc5d100 is out of bounds
=========     and is 54,906,361 bytes after the nearest allocation at 0x7fcc0c800300 of size 8 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x32e950]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x1074e]
=========                in /root/bobc/./t138
=========     Host Frame:cudaLaunchKernel [0x7095e]
=========                in /root/bobc/./t138
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) [0xb0ca]
=========                in /root/bobc/./t138
=========     Host Frame:__device_stub__Z1ki(int) [0xaf70]
=========                in /root/bobc/./t138
=========     Host Frame:k(int) [0xafa3]
=========                in /root/bobc/./t138
=========     Host Frame:main [0xade5]
=========                in /root/bobc/./t138
=========     Host Frame: [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xaad5]
=========                in /root/bobc/./t138
=========
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame: [0x47e786]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:cudaDeviceSynchronize [0x48874]
=========                in /root/bobc/./t138
=========     Host Frame:main [0xadea]
=========                in /root/bobc/./t138
=========     Host Frame: [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xaad5]
=========                in /root/bobc/./t138
=========
========= ERROR SUMMARY: 2 errors
# cuobjdump -sass ./t138

Fatbin elf code:
================
arch = sm_52
code version = [1,7]
host = linux
compile_size = 64bit

        code for sm_52
                Function : _Z3fooRi
        .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM52 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM52)"
                                                                  /* 0x001ec800fda207f6 */
        /*0008*/                   LEA R6.CC, R4.reuse, RZ ;      /* 0x5bd780000ff70406 */
        /*0010*/                   LEA.HI.X P0, R7, R4, RZ, R5 ;  /* 0x5bd802c00ff70407 */
        /*0018*/                   LD.E R0, [R6], P0 ;            /* 0x8090000000070600 */
                                                                  /* 0x003ffc001fa107f4 */
        /*0028*/                   IADD32I R0, R0, 0xd ;          /* 0x1c00000000d70000 */
        /*0030*/                   ST.E [R6], R0, P0 ;            /* 0xa090000000070600 */
        /*0038*/                   RET ;                          /* 0xe32000000007000f */
                                                                  /* 0x001f8000fc0007ff */
        /*0048*/                   BRA 0x40 ;                     /* 0xe2400fffff07000f */
        /*0050*/                   NOP;                           /* 0x50b0000000070f00 */
        /*0058*/                   NOP;                           /* 0x50b0000000070f00 */
                                                                  /* 0x001f8000fc0007e0 */
        /*0068*/                   NOP;                           /* 0x50b0000000070f00 */
        /*0070*/                   NOP;                           /* 0x50b0000000070f00 */
        /*0078*/                   NOP;                           /* 0x50b0000000070f00 */
                ..........


                Function : _Z1ki
        .headerflags    @"EF_CUDA_TEXMODE_UNIFIED EF_CUDA_64BIT_ADDRESS EF_CUDA_SM52 EF_CUDA_VIRTUAL_SM(EF_CUDA_SM52)"
                                                                  /* 0x001fc800fe2007f6 */
        /*0008*/                   MOV R1, c[0x0][0x20] ;         /* 0x4c98078000870001 */
        /*0010*/                   MOV32I R18, 0x0 ;              /* 0x010000000007f012 */
        /*0018*/                   MOV32I R19, 0x0 ;              /* 0x010000000007f013 */
                                                                  /* 0x0c1fd800f64007f0 */
        /*0028*/                   IADD32I R1, R1, -0x10 ;        /* 0x1c0fffffff070101 */
        /*0030*/                   LDG.E.64 R2, [R18] ;           /* 0xeed5200000071202 */
        /*0038*/                   LEA R8.CC, R2.reuse, RZ ;      /* 0x5bd780000ff70208 */
                                                                  /* 0x0003e800fc4007f1 */
        /*0048*/                   LEA.HI.X P0, R9, R2, RZ, R3 ;  /* 0x5bd801c00ff70209 */
        /*0050*/                   MOV R3, c[0x0][0x140] ;        /* 0x4c98078005070003 */
        /*0058*/                   STL [R1], R3 ;                 /* 0xef54000000070103 */
                                                                  /* 0x001fc400fe2007b1 */
        /*0068*/                   LD.E.U8 R0, [R8], P0 ;         /* 0x8010000000070800 */
        /*0070*/                   LOP.OR R2, R1, c[0x0][0x4] ;   /* 0x4c47020000170102 */
        /*0078*/                   MOV32I R4, 0x0 ;               /* 0x010000000007f004 */
                                                                  /* 0x001fc400fcc007f4 */
        /*0088*/                   MOV32I R5, 0x0 ;               /* 0x010000000007f005 */
        /*0090*/                   IADD32I R17.CC, R2, 0x8 ;      /* 0x1c10000000870211 */
        /*0098*/                   IADD.X R16, RZ, RZ ;           /* 0x5c1008000ff7ff10 */
                                                                  /* 0x0403d400fe2007f5 */
        /*00a8*/                   MOV R6, R17 ;                  /* 0x5c98078001170006 */
        /*00b0*/                   MOV R7, R16 ;                  /* 0x5c98078001070007 */
        /*00b8*/                   STL [R1+0x8], R0 ;             /* 0xef54000000870100 */
                                                                  /* 0x001fc400fe200ffd */
        /*00c8*/                   JCAL 0x0 ;                     /* 0xe220000000000040 */
        /*00d0*/                   MOV32I R4, 0x0 ;               /* 0x010000000007f004 */
        /*00d8*/                   MOV32I R5, 0x0 ;               /* 0x010000000007f005 */
                                                                  /* 0x001ff400fec007f1 */
        /*00e8*/                   MOV R6, RZ ;                   /* 0x5c9807800ff70006 */
        /*00f0*/                   MOV R7, RZ ;                   /* 0x5c9807800ff70007 */
        /*00f8*/                   JCAL 0x0 ;                     /* 0xe220000000000040 */
                                                                  /* 0x001fd800162007f0 */
        /*0108*/                   MOV R4, R2 ;                   /* 0x5c98078000270004 */
        /*0110*/                   LDG.E R18, [R18] ;             /* 0xeed4200000071212 */
        /*0118*/                   MOV R5, RZ ;                   /* 0x5c9807800ff70005 */
                                                                  /* 0x001fc021ffa007f1 */
        /*0128*/                   PRET 0x138 ;                   /* 0xe270000000800040 */
        /*0130*/                   JMX R18 ;                      /* 0xe20000000007120f */
        /*0138*/                   MOV R6, R17 ;                  /* 0x5c98078001170006 */
                                                                  /* 0x001fc400fe2007b1 */
        /*0148*/                   LDL R0, [R1] ;                 /* 0xef44000000070100 */
        /*0150*/                   MOV R7, R16 ;                  /* 0x5c98078001070007 */
        /*0158*/                   MOV32I R4, 0x0 ;               /* 0x010000000007f004 */
                                                                  /* 0x003ff4201ea007f1 */
        /*0168*/                   MOV32I R5, 0x0 ;               /* 0x010000000007f005 */
        /*0170*/                   STL [R1+0x8], R0 ;             /* 0xef54000000870100 */
        /*0178*/                   JCAL 0x0 ;                     /* 0xe220000000000040 */
                                                                  /* 0x001f8000ffe007ff */
        /*0188*/                   EXIT ;                         /* 0xe30000000007000f */
        /*0190*/                   BRA 0x190 ;                    /* 0xe2400fffff87000f */
        /*0198*/                   NOP;                           /* 0x50b0000000070f00 */
                                                                  /* 0x001f8000fc0007e0 */
        /*01a8*/                   NOP;                           /* 0x50b0000000070f00 */
        /*01b0*/                   NOP;                           /* 0x50b0000000070f00 */
        /*01b8*/                   NOP;                           /* 0x50b0000000070f00 */
                ..........

EDIT/Later: At least in some cases it does not seem to generate the error I reported above.

@Curefab

One SASS instruction has a size of 16 bytes.

It’s also not obvious to me that one SASS instruction universally has a size of 16 bytes, although I acknowledge it might be possible to express every SASS instruction in a 16-byte field.

[Curefab]

@Robert_Crovella

It’s also not obvious to me that one SASS instruction universally has a size of 16 bytes, although I acknowledge it might be possible to express every SASS instruction in a 16-byte field.

Kepler has generally 8 bytes (+ 1 control word every 7 instructions), Maxwell 8 bytes (+ 1 control word every 3 instructions), since Volta 16 bytes with included control instructions.

Here are some more details: maxas documentation Google Code Archive - Long-term storage for Google Code Project Hosting.
and here: Dissecting T4 via Microbenchmarking https://arxiv.org/pdf/1903.07486.pdf (Chapter 2)

[Curefab]

@Robert_Crovella

A simple test, attempting to “read” code, doesn’t seem to work for me:

Yes, and there seem to be no Cuda driver API functions taking a device pointer for providing .fatbin or similar code. That means the large kernels would always have to go through the PCIe interface after being modified or created.
Thus reducing the achievable FLOPs of the whole 4090 (with PCIe 4.0) to 63 GFLOPS from assumed 2 TFLOPS. (or depending on how one counts to a maximum of 126 GFLOPS from 4 GFLOPS if all the instructions are FMA, which combine multiply and add), if each warp has a separate instruction stream, as havlikmarcell stated.
And no fast ptxas running on AMD would change that. And also no parallel patching and merging of SASS, even when run on the GPU.

So the best remaining bets would be either
a) inventing a slim custom language, which the (normal-sized) kernel on the GPU interprets. And streaming the 100 million instructions for the interpreter like normal data. This normal data can optionally be generated on the GPU and can stay on the GPU without PCIe roundtrip.
b) not creating the 100 million instructions in the first place, but directly executing the code just depending on the input, the code generator from a) would get.

[njuffa]

@havlikmarcell

This may be of interest: The authors of the following publication created a replacement for ptxas, with support for the Pascal, Volta, and Ampere architectures, with the goal of speeding up the online query compilation for GPU-based databases. Their source code is available on Github.

Alexander Krolik, Clark Verbrugge, and Laurie Hendren, " rNdN: Fast Query Compilation for NVIDIA GPUs." ACM Transactions on Architecture and Code Optimization, Vol. 20, No. 3, Sep. 2023, Article No. 41, pp 1–25

Results indicate that our pipeline is effective at improving both compilation and end-to-end time against all comparison systems, while maintaining high speedups in cached scenarios. Importantly, in comparison to NVIDIA’s optimizing assembler, we vastly improve compilation time while showing only minor slowdown in kernel performance.

I have not tried this software or looked at the details of the Github repository, but this may provide a starting point for further research.

[havlikmarcell]

Hey @njuffa , I wanted to give feedback.
I read your post and followed the source too, just lost during that track and didn’t give feedback.

These are really great materials! So thank you for this! This can be definitely useful for me later on or anyone who face with similar problem!

I lost in the LLM space for a while, so hopefully I will get back to this topic later this year.