CUDA Programming for NVIDIA H100s – Comprehensive Course

Key Takeaways

Welcome to this intensive 24-hour course designed to bridge the gap between basic coding and high-erformance CUDA programming for NVIDIA hopper GPUs. You will learn to build efficient WGMMA pipelines and leverage cutless optimizations to perform the massive matrix multiplications that power modern AI. Beyond single chip performance, the curriculum covers multi-GPU scaling and NCCL primitives necessary for training trillium parameter models. To get the most out of these lessons, you should have a foundational grasp of C++ syntax and linear algebra, particularly how matrices are tiled and multiplied. Ultimately, this course helps you build a strong mental model of asynchronous execution so you can write actual kernels for the world's most advanced hardware. This is a 24-hour course for CUDA programming for Nvidia Hopper GPUs. Not just the basics, not just an overview, but in this course, we are going to learn how to write really efficient WGMMA pipelines for doing fast matrix multiplication operations. We are going to learn what are the different optimizations hidden behind the code for cutless and how to use that to write actual kernels for ourselves. We are going to learn how these multi-GPU setups work and what are the nickel primitives which enable us to train these neural networks for these trillion parameter models. There's nothing like this which exists in the internet. I looked it up and that's why I created it. Now, before we start, I want to tell you that this course is free and I intend for this course to be free, but it took me hundreds of hours to create this whole course. And the reason why I built this is because I don't think that this whole knowledge should be logged behind pay walls or scattered across different code and different documentations. So if you are an engineer then I hope that this course changes your career. If you are a company or some executive or some senior engineer in a company and you want me to go further or if you want to support this whole work then find my contact details on the description and if you are someone who who likes this or if you are someone who just learn something new from this send this to one person that's all I ask let's begin so there are three or four major prerequisites for this course. First is a basic understanding of C++ C syntax and this course would be built on KUDA C++ and PTX. So I at least assume that you have some basic amount of knowledge for of how C++ and C works. The second prerequisite for this course is a functional understanding of older architecture and a synchronous CUDA execution model. So things like how how the whole thread grid hierarchy works things like shuffle sync and how these different you know CUDA specific functions work. So if you cannot write a standard kernel then it would be really difficult to understand the asynchronous model of programming which H100 is based upon. The third prerequisite is basic linear algebra specifically how matrix multiplication operations work and some of the basic memory layouts. Tile matrix multiplication understanding is a must because that is the basic layer on which we are going to build this whole course. Also a basic understanding of transformers and multi-layer perceptrons would help a lot. This is not a compulsory requirement but we would be using these terminologies a lot during this course and knowing why we are optimizing these specific matrix operation is really important because it just accelerates your understanding of the bottlenecks and of different operations which we are trying to do. Now there are two main things which you have to have open for most of the time in this course. First is a notebook and a pen to create mental models. And the second thing is an AI assistant. What are mental models? Mental models are just your interpretation of how these different components work. They are your compressed way of seeing the reality. And when you see a messy situation or something which is really complicated, which would happen a lot of times in this course, then you can predict what matters without memorizing everything. If you have a strong mental model, an AI assistant and myself are there to refine that mental model and to help you create it. Every single time if your mental model does not match the reality or if there is any kind of doubt then you must go to AI model or you must you know rewind your whole thing to validate your mental model and to fix things if there are something wrong. Now there are a few recommendations which I would give you for completing this course. Let's go through them quickly. Now most of the problems when you are learning hopper arrive either from having bad fundamentals or mental model. So whenever you are struck somewhere just throw all of your question on an AI model and try to understand what is happening. The second reason why a lot of people just is stall during their learning journey is because they try to understand everything at once. This is a trap. Your job is not to master hopper, cutlas, fast.co and thunder kittens in parallel. Your job is to build a staircase. So get through it gradually and don't just stall if you don't understand a single thing. Just try to refine your mental model and keep moving ahead. In the first few lessons, you are not supposed to fully understand everything. You are not supposed to have a full grasp of how everything comes together and what are the different ways in which you know everything is integrated. First you are supposed to get the map. First you are supposed to get you know that this component works this way, this component works that way. All of that. then later lessons would add the mechanism and then the code would give you the concrete reality. So if you think that lesson one must already feel complete then you would panic the moment as lesson 9 it gets gnarly. So it's really important to have that patience and to go through everything slowly and of course use the AI agents to learn about code bases. If you have something like clot code or codeex then launch it in cutlas or thunder kittens or fast.cu and ask questions from there. That is the closest you can get to the accurate information for whatever code we are learning about and the other primitives which we would see. These are the table of contents for our whole course. We would start by looking at the architecture of the hopper GPU of the H100. Then we would go to clusters, data types, inline ptx and pointers. These are the things which we are going to learn next. Then we are going to learn a really important topic called asynchronicity and barriers. And then we are going to learn about tensor map which is a descriptor for doing these large copies from global memory to shared memory using a unit called tensor memory accelerator. Then we get to cp.asing.bulk bulk which is the instruction for executing these large copies. Then we get to the WGMMA part which is the instruction for doing the actual matrix multiplication operations really fast on tensor course. Then we get to the second lesson in WGMMA where we will be discussing how to patch all these operations, how to store them and what are the different other forms of WGMMA operations like WGMMA on is sparse tensor course and FPA WGMMA. Then we get to kernel design and with kernel design there would be the code part. So first for like around 3 4 hours we would be just learning what are the different ways in which we put all of these things together to write a really efficient kernel and what are the different components for writing an actual good kernel and then what we would do is that we would start looking at the source code of cutlas the source code of fast.cu CU to validate all of that and to try to understand how these things are actually you know implemented in the code and then finally we we have these two lessons on multiGPU and the first lesson would talk about different multiGPU setups which we have for H100s and we would talk a bit about things like MPI and generic CUDA P2P programming and in the second lesson we are going to learn about nickel and how to initialize nickel and pmix slum and what are the different data parallelism and model parallelism and pipeline parallelism strategies which we have things like that we are going to study. So let's talk about what we would be going through in the first lesson. So the first lesson would give you a technical introduction to the H100 GPU. you would be focusing on its internal architecture, hardware components and the significant shift towards asynchronous execution. So we are going to learn about key architecture specifications. We are going to learn about all the technical innovations about memory and executional hierarchy and the operational logic. The second lesson would dive into the programming and memory models. So it would cover four primary areas. First is thread block clusters, inline PTX, memory state spaces and the hardware synchronization. So these are the things which we would be studying. We would be learning about thread block clusters and what are the different primitives for them. We would be learning about parallel thread execution or PTX and inline PTX which is the virtual assembly for our CUDA programs. We are going to learn then about state spaces and different data types and different pointer operations which we do in CUDA. Then once we are done with that then we get to one of the most important lessons in the course which is on asynchronicity and barriers. And this lesson would purely focus on asynchronicity and hardware barrier. And it would detail how Nvidia H100 manages this massive speed gap between its fast computational units and the relatively slow memory accesses. So we would be learning about the philosophy of asynchronicity. We would be learning how to prevent the memory hazards. We would be learning about M barriers. And this is the topic which is going to take most of our time. And then we are going to learn about advanced synchronization patterns and other synchronization primitives which we have. Then we get to the fourth lesson and this basically provides a deep dive into the tensor memory accelerator and the descriptor for tensor memory accelerator which is cool tensor map. And this tells us which what is the way in which TMA works and uh which are basically what are the different hardware and software components that enable this whole H100 signature asynchronous data movement. So that's what we are going to learn in this lesson. We are going to learn what is TMA. We are going to learn about tensum map descriptors and we are going to learn about all these different optimization features which you can encode in the censor map like swizzling for shared memory like L2 promotion like interle all these different optimization features which you can enable in the tensor map to to be executed when you launch your bulk asynchronous copy using TML and if you don't understand these these technical ical terms right now. Don't worry, we are going to discuss them in a lot more detail later. The next thing which we would learn after that is the cp.acing.bulk instruction and this lesson would be on the main PTX instruction which we have and this is the instruction which utilizes the TMA and the coupensor map to do bulk asynchronous copies. We are going to talk about the innovations which we have in asynchronous data movement. the core copy instructions which we have and the different variations of the cp.asing.bulk instruction which we have here. So the core concept which we would be discussing would be the instruction cp.asing.bulk and we would be discussing about the tensor memory accelerator. Then we would be learning what are the difference in which differences between the asynchronous copies in hopper versus ampere. And we're going to learn about different data layouts for doing these bulk copies. So unstructured bul copies versus structured bulk copies for for tensor operations or convolution operations. Then we are going to learn about synchronization and completion mechanisms which are there for these bulk asynchronous copies. Then we are going to learn about L2 cache hinting and multiccasting which are really important part of the cp.async.bulk bulk instruction and finally we are going to learn about pre-fetching data and asynchronous reductions for these bulk asynchronous copies and these these are the operations which use the TMA but they are not exactly bulk copies so that's what we are going to learn and then after learning all of the data movement side of things we would get to the main computational part so this This is the first part of WGMMA operation and this lesson would focus on executing the actual math for the tensor course using the W group matrix multiply accumulate operations on the hopper GPU. So this is the operation which is there for doing the MMA operations on on the tensor course and basically in this whole lesson we are going to learn about the paradigm shift in the MMA operations from WOPS to WO groups. We are going to learn about the core instruction which we have that is WGMMA MMA async. We are going to learn how to fetch matrix A for doing these WGMMA operations. And then we are going to learn how to manage these registers for if we fetch matrix A through registers. And then finally we are going to learn about the WGMMA descriptor which is used for these MMA operations. So that's what we are going to do and once we are done with that then we would get to WGMMA part two and this is the lesson which would wrap up the whole hopper math pipeline by covering synchronization is how we store results and the nuances of low precision and sparse operations. So we are going to learn about pipelining with WGMMA groups. We are going to learn how we store output of these WGMMA operations. We are going to learn how things have changed in the FP8 paradigm and how you know with MP and the previous generations of the GPUs, we didn't have the support for doing FP8 MMA operations in the GPU and in the tensor course. So with Hopper, one of the biggest focuses was on enabling the use of FP8 data types for doing WGMMA operations and using these really fast operations for FP8. So we are going to learn about that. Then we are going to learn about some of the FP8 quirks of the H100 GPU. things like packing and the K major layouts and basically the behavior of tensor codes when we are using FP8 and accumulating in FP32 and finally we are going to learn about is sparse WGMMA which is WGMMA on structurally sparse tensor course so that's what we are going to do and finally we would get to the most important lesson in the whole course. This would be the lesson on kernel design and this is the lesson which would basically enable you to use these different ways of of writing kernels and it would teach you all the different specific design patterns for optimizing computebound metal kernels. So that's what this whole course would be based on and the core philosophy of this core of this lesson would revolve around W specialization pipelining and persistent scheduling to ensure that your GPU stensor codes are utilized near 100% efficiency. So we would start by classifying the ridge point for computebound kernels and how at some arithmetic intensity a kernel might be classified as as a memory bound kernel and how at some arithmetic intensity a kernel might be defined as computebound kernel. We are going to learn about that. We are going to learn about W specialization. Then we would learn about pipelining and circular buffering and how these different things work. Then we are going to learn about intelligent scheduling which would be one of the most complicated part of this whole course and finally we are going to learn about epilog and some of the multicernel optimizations. So these are the different things which we are going to learn and then we would get to code and we are going to walk through all of these different files for code in this whole course. So we are going to learn about SM90 pipeline.hpp. We're going to learn how this file works and how this file basically organizes the whole ping-pong pipeline for the gem operations. We are going to learn about cooperative pipelining. We're going to learn about how MMA operations and WGMMA math works on RS and SS layouts and how all of these things work. Then we are going to look at the tile scheduling and we are going to look at tile scheduler for group group matrix multiplication for stream k operations and for general you know persistent scheduling and how tile scheduleuler works for all these different things. Finally, we are going to look at SM90 epilog pop specialized uh file which basically concerns itself with how epilog happens and the whole pipeline for epilog operation and we are going to look at SM90 builder.inl which is the file which concern is concerns itself with how we build these epilog and how we basically build the collective epilog object. And once we are done with all of that then to summarize it all and then to wrap all of it around we are going to look at fast.co which is a repository by by prangel ssh of xai and we are going to learn about this this whole kernel for from scratch this matrix multiplication from scratch which gets 107 per% performance of kublas. So that's what we are going to learn and uh once we are done with all of that then we would get to our multiGPU lessons. So in multiGPU part one we are going to learn about specific technologies which are used to fuse multiple GPUs into a single massive execution engines. We are going to learn about the compute bottleneck of a single GPU. We are going to learn how the internal node mesh works. So how inside a a node of H100 which contains eight GPUs how everything works and what are the main innovations there and then we are going to learn about the scaling fabric that's when things would go a bit crazy and we are going to learn about you know how these these data center works with multiple scaling units of 32 DGs and how it all works. Then finally we would get to CUDA P2P and MPI and we would see how the software part of all of this works. Then we get to the last lesson of this course which is on multiGPU. And this lesson would basically deal with the orchestration and the communication software which is required to make this raw hardware act as a single execution unit. We would map how different task managers allocate these resources and we would introduce and introduce isolated processes to each other. So we are going to learn about slurm and pmix which basically are responsible for the whole orchestration of the whole cluster of hardware which we have. Then we learn about a specific primitives and mathematics which is used to move the tensor data across the whole cluster during the training. So we would start by learning the cluster orchestration. Then we would learn how we initialize the whole network and we would learn about nickel and how nickel is initialized with slum and pmix. Then we are going to learn about four major parallelism paradigms which we have. So your data parallelism, your tensor parallelism, your pipeline parallelism, expert parallelism, all these different things we are going to learn. And finally we would learn about the nickel communication primitives. The primitives which we use for doing this massive all reduce operations for these major models in AI. So this is everything which we are going to learn in this whole course. Welcome to the first lesson in our CUDA programming series for H100s. And in this lesson, we are going to learn about the hardware architecture of a single H100 GPU. Now, you might be asking Pratik, why do we even need to understand about the hardware architecture of a single GPU? And uh can't we just move on to writing code for these GPUs? Can't we just move on to learn how to work with these GPUs? The reason why we are learning about the architecture first is that it first of all gives us a good idea about how the whole GPU works. It gives us a glimpse on every single unit which we are going to learn in this whole course. But more importantly, it informs us about different bottlenecks which we have to consider about every single unit when we are writing code. And uh learning about hardware gives us a really good idea about how to write software which you know which works really well with the hardware. And especially when we are writing you know really high speed kernels which get us to really high levels of utilization then we need to understand every single component of hardware really well and how to optimize for specific components. So this is why we are learning about the architecture of our GPU before going on to the programming details of everything. So before discussing all of that, I first want to get into the specifications of the GPU. So H100 was released in September 2022 which is 4 years from now in 2026 and uh it was based on TSMC 4 and process node. It comes with hopper architecture and uh it comes in two form factors. First is a PCIe form factor which is 30 which is 300 W and the second one is the SXM form factor which is 700 W. It comes with 80 GB of HBM3 memory which comes with 3.35 terabytes per second of memory bandwidth and u it comes with 132 streaming multipprocessors. You can think of it as the heart of the GPU. So every single streaming multipprocessors has a lot of you know different cores compute cores which enable the massive parallelization which we see and uh every single H100 has 528 tensor cores which basically says that there are four tensor cores per streaming multipprocessor or SM and u every single uh H100 comes with 40 with 80 billion transistors on a custom 4 nanometer process node. So there's that. And uh the biggest improvement in the H100 GPU is the introduction of asynchronous features and the introduction to all these different revolutionary units which enable uh the latency hiding and all these new really flashy asynchronous features. So what are these new innovations? What are these new features which come with H100? The first really big and really important asynchronous feature is the tensor memory accelerator or the TMA. So every single streaming multipprocessor has a TMA unit and uh this is the new unit to offload all the tensor copy operations from the SMS. So you guys have done you know some form of CUDA programming previously and u if you remember for a lot of kernels what we do is that we move the data from global memory into the shared memory and then do some kind of computation on that data. Now, previously the way we would do that is that we would have the threads calculate some form of indexes in the data in the global memory and uh and then we would send some load instructions or store instructions and the data would you know arrive from the global memory into the shared memory u for you know whatever computation which we are doing. But with TMA what we do is that we just offload all the tensor copy operations from the SMS into this new unit which is tensor memory accelerator. And uh the reason why we have this new unit is that we want to avoid all the indexing and all these things and we want to speed up all the memory bound operations and uh this whole thing is done by multiple new ways. So we have a lot of asynchronicity features which basically means that we can now launch a bunch of operations and while they are working in background we can work on some other things and uh there are a lot more other things which we will discuss in the later sections of this course about the TMA and asynchronous P copies but for now it's important to understand that it is a new unit which is created to speed up all the memory bound operations. conditions and uh as it is written here that previously we had to do per thread math. We had to loop over every single region. We had to launch multiple instructions to work with global memory to shared memory transfer. But with the tensor memory accelerator we get the descriptors to do things asynchronously. So a single thread launches an operation and the hardware handles all the things in background and the thread goes on to work on other things. So this is how this whole tensor uh tensor memory accelerator unit works and uh we are going to learn about these things in a lot more detail in the later parts of the video. But just to get a basic intuition of synchronous copies versus asynchronous copies, I have this small slide which basically says that a copy is synchronous if a thread which issues the copy cannot make forward progress until the copy is done. So if I'm trying to load a whole you know let's say a four byte data from global memory to the shared memory then if I'm doing that using the synchronous copies then what would happen is that I would not get the control of my thread and I would not be able to launch any other instruction with that thread until the whole copy is done. And it is simple. It is you know you have the guarantee that once the data is moved then only you would get the uh control of the thread and uh it is good but the problem is that it has a lot of bookkeeping. It has a lot of you know it has a lot of complexity and uh there are a lot of you know a lot of latency because you work with a lot of different memories and uh we are going to talk about that as well. But what is an asynchronous copy? So an asynchronous copy is basically a copy where once you issue the instruction for the data transfer then the thread can work on other things when while the data transfer is on the flight and uh once it completes its work on different other things then it can come back and check if the whole operation is complete before you know using all that data. So basically instead of you know waiting until the whole thing is done we use that time while the copy is being carried out to do some other useful work and uh this is awesome because it enables overlap and uh we can basically have all these pipeline kernels where while computing on data I you can fetch the data I + 1. So it is really good for you know all kinds of latency hiding operations and uh the only problem is that now we have much more complexity with synchronization. So we need to have some way to you know know once when the copy is complete and when we can use the data. So there is a lot more complexity with the whole bookkeeping part of this process. Another really important innovation which we get with Nvidia H100 are the fourth generation tensor cores and these are also one of the most most important you know innovations which we got with H100 as compared to you know the previous generations like A100 and uh what are tensor cores they are just really fast units for doing MMA operations. So they are they specialized in doing the matrix multiply accumulate operations really fast on the data and uh we have four tensor cores per SM and uh we have the WGMMA or WO group matrix multiply accumulate instruction which enables the use of these tensor cores and uh in a W group matrix multiply accumulate instruction four WOPS work together on an MMA operation And uh that is why it is called a W group matrix multiply accumulate because a group of four WOPS or a WOB group operates on these instruction. And um we have sparse tensor codes which can do really fast MMA operation on structurally sparse data which is the data which is uh which has zeros in a predefined layout. So it is structurally is passed and uh we have instruct special instruction and support for FP8. So in A100 for example we didn't have tensor cores which supported FB8 instructions. But with H100 with the fourth generation tensor cores, we finally have the we finally have the tensor cores which support really fast FP8 operations and uh now we can have 1,979 teraflops for FP8 performance. So this is great and uh yeah this is a big improvement over previous generations of GPUs which didn't even support FP8 for tensor cores and we have the really big number of 1,024 flops per cycle for dense FP16 and 148 flops per cycle for dense FP8 which is a big number and uh yeah and the other innovation s in include the transformer engine, the fourth generation NVLink interconnect which is really good and the DPX instructions and the massive massive 50 MBL2 cache. So these are the other new innovations with the H100 and uh this is the whole way in which the GPU is divided. So we have the giga thread engine which is responsible for all the scheduling of thread blocks and allocating the thread blocks to different you know GPCs and uh SMS and uh we have eight GPCs which are basically graphic processing um cluster I think and uh we have HBM3 stacks we have HBM3 memory controllers we have PCI PCI Express 5.0 Z host interface we have NV link switch ports and hub. So NV link switches basically enable these really fast uh really fast uh communication between different GPUs and uh we have the NVLink ports and we have the NVLink hub and uh we have the L2 cache slices in the GPU. So this is uh illustration which I created to just you know have a better understanding of how a GPU is laid out the architecture of a single GPU and we can see all the different units which are there and different SM different GPCs and uh yeah so what is the giga thread engine? So it is the hardware which takes a kernel launch and it hands out all the thread blocks to SMS and uh it tracks which which u which blocks are not yet started and uh which ones are running which ones are finished and uh once an SM has capacity for launching another thread block then the Giga thread engine assigns that next thread block to that specific SM and uh it enforces different constraints. things like occupancy limits and on hopper it also understand something called thread block clusters which are another really important innovation with hopper and we'll discuss that in the next lesson. So yeah the next thing which we have is the HBM3. So we can see the six HBM3 stacks in here in the H100 SXM GPU. there are only five HBM stacks enabled because you know when these chips are being created then there are a few defects and Nvidia turns off some of the some of the features of this GPU like some of the SMS which you can see here and uh and HVM3 stack and all that. So there's that and uh we can see that HVM3 is the offchip device memory. So it is not in the main chip it is off chip memory. We have 80 GB of HBM3 at 3.35 terabytes per second. And uh the HBM3 is divided into five stacks. Uh actually it is divided into six stacks but it is disabled for yield reasons. And we have the 5120 bus width which enables the TMA to move 128 bytes of data at once. So we have a large bus width for moving a lot of data at once and um it is connected via 10 independent 512bit memory controllers. So you can see that every single HBM3 stack has two memory controllers. So for five HBM3 stacks that would be 10 H uh 10 HBM controllers memory controllers and uh yeah and uh generally when we request some data for uh from HBM3 then first of all from SM it goes to L1D or coacer then it goes to L2 cache then it hits the memory partition crossbar then it hits the memory controller and then it hits the HBM stack. And um what are memory controllers? So memory controllers are these specialized things which do a really important critical work which is basically first managing all the DM protocol. What is DM protocol? So basically the way it works is that this DM is like a huge grid of you know different cells in which data resides and uh when we request some data in you know this row and that column then what happens is that you know the whole uh so what we do is that we navigate into a specific row in that whole DM then we activate that row. So we give basically energy to the whole DM row. It is a bit silly way to understand it but it helps. So we activate this whole DM row and uh once it is activated then we can just access that data really fast. And once we access that then we you know get back and uh you know move that data into the L2 cache and supply that data to whichever SM request the data and u the job of uh you know job of memory controller is to make sure that this DM protocol works. The second job of the HBM3 memory controllers is scheduling. So they basically activate and pre-charge and they make sure of readrite timings and uh they reorder the operations to maximize row hits. So what does that mean? So basically if you know we we request a bunch of data in a non-co way where we are requesting data from multiple rows then what would happen is that the u first row would be you know activated in the HBM3 stack. It would take a piece of data from the first row. Then it would have to activate the second row and then take the data from the second row. And this is a lot more energy intensive and slow. And u the whole job of the memory controllers is to basically reorder the operations so that you know once we activate a row then we can get in maximum data from that specific row. Also that is the reason why we need you know coalaced memory access and why it is so critical because you know if we are you know activating multiple rows in the in the DM then the accesses would would be a lot more slower because we need to activate all the different rows but if we are accessing the data in a single row then we just have to activate a single row and take all that data and then move it. So this is one of the reasons why you know it is really important to it is really important to have collaged memory accesses to access the data in a single row because it is much more faster and much more energy efficient. And uh yeah and the third job of a GPU memory controller is to map all the addresses to whatever channel banks and rows and columns which we have in the DM. So yeah that's how the HBM3 works. The second really important component of the GPU which we need to understand is the L2 cache. And uh this is also like one of the really important you know one of the really important components of the H100 GPU. And uh how you utilize the L2 cache um you know tells if you are going to get really high performance in your kernels or if you are going to get really bad performance. So it is really important to understand what L2 cache is. So we have 50 MB of L2 cache and it is divided into these 25 mgabytes of part partitions and uh these partitions are connected to a crossbar and u every single GPC is connected with some part of L2 cache and you can see that some GPCs are farther from an L2 cache strip than other GPCs and this is also one of performance defining factors in a GPU. But basically SM in a given GPC have closer part to an L2 cache partition which they are directly connected to. And uh when you access these partitions then um you mostly hit lines that live in the nearby partition. And when you hit the lines which live in the nearby partition you get better performance, better latency or bandwidth. And if you hit uh the data which or hit the lines which live in the farther partition then you get worse. And uh you have 128 byt lines and 32 byt sectors. So whenever ns requests some data the L2 cache move that data in either in you know sectors of 32 bytes or combined line size of 128 bytes. So if you know it is like um if you access like two bytes of data from HBM3 then the whole GPU won't move that single two byte data into that SM but what would happen is that it would it would move the data into a 32 byt sector and it would move that data from the HBM3 into L2 cache and uh it would move some other data along with that so that you know if you need that other data then you can access it. But the point is that whenever the data is moved through L2 cache, it moves in 32 byt sectors. It moves in a packet of 32 bytes and a memory request can touch one to four sectors of a single 128 byt lines and you can be coal in a way that explodes all these sectors or request. So there's that. And um the another really good feature of L2 cache is that it absorbs and coalleses messy small writes from the SMS and it turns them into clean large efficient rights into HBM. So, so then we have the GPC or the graphics processing cluster and this is a bit funny but it has nothing to do with graphics processing in the H100 because generally H100 is not used for processing graphics but anyways it is just a group of 18 streaming multipprocessors. So in this whole illustration, it is the thing which holds up the group of 18 streaming multipprocessors. And uh basically this is how this whole thing works. And uh every single GPC is connected to its own chunk of L2 cache as we can see in this in that illustration. And uh whenever there is a request to fetch something into the shared memory, the data flows from HBM to L2 to L1 and every single GPC handle that stuff. So that's how this whole thing works. And uh GPCs also enable the use of distributed shared memory between SMS. So the only real you know significance of a GPC is that you cannot have a thread block cluster outside a single GPC. Every single thread block cluster which you create is inside and is constrained to a GPC. So what is distributed shared memory? What is a thread block cluster? Basically for now you can imagine it that you know in the previous generations we had a thread then a thread block which was you know multiple threads and then we had a grid which had multiple thread blocks. But with hopper we introduced a new hierarchy which is thread then thread blocks and then thread block clusters which are basically a group of thread blocks inside a grid. So it's threads, thread blocks and then thread block clusters and then uh a grid. And the main features of these thread block clusters is that they pull their shared memory together. So every single thread block in a cluster can access any other thread blocks shared memory and uh that is the main feature of the thread block cluster and it is called distributed shared memory. So GPCs are the units which enable the use of this distributed shared memory between the SMS and uh there is no distributed shared memory outside a GPC because you know you cannot have a thread block cluster outside a GPC. So there is that and uh then they have the TPCs which hold two SMS and their job is to only make sure that the communication between the SMS inside the TPC are really fast and this is the connection for the distributed shared memory. So TPC basically stand for a texture processing unit and uh a single TPC contains two GPU uh two SMS and uh the job of TPC is to basically have a shared SM block where the communication between these two SM is really really fast. So if we um create a thread block cluster with these inside these two SM then the communication between you know the these two thread blocks uh the thread blocks inside the thread block clusters in these two SMS would be really fast because inside a TPC the communication is really fast. So yes there is that and uh yeah so that's basically all about a GPC. We are going to learn about SMS a lot in a few minutes. And yeah, then we have the NVLink 4.0. And uh this is arguably the most important feature because NVLink is what enables interGPU communication. And uh the reason why H100's are used in you know in AI is because we can connect a bunch of H100 GPUs with NVLink to get you know really high speeds of data transfer almost like the speed like 800 MB 800 uh Gbps and uh 900 Gbps and crazy speeds like that are enabled by NVL link and uh yeah so they give 900 Gbps of birectional uh GPU toGPU bandwidth and uh the these are basically organized into nine sublinks. Every single link provides 100 Gbps of birectional bandwidth which basically means 50 Gbps per direction and uh it provides us with PAM force signaling which is basically using four different voltage levels instead of two. So instead of just having two voltage levels to uh share the data, it has four different voltage levels. And even though this increases the amount of noise because now you have you know four different uh different voltage levels which can you know uh create some confusion but still there are things to mitigate this and uh basically what happens is that we double the bits per second between GPUs uh which are connected through NVLING. So this is one of the biggest innovation which with NVLING 4.0 zero and uh you can now directly connect eight GPUs and uh use something called NV switch fabric to connect to up to 256 GPUs and uh we'll learn about them when we will discuss some of the multiGPU stuff but basically we can directly connect to 8 GPUs and uh we can use NV switch to connect to a multiple GPUs and uh we can also use TMA uh for bul copies inside a single single DGX node. So basically what happens is that inside a single DGX which is just a whole machine with with eight H100s connected together using NV link and NV switch and what we can do is that we can use the tensor memory accelerator to send the data between two GPUs in you know which are there in a single DGX box. So we can send the data between these eight GPUs using tensor memory accelerator and uh yeah so finally we have the time to talk about the SMS. So we learned about the HPM3 stacks which are here and we learned about the L2 cache and the way you know the data is provided from HBM3 to different SMS and uh we learned about NV link we learned about GPCs we learned about the gig thread engine and uh now it's time to learn about these SMS and what they do so this is the way a single SM looks like and uh basically if I zoom out a bit then you can see that a single SM contains four identical uh four identical subpieces. So basically these are called SMPP or SM subartition and uh basically a single SM has L1 instruction cache. It has u tensor memory accelerator. It has the unified L1 data cache and shared memory. It has some texture processing units and then it has four SMSPs with all the compute cores, pop scheduleuler, dispatch units, register files, L0 cache, load store units and the special function unit. We are going to talk about each of them in a lot more detail. So basically the SM are the fundamental execution unit of a GPU. These are the units which are responsible for all the computation. The they have all the computational cores. They have the tensor memory accelerator and uh basically these are the different components which exists in a single SM. So there are FP32 CUDA cores and there are int and FP64 units. Then there are the tensor cores. There are shared memory and L1 cache. So we have the unified shared memory and L1 data cache. There is L1 instruction cache. Then there are scheduulers and we have the dispatch units, register files, L0 instruction cache and we also have the tensor memory accelerator in every single SM. So there is all that and uh every single SM is divided into four identical subi divisions called quadrants or SMPS SM subartitions and uh every single quadrant have 16 FP64 units 32 FP32 CUDA cores 16 INT32 units a single tensor course eight load store units and four special functional unit and uh apart from these This every single quadrant have their own scheduulers, L0 cache, dispatch unit and register files. And uh every single quadrant can dispatch instructions to its own local units at every single cycle and it can juggle up to 16 Ws per quadrant. So okay. So let's now talk about the general compute units which we have apart from the tensor course in an SMS. And basically we have three different types of general compute cores. And uh the first one of them are the int or int32 cores and then there are the FP32 cores or CUDA cores and then there are FP64 cores for you know different purposes. So intunits are the general integer alus for memory addressing loop controls and integer math. And uh we have 64 in 32 cores per SM. We have 16 of them in an SMS and uh they can execute simultaneously with all the floatingpoint data paths. So we are basically it allows us to you know have things like computing memory addresses while processing the FP32 data and this all happens without installing because they execute simultaneously. And uh we have total of 8448 units in our SXM5 GPU. So these are the int 32 units. And uh that's how they work. And uh yeah, so they are just used for all the integer operations, all the logical operations on the memory addresses and bit manipulations, all that stuff. And uh yeah the next units which we have are the CUDA cores and uh basically they are the FP32 floatingoint units and uh they are primarily for the general purpose graphics and compute sharing but we don't use these a lot more right now with all the AI operations generally we are using tensor coursees for doing most of the MMA operations but previously we we used to use these CUDA coursees for all kinds of stuff and every single SM has 128 CUDA cores in H100 and every single SMS SMPP has 32 CUDA cores. So basically 16 and 16 32 CUDA cores per unit and uh yeah basically they handle 32-bit floatingoint calculation which are essential for most of the simulation and it is the default precision for general compute tasks. So yeah and the last one are FP64 units. No one uses FP64 a lot now apart from some applications in you know computing in scientific computing or whatever they are not used a lot but they are generally you know for extended range and precision which is required for some of the scientific computing and uh complex numerical simulations and financial engineering. So that's what FP64 units are used for. And uh yeah, so there are not a lot of uses for FP64 units. So these are the three major integer cores which we have in H100. In 32 units, we have 16 of them per SMS, 64 of them in an SM. FP32 units which are used for the general FP32 computation. we have uh 32 of them per SMS and 128 of them in a whole SM and then we have the FP64 units which are responsible for all the FP64 computation and we have 16 of them in a single SM. So yeah and now we are going to just learn about all the different units which are there inside a single quadrant. So let's go into the SMS part of our quadrant and uh the first thing which there is to talk about is the special functional unit. So SFUS and we can see that here is the SFU and uh every single SMS has four SFUS and we have 16 SFUS per SM and u a single SFU can issue a single instruction per cycle. So we get 60 16 16 operations per cycle per SM with our SFUS and their job is to just handle all the complex mathematical functions like s cosine logarithms exponentials square roots and reciprocals which would be a bit computationally expensive on you know tensor coursees or cuda coursees and uh these sfus sacrifice some precision for massive throughput gains. So they use polomial approximation and lookup tables combined with some interroporation and u if sfus is stall then you know if every single thread is calling math functions then sfus might stall. So it is you know it is wise to use sfus in you know in a better way in a more planned way. So we have things like flash attention three which uh you know which overlaps the matrix multiplications with with soft mac max operations. So there are different ways to you know to avoid the sfus stalling with you know being overwhelmed by all the threads calling sfus for doing exponentials or whatever. And uh yeah then we have the load store units which are here and we have basically eight load store units per SMSP and uh they execute per thread memory instructions like loads stores and atomics and um every single quadrant contain eight dedicated load store units. So we have total 32 LSUs per SM. And uh these LSUs directly connect into the L1 and the L2 cache. And uh so when your is executing some load store or atomic operation then what happens is that LSU coaleses the 32 thread addresses it forms a cache line and it queries the L1 data cache for you know that load or that store. And uh if it hits then the data comes back quickly into the register. And uh if it misses then the request just goes downward to L2 cache or DM cache. And uh it may fill back into the L1 cache for you know any future future use of that data. And uh when accesses are coal then LSU merges the WOPS 32 addresses into fewer fewest possible cache lines. So it can issue fewer instruction to get more data and this basically reduces the replays and stalls and it boosts the effective bandwidth. So there's that and uh this is where you know the benefit of uh using something like tensor memory accelerator comes into play because what LSU does is that there is a lot more scheduling there is a lot more you know accessing the L1 cache and then accessing the L2 cache than going into you know the DAM and all that. But what tensor memory accelerator does is that it bypasses the whole L1 data cache path. it directly goes into L2 cache or into the DAM and it takes the data and it puts back into the shared memory. It does not even touches the L1 data cache. And that is one of the important things to understand about the tensor memory accelerator. And then we have the unified shared memory and L1 data cache. It is 256 kilobyt in size. So here is the shared memory and L1 cache and uh yeah so it is 256 KB in size. There is 33 tabytes per second of bandwidth per SM and uh the whole shared memory is divided into 32 banks and every single bank is four bytes wide and uh the shared memory size which we can use is 228 KB and uh the per thread block limit is 227 KB and uh 1 KB is reserved by CUDA for some other applications and There is 128 bytes per cycle which we get for loads and stores when we are working bank conflict free. What is bank conflict? We will learn these things in a lot more detail but you know the thing is that the whole share uh the whole shared memory is divided into 32 banks and you can think of them as like small gates into a stadium. So let's say there is a massive football match and uh people are going through these gates. Now if you know the whole crowd goes into a single gate then there would be a lot more serialization. So a single person would only go through the gate at once and u you know the rate of people filling the stadium would be slow. But if we have uh 32 banks, if we have all the 32 gates open and people go from each gate into the stadium, then the stadium would fill really fast and the accesses to the stadium would be really fast. So banks are like that. So when you are requesting data from a single bank and you have multiple requests then um basically every single request would serialize and uh this would become really slow for accessing the data and uh this whole thing is called bank conflict where you access the single bank uh in every single access. So whenever you try to access the shared memory you go and hit the same bank and that just serializes your accesses and makes everything slow. So when we are working with you know so when we are working with shared memory and uh if we you know mitigate this bank conflict then we can store 128 bytes of data in a single cycle for uh stores and can also load 128 bytes per cycle and u when we are you know using tma async copies with something called swizzling which is a way to avoid these bank conflict then uh we hit the near peak bandwidth and we also have the option to overlap the compute with the data movement and uh banks generally interle consecutive four byte birds across sequential banks. So yeah there's that and uh the next thing is about L1 cache. So L1 cache acts as a coaling buffer. So if there are any you know rights through the load store units then they first get you know stored in L1 cache and they then they are actually you know enforced and uh what happens is that if you know is requesting some data then it gathers that data into that L1 cache buffer and then it delivers it effectively and that's how the L1 data cache works and uh a single cache line is 128 bytes a single sector is 32 bytes and just like we learned with L2 cache the cache can be filled sector by sector and uh yeah maximum configurable shared memory capacity is 228 KB per SM but we can only use 227 KB per SM. So yeah then there is some more information on shared memory. So when we are dealing with some address in shared memory then that specific byte address will contain information about banks. So which bank we are using and uh bank rows which basically tell us about the rows inside that bank and uh generally when we are using using static shared memory then uh it is limited to 48 KB to maintain compatibility with older compute capabilities. But if you want to exceed this limit of you know 48 KBs in shared memory you must use dynamic shared memory using extern shared keyword and uh optimal tile size for the shared memory is typically 64 to 128 KB per block for massive gems and uh when we are working with massive gems that then these transfer might hurt the latency heading. So generally we use you know smaller tile sizes and uh once you increase the shared memory then you are shrinking the space left for L1 data cache and vice versa. So that is also one of the things which one should keep in mind while writing the kernels. Okay. So now let's talk about registers and uh as you guys might know that every single thread gets a private set of onchip registers and uh these are the private set of registers. So that basically means that you know uh no other thread can you know read or write on a thread's registers and uh they have the fastest memory bandwidth and they have the lowest latency and they are the most power efficient uh memory. So it takes a lot less power to access and work with registers. And uh we get 256 kilobytes of registers per SM. And um we get up to 128 reads per cycle per SM with the registers. And uh if you think that this 256 KB per SM is a huge number, then you are wrong. registers are generally a really frequent limiter and um if your kernel uses R registers per thread then maximum resident threads would be 65,536 by R. So we can have maximum of 65,536 registers in a single SM and uh because every single SM is you know 32 bits. If you multiply this by uh four bytes then you get to this number and uh yeah so maximum resident threads if we just look at the you know the total number of registers which we have would be this number and uh then you also have you know block granularity and all the other limits which which you know which takes registers. So the actual number could be much lower than this. And uh if we have you know 128 registers per thread then you will be basically limited to 512 threads per block. And uh this is one of the things to take care of. It's really important to track how many registers you are using and uh to make sure that you know you are not overusing and uh you are doing things efficiently. And uh 32-bit is the base unit for these registers. So every single register is 32-bit. And uh the hardware register file is effectively 32-bit slots. And uh when we are working with things like FP16 or FP8 then what we do is that we often just uh we often just uh pack these data types into registers. So if you're working with FP16 then we would pack two FP16s into a single resistor and uh this helps a lot and uh yeah and as discussed earlier that they have they are a lot more power efficient and uh they are 30 to 50x more power efficient than shared memory and they are thousandx more power efficient than HBM 3xs. So that's a really really huge number and uh this thing comes into use a lot when we are working with our kernels and yeah and uh register usage is determined at compile time and not at runtime. So this is an important point and because it is determined at compile times when compiler runs out of register then it basically spills all the data you know rest of the data into the local memory which is much much slower than registers and uh this problem was kind of mitigated by GUDA 13.0 zero which basically added the support for spilling into the shared memory first and then only falling back into the local memory if the shared memory is insufficient. And the last point and really important point is about occupancy. And uh basically what it is saying is that even if you change you know the number of registers in a single thread by a small number then uh what might actually end up happening is that you may drop the number of resident WS because allocation is rounded to an internal granularity and basically you would have to deal with a much less occupancy and you will see fewer active Bobs and this would lead to lower performance. So what am I talking about? So basically the thing is that if a W tries to you know tries to get some resistors for doing some kind of operation and if it does not get those resistors then it might stall and u the thing is that once the is installed then the hardware is idle and we cannot do a lot of work. So basically if you know if the resistors which we demand gets a little bit high then what would end up happening is that you know the number of active WS which we have. So the number of WS which are actually executing and everything would decrease and u that would basically result in the in lower performance. So this is just a simple math which basically says that if R are the registers per thread, T is the number of threads per block and res SM is the total number of 32-bit registers which are available in a single SM. Then uh total resident blocks from that register would be total number of regist uh registers which are available in that specific SM upon the total number of threads in the block times the number of uh the number of registers in that uh which are you know which are owned by that thread. And uh this would be basically the total number of res resident blocks which we can have from the point of view of you know the total number of resistors and um basically the number of active bops which we would have would be the total number of resident blocks which we have times total number of threads in that whole block divided by 32. So if we have more number of registers per thread then basically this whole total number of resident block thing would decrease and if this decreases then the number of active W decreases and if we have lesser number of active WS and more WS stall then we would have much less performance. So register usage is something which one should really take a look at while writing really high performance kernels and yeah and then there is some talk about register layout. So basically tensor core operations require the memory to be in a specific location. So A can be stored either in shared memory or registers and b must always be stored from shared memory and never from registers. And uh once we you know do the operation and we get the result then that whole result is distributed across four whole ws in the resistors. And uh yeah so basically we require a lot of um resistors for the WGMMA operation and uh this is one of the things to look out for when we are working with registers and uh when we are worried about the occupancy which we might have and uh this is just a small point that if we are working with FP8 or FP16 then WGMMA would generally expect the data to be pegged in register. So 2 FP16 in a single 32-bit register or 4 FP8s in a single 32-bit register. So that's how basically things work. And uh especially with WGMMA, you must uh really worry about you know what uh registers you are using and stuff like that. And then we have the L0 instruction cache. So in our SMS we have the register file and then we have the L0 instruction cache and as it is written here that it stores the recently fetched instruction for immediate reuse. So basically it acts as a micro buffer for the whole instruction stream. It basically filter requests to the larger L1 cache and it allows the scheduleuler to issue instruction at really fast speeds and um yeah basically at the speed of hardware without saturating the shared memory bandwidth. So that is the main use of the L0 instruction cache. We don't have a lot of con control over what goes into L0 instruction cache but basically it is very small cache which is designed for a speed. So yeah that is written here and uh it can hold a very small number of instructions in it and uh you know when we are working with something like aggressive loop unrolling or if we are inlining a lot of function then uh what might end up happening is that we might exceed the L0 cache size. So this is one of the things which we have to keep in mind as well. And uh yeah code running on SMSP0 cannot use L0 cache of SMSP1 because you know the L0 cache for every single SMS is isolated. So there is that as well. And the next thing which you know comes up is the and the scheduleuler. So what is a You guys probably know what a is. You are CUDA programmers. So, uh basically a is a group of 32 threads which are created together from a thread block and they share a scheduleuler's attention on per cycle basis. So what is a scheduleuler? It is just a unit in SMSP. We will talk about that in a lot more detail. And uh a have some private registers. So every single thread in the W have some private registers but they generally have a shared instruction string and uh you might think why we have 32 threads in a single bob and uh the reason for that is that you know if we have a lot of threads in a single W then basically it would be really difficult to control every single thread and then you know in the cases of you know thread divergence or W divergence then you know things might get out of hand and uh you know and with 32 threads you know we can do satisfactory work per instruction. So it basically comes at a sweet spot of having control of all the threads and you know having a good number of good amount of work done per instruction. So this is one of the reason for having 32 threads. Another reason is that it just you know a lot of common access patterns become naturally nice when we have 32 threads in a warp and it works really well with memory coallesing lines and uh yeah it it works really well with accessing memory and working with all that stuff. So this is another reason why we have 32 threads in a single That brings us to the WOPuler. What is WOPuler? It is just a unit which manages all the instruction issues. So basically when we launch a kernel then every single uh WO scheduleuler is assigned a few WOPS and um what the job of scheduleuler is is to take a look at all these WOPS and track their uh their completion and u make sure that you know if a W is ready for an instruction then it is handled that specific instruction. So it manages all the instruction issues. It selects all the ready WS for instruction issue every single cycle. So scheduleuler issues instruction at every single cycle and uh that is one of the reasons why L0 instruction cache is really important because it enables the WO scheduleuler to um fetch the instructions really fast to issue them really fast and uh there's that and uh it also handles the W level branching and divergence. So basically if you are trying to write some W divergent code then basically what it will do it is that it will mask some of the threads and uh it will handle all the things related to W level branching and divergence and independent thread scheduling which is another really interesting concept which if you want to learn you can look it up and uh it also manages a scorecard which tracks which are truly ready. So basically once it assigns some you know some uh to do some kind of job then it manages a scorec card which tracks the completion of the instruction in that VO there are I think u there is a fixed uh number of cycles which every single operation takes and uh you know in this case we are not talking a lot about asynchronous instructions but you know if I'm loading something or if I'm adding two values then the cycles which are required for the completion of that instruction is you know pretty evident and uh basically what scheduleuler does is that it uh it uses this math to make sure that once the is ready then you know the next instruction can be issued. So uh there's that and u it ensures that the source registers are read from the register file or forwarded from the pipeline before issue. So basically it manages that uh it makes sure that the resources which are required for whatever instruction which we are launching are available to that pop and uh everything is good to go and uh it also takes cares of you know structural limits like how many long latency ops within a warp we can have in flight pipeline availability for execution etc. So there's that as well. And uh the next thing is about how the WOPS execute. So basically the way it works is that the scheduleuler scans all the 16 Ws which are assigned to the SMSP and based on its scorecards and everything. It picks up one which is ready and then it fetches the instruction. It selects the ID and it selects the program counter to dispatch unit and uh it basically sends that instruction to the dispatch unit so that you know the dispatch unit then can route them to ex appropriate execution pipelines like CUDA course, tensor course and everything. So it fetches the instruction, it selects the and then it sends the instruction to whatever you know whatever execution pipeline there is and uh once that is done then it waits for the dispatch feedback which tells if the execution is done or if the instruction is issued and uh once that is done then basically it works for the next instruction and uh the Wululer would always try to launch as many instruction as possible. So every single cycle the scheduleuler would try to launch an instruction and u it would try to uh turn all the free WOS busy at every single uh every single cycle and uh there's that and uh yeah this is an important point. So if we want to maximize the hardware utilization then you generally want your block dimensions to be multiple of four WS for you know whenever you are working with things like WGMMA so that the whole work is evenly balanced across the four sub partition. So basically the thing is that you know if you are looking at NSM then NSM has four SMSPs and uh in every single SMS there we have a single scheduleuler. So when we when we basically uh when we make you make make our block dimensions to be multiples of four WS then the work is evenly balanced between all these scheduulers and things work really well that way. So it is generally advisable to um have your blog dimensions to be divisible by you know uh 128 threads and all that. So yeah and uh once the uh once the instruction is launched from the scheduleuler it goes into the dispatch unit and uh what the dispatch unit does is that it is responsible for the physical act of sending the WS operation to appropriate functional unit and uh this happens at every single clock cycle and this is done by issuing instruction to various execution pipelines. things which exist and uh basically every single execution unit. So these int32 units, FP32, CUDA codes, FP64 units, tensor course, load store units, all of these have a fixed execution pipeline which routes them from uh the dispatch unit. So what this whole thing does is that it sends all those instructions to the physical data paths to whatever functional units there are via dispatch ports. Dispatch ports are like you know the ports through which these instruction are routed to different execution pipelines and uh when we are working with shared SM resources like uh LSUs, SFUS or even shared memory then uh the dispatch unit sends instruction directly to the execution units dedicated instruction queue and every single execution unit has an instruction queue and dispatch unit sends the instructions directly to that queue. And u the next thing which there is is the dispatch port which is the physical connection point through which the dispatch unit sends the instruction to a specific execution pipeline. So it is the pathways to tensor cores and these FP64 units and FP32 units and IN32 units. All of these units have specific pipelines to which we send the instruction and uh multiple execution pipelines compete for the same dispatch port. So they cannot receive the instruction simultaneously. So they try to compete and to get the instruction and all that. And u only one instruction can use a shared dispatch port per clock cycle. So if there are multiple um instructions then the dispatch unit must serialize the issuance. So there's that and uh the port is only needed for instruction launch. Once the instruction is done then basically we send the dispatch feedback and uh you know the scheduleuler can launch other instructions as well and the port is not needed for the full execution duration. So yeah and the next uh part which we have next cache which we have is the L1 instruction cache and u we have a L1 instruction cache for the whole SM and uh its job is to basically buffer all the executive machine code which is the SAS instructions to keep the scheduulers constantly fed. So that's the role of L1 instruction cache and it decouples execution units from the memory hierarchy. So without it the SM would have to fetch instructions for L2 cache or HBM3 and uh that would take hundreds of cycles. So basically we just have a L1 cache for doing all that and uh it is optimized to handle jumps, loops, function calls, all of these things and uh when the kernel loops the instructions are served directly from L1 instruction cache which saves the bandwidth and power and a lot of time the the instructions from L1 instruction cache goes to L0 instruction cache for you know for being more accessible to the course and the execution units and the WOB scheduleuler. So yeah and uh the next thing is about how thread blocks are distributed across SMS and SMSP. So basically your giga thread engine selects the SM which has enough resources to handle your thread block. Then the hardware does not keep a single block inside a single SMSP unless you know the block size is really small. So what happens is that the hardware divides the block into WOPs and then assign these WOPS in a roundrobin fashion to every single SMS and uh when we use things like sync thread or bar dot you know whatever then um the synchronization is handled by the barrier logic in the shared memory where all the SMS share. So basically all the synchronization at that point is handled by the logic in the shared memory and u other WS are put into sleep until every single W have completed their instruction. So that's how basically sync threads and the thread block synchronizations work. And the last part which we have is the connection to the CPU the PCI Express 5.0 0 host interface and this is for CPU's connectivity with the GPU and u basically it acts as the primary 128Gbps data highway which connects the H100 to the host CPU and the system memory and uh when we are working with like large clusters then it also serves as the critical phys physical bridge which connects the GPU to the network interface cards and you can think of them as you know as the cards which connect this GPU to all the other GPUs. And it also enables GPU direct RDMA which is basically a fancy way of saying that you know it enables the GPU to communicate with other GPUs and other nodes uh without you know the interference of the CPU and uh yeah that's what it's saying I think. So it allows the network cards to read the GPU memory directly without burdening the CPU. So it is a really important part of the H100 and uh yeah that's how it works. So just to conclude the defined characteristics of H100 is the shift to fully asynchronous execution and uh we moved from synchronous copies to using the TMA to handle the data movement u which allows the SMS to focus entirely on compute and other stuff and uh now the architecture relies on a specialized unit for every single bottleneck. So, TMA for memory bandwidth, fourth generation tensor cores for all the MMA stuff, SFUS for complex math. So, there's that. And uh the data can now flow efficiently from the HBM3 to L2 cache to shared memory to resistors. And uh yeah, that's how the whole thing works. And uh we learned about every single thing which there is to learn about the H100. And uh we learned about NV link. We learned about HBM3. We learned about L2 cache and GPCs and all these you know really cool stuff. And we learned about how a single SM works. We learned about all the execution units. We learned about tensor coursees. We learned about a lot of stuff. And uh with that I would like to conclude this lesson. And uh yeah, welcome everyone to the second lesson in our CUDA programming series. In this lesson, we are going to talk about thread block clusters, data types, inline ptx and pointers. Now these concepts are really important for the rest of the course. The rest of the course we can say will be built upon these fundamentals. So it's really important to grasp them really well. So let's start this lesson now. What are thread block clusters? So if you have programmed in CUDA before, you know that there is the traditional CUDA programming hierarchy which consists of threads, thread blocks and grid. So a thread block will contain multiple threads and a grid would contain multiple thread blocks and this is how things will work. And this was great for a long time. But then we get into a bottleneck which is that the threads in a thread block can only access the data in the shared memory which is allocated to that thread block and if it wants more data then it needs to go to L2 cache or worse to HBM which is expensive. So we solved that using thread block clusters. So what are thread block clusters? So instead of having this system where each thread is basically operating on its own allocated shared memory, we get this kind of system where we group a bunch of thread blocks into a thread block cluster. And now a thread can access the shared memory which is allocated to any thread block cluster. And this is the main feature of the distributed shared memory it's called. And this is the main feature of thread block clusters. So basically each thread block was previously constrained by the limited shared memory and the limited compute and u in operations like top k or matrix multiplication where we needed a lot more data we couldn't use the data from different thread blocks and uh this was solved using our distributed shared memory in thread block clusters. So what do we do? All we do is that we take all the thread blocks and we let all the threads in the thread blocks which are in the same cluster use the shared memory of other thread blocks and uh this pool is not pool like we are pooling resources but we are not doing it in the sense that you know all the thread box just pour their shared memory into one single thing. Each block still owns its piece of shared memory. The difference is that now you know a thread can go out and use other threads shared memory. And generally this thread block clusters and this group of threads are constrained inside a GPC. We discussed what GPC is in the previous lessons. So yeah, now there are a few important points about thread block clusters. A thread block is still runs on a single SM. We don't have now a thread block which can run across multiple SM and you know anything like that one to one relationship between an SM and thread blocks is still not changed. Now an SM can run multiple thread blocks simultaneously. Yeah. But clusters are software concepts. GPC is a hardware concept. Now what does that mean? So we have GPCs and TPCs. GPCs are basically a group of 18 SM and TPC is a group of two SMS and the SM inside a TPC are interconnected with really fast data transfer and the data transfer can happen really quickly and the other thread blocks inside GP other SMS inside GPCs can also communicate with each other but then the transfers are slower. So clusters are a software concept. We cannot there is no such thing as you know a cluster which is hardcoded in hardware like we cannot have clusters which actually manifest themselves in hardware. We still have them as a broader concept and we basically just use the interconnect between different SM inside GPC to create these clusters. And the last point is that there is no automatic assignment. You need to explicitly define thread block clusters in code and you need to use your brain in order to you know define the cluster size and all that and then you need to write it in your program. There won't be thread blocks which will thread block clusters which will be created automatically by the compiler or anything like that. Great. So now let's talk about cluster sizes. So as discussed earlier the thread the clusters inside the the SM inside TPCs are created to have really quick transfers. So this is a really important part. If we have a really large cluster size, let's say we have eight thread blocks or we have 16 thread blocks, then what might happen is that the communication between the thread blocks will be slower and uh that is that will cause some problems. So we can have a cluster size of 16 threads thread blocks but then what will happen is that you know we can have around 64 thread blocks inside an SM and if we are having a cluster size of 16 thread blocks and they span across multiple SMS then what might end up happening is that a few SMS get idle and we are we might not be able to use our whole GPU. do efficiently. So generally what is ideal is a cluster size of two or a cluster size of four because it's really easy to communicate and it's really optimal and if we have a larger cluster size like four plus cluster size then it costs us around 87 cycles for communicating between them and 150 cycles for a cluster size of 16. So be aware about cluster size. They come with a cost. Great. So now let's talk about distributed shared memory. So distributed shared memory is basically a feature in H100 which which is basically the main feature of thread block clusters that we can now have um all the threads inside all the thread blocks inside cluster pull their shared memory so that any any thread in any thread block in the same cluster can use that shared me memory and uh as talked before H100 implements a dedicated SM2SM network for these clusters so that we can have faster loads and stores operation and all that and uh distributed shared memory can also be used simultaneously with L2 cache accesses. So we can use both of the things and we can get the combined bandwidth. So that's great. Now let's talk about multiccast with TMA and distributed shared memory. So we can use TMA for asynchronous copy operations with distributed shared memory. That's one of the really cool features and uh what it basically means is that instead of you know getting back and forth between global memory and shared memory what we can have is that we can fetch a data from global memory to our SM and then what we can have is that we can distribute this data to multiple SM. So you can have something like a data going like from global memory to shared memory and then it gets distributed between all the SMS. So this basically it speeds up and reduces a lot of latency of going back and forth and you know each SM getting the data same data from global memory which is really expensive and uh what happens is that a single thread per cluster issues a TMA operation and basically that TMA operation just happens using you know asynchronous hardware and it just works. Okay. So now let's talk about how to program these clusters. So we can define cluster dimensions at compile time using this syntax and uh we use this in our kernel declaration like this. So yeah and uh once we define it then we can just launch the kernel normally. But an important part of this is that the grid dimensions must be a multiple of cluster size otherwise you know we might have some problems with the kernel launch. Okay. And there are a few cluster handlers. So we have cooperative groups. These are all C++ handlers. So we can basically use cooperative groups to do that. And uh we can basically point to a single cluster using this and uh we can create basically this creates a handle to thread block cluster and then we can operate on this cluster. So we can use something like cluster.mmap.shared rank and uh using this we basically get the remote address of that you know that target thread block and you know the the memory address of the element inside it. So if we basically have some kind of an array inside that inside that uh inside that thread block in that cluster then we can use this to basically get a address which we can you know we can use in other thread blocks to access that value. So basically what we are creating is a global memory address which can be used to access the shared memory of any thread block inside a cluster. So we all know that generally how it works is that if we have a shared memory address then only the uh threads inside the thread block can make sense of that address and we don't have like shared memory address for which is which can be used by all the threads in a grid. So by using uh by using this map shared rank thing what we can have is that we can have a memory address to the to the shared memory of the target block inside that cluster which any other thread inside that cluster can use to access the data in the shared memory. So that's how it works and then we can get you know cluster size using this and cluster rank using this. So yeah, we have some important things in here. Okay. So let's now talk about the representation of these in PTX. So we have PTX special resistors to represent these and uh basically we have cluster CG ID to represent the CTA ID within the cluster. So this is basically in the form of X Y and Z coordinate of the CTA in the cluster. Then we have cluster nct ID which basically gives us the dimensionality of the cluster. So basically this whole thing it gives us it returns us this and uh then we have cluster CTA rank which gives us the linearized CTA rank in the cluster. So what we basically get is the is the x coordinate of the CTA plus the ycoordinate of the CTA times the x dimension of the cluster plus the z coordinate of the CTA times the y dimension of the cluster time x dimension of the cluster. Then we have cluster and rank which basically gives us the total number of CTAs in the cluster. And then we have is explicit cluster which is basically a boolean. And uh this represents uh this basically boolean tells us if we are explicitly launching a kernel using uh explicitly launching a cluster using this or we are just using the default cluster dimensions. Uh yeah or if we are just launching a one by one by one cluster which is not a cluster like okay now let's talk about PTX. So what is PTX? PTX is basically the assembly language of Nvidia GPUs. Why is PTX? We use PTX why we are learning PTX. We are learning PTX because we want much more control over the code which we are writing. So if we want to really optimize a lot of our operations then we need a lot more control over how our code is written and how the code is executed and everything and for that we need PTX. So what is PTX? PTX is the instruction set architecture. So basically when CUDA C++ files are compiled the then they translate into PTX when they are compiled and then PTX is compiled into shader SAS and uh then SAS is basically assembled into binaries. So yeah, we are learning PTX because first of all it gives us a lot of low-level control and because it is widely used in almost all the famous repositories which use CUDA kernels and uh the reason is again that they give us a lot of control. So yeah and most of the times we won't be you know tinkering around with actual PTX code. We will be using something called inline PTX which is basically a way to write PTX inside our CUDA C++ files and basically embed PTX code in the CUDA C++ code and uh this basically takes a lot of weight out of our shoulders. So we can basically leave a lot of heavy lifting to the compiler and we can write PTX code where it is needed. So that's one of the reasons why it is really important. Then we have templates. Templates basically enable us to embed the compile time parameters. So we have template parameters which during compile time are replaced by actual values. And uh what we can do is that we can have ptx instructions which use these template parameters and when the C++ file compiles they are placed there. So yeah. Okay. Ne next is the format of ptx instructions. So we use the ASM keywords to insert PTX code in our C++ CUDA C++ program. And uh we can use the volatile keyword after ASM to basically prevent the compiler from manipulating the PTX code or you know reordering the code in any way. So basically what happens is that when if we are not using volatile and when we compile the CUDA C++ files then what might happen is that your PTX code might be reordered which we definitely don't want or worse it might be deleted. So yeah and uh yeah basically that's how it works and uh the general format of a PTX instruction is this. So we have the ASM keyword. Then we have the PTX instruction which is separated by output operants. What are the output operants? These are the things which get we get as outputs after the PTX instruction is you know launched and everything. You have input operins which feed into the PTX instruction. And then we have clobers which are really interesting pieces of things. We'll talk about them a lot later. Okay. So what are input and output operants? Input operants are basically the inputs which we give to the PTX instruction. So if you will see the normal way a PTX instruction is written and uh maybe we have an example. No we don't. Okay. So the manual way a PTX instruction is written is that instead of having all these you know the inline PTX format we basically have this instruction and after this instead of these placeholders we actually have the registers in here. So this is a really weird way to write it. Not actually weird way but it's a bit complicated to write it. we have our C++ variables and uh we can just use them instead of you know having you writing all these things by ourselves. So input and output operations are basically a simple way to do them do that whole thing and uh yeah so output operands are declared first as you can see in here output operands are declared first then we declare input options then we declare clobers. The next point is that operate operands are numbered sequentially across outputs and then inputs. So you can also see this really well in this example. So we are using these placeholders and uh these placeholders basically number the operands. So the first placeholder basically points to the first operand, the second one to second operand and as you can see these operands are numbered sequentially. So yeah. >> Yeah. So now let's talk about Yeah. Another point that if we don't have any input operants or output operins then we can just leave the fields empty. So like this we can just have nothing around here and yeah it works perfectly. Now let's talk about constraints and constraint modifier. So if you look at this PTX instruction then you can see that in the operands field we have this equal sign then we have this r and then we have the variable name which is declared in C++. So what are these? So the equal sign are our constraint modifiers. The R sign is the constraint and then we have the variable which we are working with. So constraint modifiers basically tell us the the nature of that operand. So basically they tell us if the operand is right only. They tell us if the operand is read and write and uh if the operand is an early clober which basically early clober helps us to perfectly define the reusable register. So it prevents us uh it prevents basically the hardware from uh reading the operand before it is modified. So yeah that's how it's works and then we have our constraints. So these and uh these like we have f in here and we have different types of constraints. So we have h which represents the 16 bit unsigned in integers. Then we have r which represents 32-bit unsigned integers. These are used a lot in addresses and you know representing unsigned integers. Then we have L which are basically 64-bit unsigned integers and 64-bit addresses. We have F which represents the floatingoint numbers. D which represents 64-bit floatingoint numbers and N basically represents the compile time constant which we discussed in here. So yeah these are basically evaluated at compile time and yeah so here are some more important points about the operands. So we represent operands in text order like uh these. So we use these placeholders to represents the op uh to represent the operand. So the first placeholder will represents the first variable and so on and uh we represent them using constraint modifiers and constraint like these and uh we use these for inputs but we can also have you know constraint modifiers for inputs too. So yeah and we use curly braces for local register scoping. You'll see them a lot. They are just pieces. So yeah, so now let's talk about the instruction string. What is instruction string? Why it is important? All that. So instruction string is basically the string which constrain which contains the actual Ptx instruction which we are going to use. So it is the PTX part of the instruction. It is this WGMMA async or it is the CP.ync sync or whatever PTX instruction you are using and uh instruction can be written in single line or multiple lines using new line. So we can have a single line PTX instruction like like this or we can have an instruction which spans for multiple lines like this. So yeah and uh yeah and we can use placeholders to represent the operants. We discussed them earlier and these are replaced by actual operants which are used in assembly. So yeah, sorry for that. Then we have output operants. they come after the instruction string and uh we use a special modifiers and constraint constraints to represent the output operants. So like this and uh the output operants are comma separated to represent multiple outputs. So if we have a lot of outputs we can just separate them using comma and yeah it will work like that. Also in the ptx instruction string we have comma separated operants to represent multiple operants. So yeah that's one of the things to keep in mind and yeah basically that's all about the output operants they basically represent the output of that you know the things which are modified with that ptx instruction or whatever. Then we have the input operants. These are represented after output operants. So after we have the instruction in string which is separated from output operants using a colon and then we separate the output operants from the input operands using another column. So you can either have the only read only inputs operands or we can have read and write operands. So yeah and uh we can use plus n constraint to get compile time operants. So these are evaluated at compile time. We discussed them a lot in class. So yeah and uh we can use multiple input operants by again separating them with commas and putting them after output oper. So yeah then we have clobers. So what are clubbers? Soers tell the compilers which resources are implicitly modified and we are not specifying them in our output oper but these are like you know internally modified. So this can be anything from you know uh a temporary register which is created inside the um instruction or whatever and uh this prevents the compiler from you know making wrong assumptions and modifying the code incorrectly or you know maybe corrupting the variables or whatever. So these are really important and yeah. Okay. So now let's talk about another really important concept in CUDA which is state spaces. So what are state spaces? State spaces basically refer to different memory regions which we have. So we can have you know a state space for our global memory. We have them for shared memory for clusters. We can have them for shared memory for CTAs and so on. We'll discuss all of these in detail. And uh these basically specify on which state space we are launching any instruction. So most of the time when we launch some PTX instruction it is uh we need to specify that okay we are working on this particular status spaces and status spaces help us specify that. Okay now let's talk about the first state space which is the register state space. These basically represent the register memory and uh these are one of the really important status space which get often gets overlooked because they are not a lot of time used in the main instructions but these are really important because they help us create temporary registers and uh we can use these temporary register in any way we want and yep if we create a lot of these temporary register um you know variables then uh we can just spill our register memory and uh that can really impact performance in the negative way. So take care of that. Then we have our global state space. This represents the global memory and uh we can use these for different operations. So here we have a PTX instruction for asynchronous bulk copies and uh as you can see that we have the global shared space. I specifically defined in here. Uh so yeah we use them a lot. Then we have the local status space. These are the per thread private memory. Each thread gets its own private local memory state space for storing data which does not fit in registers. So when you know register spilling happens, it generally goes down to local state spaces. These are rarely used. So yeah, then we have another rarely used state space. It's called param and uh it serves as kernel input parameters a lot of times and uh it is also used to pass arguments to basically it's basically used to pass arguments to kernels and functions. So yeah then we have one of the most important state spaces. These are the shared state space and um these comes with two subqualifier uh which is CTA and clusters and they are represented like these and uh shared status space is really important. They represent the shared memory and uh these can be accessed by the threads in same CTA or in the same cluster and we'll discuss about them a lot during this course. So yeah then we have our data types. So we have a few categories of data types. We have the unsigned integers which are represented by U8, U16, U32, U64. Then we have our signed data types. These are represented by S8, S16, S32, S64. Then we have our floating type floating point data types which are represented as F16, F32, F64. so on. Then we have the raw bit patterns. So they don't represent any specific data types. These are just bit patterns which are represented using these formats B8, B16, B32, B64, B128. And then we have the predicate type. These are represented by dot red and they store a boolean. So yeah then we have some really important data types in deep learning. These are one of the most important data types which we are using a lot in modern deep learning. These include BF-16 which is used a lot more like a lot more than all the traditional data types and uh these are bas these basically represent the 8 bit exponent 7 bit mantisa in the floating point you know floatingoint world. Then we have E4 M3 and E5 M2. uh these are basically our 8 bit data types and uh E4 M3 represents four bit exponent three bit mantisa and E5M2 basically represents 5 bit exponent and two bit mantisa so yeah and then we have TF32 which are stored as B32 these are 32-bit format with same range as F-32s but reduce precision so these have the same precision as of BF16 um I'm sorry as of U F16 I think and uh they have a um same range as F32 so yeah and then we have four bit floats E2 M1 and yeah and we can also have packed data type so these are basically multiple data types which are packed into a single value And we can use them for you know parallel operations and memory transfers and all that. So we have U16 X2 and S16S2. These basically hold the two 16 bit integers in a 32-bit register. We have F16 X2 which basically holds the two F16s. We have BF16 F2 X2 which basically holds two BF16. We have E4 M3 X2 which holds two E4 M3s and E5 M2 X2 which folds two E5 M2 4bit floating point f. Okay. So now let's talk about pointers and state spaces. These are another one of the most important concepts in CUDA. So basically in our whole programming space we have a single unified virtual addressing system which maps each state space to a non-over overlapping region. So basically when we are working with reserved state space it's working in an address which cannot be overlapped with the state space for global memory or for shared memory or whatever. So we have specified addresses for all of these overlapping region like this. So we have a section which is reserved for reserved like for system operations or whatever. We have separate ones for global memory and local memory and shared memory and so on. So we have different pointers for different state spaces and the first thing which comes along are the generic pointers. So when we create any pointer we say that all pointers are created equal and what that basically means is that all pointers are generic by default. They are basically 64-bit addresses in a single unified virtual address space and they span across all the state spaces. So there are there's no runtime checking to determine what space we state space we are in with generic pointers. And we can do all the normal pointer arithmetic like we do pointer arithmetic in CPU with generic pointers. So yeah and uh when we dreference a generate pointer the hardware checks the right the high bits to find out which you know which state space we are working with and then you know it takes that to that spa state space to you know work with them. So these are generic pointers and that's how they work. Then we have our state space pointers. So when you write inline ptx for state space pointers or when compiler generates them we get these state space pointers and uh they these are the pointers which can be qualified with their state spaces. So these are not like actual meaningful pointers. These are just raw bits which are meaningful to the specific hardware instruction. they can they cannot be dreferenced in C++. So if we try to dreference any state space instruction which we'll talk about how we can create a state space pointer but if we try to dreference it then uh it will basically crash or it will give us some kind of carvage or yeah it will be meaningless and there is no runtime space detection. So we if we have a state space pointer then hardware already assumes that it is an address in shared memory and u basically this enables specialized instru instruction which requires spe explicit space qualification. So we have a lot of instruction which require us to you know make sure that the pointer which we are using are defined in that particular state spaces and we'll discuss these with in a lot of detail when we are actually you know writing code for you know different operations and whatever. So now let's talk about how to convert any generic address to a state space address. So to convert any address we use cvta and uh it's basically a ptx instruction which converts pointers between generic addresses and specific memory space addresses and uh we can also use the C++ version of these which is basically CVTA generic to global CVTA generic to shared and for local and con constant shared space too. So we basically have these to convert any generic address into a specific state space address which we can later use and uh we need them in tensor core instructions. So when we are creating a tensor uh when we are creating a WGMME descriptor which is basically a descriptor for tensor code instructions, we need to specify um is uh basically specify a pointer in shared memory which contains the you know source address for our tiles and uh we need them for TMA. We need them for uh thread block cluster. So we have whenever we are using thread block cluster for any specific reason we need to have um a specify address specified address for that CTA inside that thread block cluster so that we can access that CTA and access that CTA's distributed shared memory and we need them in MIG multi-instance GPU configuration so that we can you know proper we can have proper memory isolation between GPU instances. So these are really important. Okay. So now let's talk about the their use in thread block clusters. So basically how it works is that if I actually want to access some data in some specific thread uh block in a cluster then I cannot just use the address inside that cluster. So I cannot use the shared memory address inside that cluster because remember when we are using a specific shared memory address it is only meaningful for the threads inside that inside that thread block. So whenever we are using shared memory address in a block it's meaningful for the threads inside that block. So if we are trying to access any other block and we are using it for in any other thread block in any other thread in different thread blocks then we cannot access the shared memory in another thread block. So if we have two thread blocks and we try to access each other's shared memory using you know the shared memory address then it will basically be meaningless. So the way memory addresses work in a cluster is that we have this specific ptx instruction which gives us the memory address which is specified to that specified to that thread block in the cluster. So basically we can get we get an address which can be used by any thread block inside that cluster to access that part of shared memory of that particular thread block. So basically what we have is the rank of that thread block inside the cluster and we have the you know we have the source address. So we have the shared memory address which we need to you know transform to an address which is which can be used by other thread blocks in the cluster and then we have the destination which is basically the final output and uh this instruction basically converts a shared memory address from the current CTS shared memory space to the space in any other CTS shared memory within the same cluster and it takes a rank not a block ID. So what does that mean? That means that it does not takes in here the block ID inside the grid. So we are not talking about the block idx. We are talking about the rank inside the cluster. We are talking about this. we are talking about the linearized rank of the CTA in the cluster. So that's one of the things which we need to keep in mind because a lot of time there's this confusion between the position of the block in the grid and the position of block in the cluster. So yeah while using map we need to take care of this and with that our lesson ends. Thanks a lot. So now let's talk about asynchronicity and barriers. And it's really important to learn these things really well because these make the fundamentals of two of the most important operations which we are going to perform on our GPU and those operations are bulk asynchronous copy using cp.async.bulk bulk and WGMMA which will be doing computation on tensor course. So let's begin. So any system has two major operations. First is doing operation and second is fetching operation. So if you are processing some kind of information then the fetching part is retrieving the data and maybe pre-processing the data and the doing part is doing computation on that data. If you are solving an equation then the fetching part might be reading the equation or getting the equation and the doing part might be solving the equation using your reasoning and you know coming to a solution and if you're cooking some kind of food then fetching part might be bringing the ingredients and preparing the ingredients and the doing part would be actual work of cooking the food. So in a synchronous world these happen sequentially which means that you are blocked until your current task finishes. So you can do one thing at a time and you are blocked from performing two task at once. And the synchronous way of cooking food can be thought of as one where where you are cooking food and you are not doing anything other than cooking food. So if someone calls you or if you have some other work to do, you cannot do this because you cannot you are blocked from doing anything else. So how does it work in asynchronous operations? What is asynchronicity? So you can think of asynchronicity in the context of cooking food as you are cooking food and while the food is being cooked in whatever utensil you can just go out and work on something else and after some time you might get some kind of a signal uh or you can see some kind of color in your food and then you can come back and uh you can use the ready food. So basically you got some kind of signal which um told you that the operation is completed and now you can come back and use the results and then you completed that task. So you accomplished two task in the time it took to do one and you achieved something called latency hiding. So if the cost of fetching a resource is higher than the time it takes to process it then you must overlap the fetching of the next item with the processing of current item. And this is basically latency hiding overlapping the fetching of the next item with processing of current item. And the key element in latency hiding which enables this whole behavior is not being blocked by an instruction and the ability to put the current job in the background and work on something else. So in the case of cooking food, you were able to put the current work which was cooking food in the background and you could do work on something else. And once the food was cooked, you got some kind of signal and then you came out and you used the ready results. And uh this blocking and non-blocking nature is something which is actually a technical term. So a blocking or a synchronous operation is an operation which pauses the execution of the thread which launch the instruction for the operation. So until the operation is fully completed, your launching thread cannot launch another instruction and uh it just sits there ideally and uh this basically creates an implicit synchronization point where you are guaranteed that the instruction launched by the thread is completely finished before launching another instruction with the same thread. But a non-blocking instruction or a asynchronous operation on the other hand returns the control immediately to the launching thread. So once the thread launches the instruction then the control is uh returned back to the thread to launch other instructions in its cues before the operation has actually finished. So so this is what enables latency hiding. So let's talk about where we need latency hiding and stuff in GPUs. So in H100 and in most of the modern GPUs doing math or compute is incredibly fast than fetching the data. And if you remember when we discussed about the architecture of H100, we saw that moving data from HBM to shared memory takes a lot of time. And one of the reason for that is just the physical physical reality of moving data from a far away point in HBM to shared memory. Not to mention the reduced bandwidth as compared to the bandwidth which we get in shared memory. And because the tensor cores or the other compute cores are really close to shared memory from which we generally perform computation on the data the operation of computing is incredibly fast than operation of fetching. So this basically means that the goal of your high performance architecture is not just to make the math faster but it is to ensure that math never stops so that you can work on maximum data in the same time. And this is one of the best things about H100 which is that it is not only it not only gives you really high computation speed but also it gives you the power of asynchronicity where you can overlap a lot of compute operations with moving data and uh you can make sure that the computation never stops and this is powered by using non-blocking instructions for tensor coursees and TMA. So now every async in operation which you will launch and which you will perform will follow a kind of pattern which you see here. So the stage one is initialization where a thread will fire an asynchronous operation and then it will immediately go off to fire the next instruction in it in its queue. Then the second stage is tracking and parallel execution. So there will be some part of the system the hardware which will be basically responsible for tracking the asynchronous operation in the background. This is really important because if you remember the cooking example while we were preparing the food we need some way to realize that the food is cooked otherwise what will happen is that the food will overcook by the time we get there or we will just forget and the food will just burn. So there needs to be some way that the system tracks the operation so that once the operation is complete there is some kind of a signal which is given to the hardware and u there's some form of synchronization which can be done and while this operation is working in the background the other units are working in the parallel because the same thread fired other instructions for this hardware. And in the end there is synchronization where once the operation is done operating in the background there is some way the hardware tells the thread that okay now this operation is done and you are good to go and use the result of this operation. Now one of the things I want you guys to truly understand is that while we are doing computation on buffer zero, we are still loading data in buffer one. And this will be really important because in the future we will be doing computation on buffer one while the data will be stored on buffer zero. And the reason for all of this, the reason for having a warp specialization, the reason for making the computation and the uh data loading operations non-blocking and asynchronous is to make sure that we are spending most of our times in doing computation and we are keeping the tensor coursees busy at all costs. And this is what we are going to do during this whole operation. So in the next step the buffer zero computation is done and the data is ready for buffer one. And once the computation on buffer zero is done then we will start loading the result for buffer zero and we will start doing computation on buffer one. And uh since the data is computed really quickly, the results will be we will have the results of buffer one before the the results of buffer zero are even completed storing. And uh we will start doing operations for storing the data of buffer one. And uh once the data for buffer zero is fully stored then we will clear the buffer and we will start reloading the data into buffer zero. And same thing for buffer one that once the data is stored for buffer one then we will start reloading the data into buffer one and the computation will start again. And this is how this whole pipelining thing works. Now here for a moment uh the tensor cores were idle and this is not what we want if we are optimizing for maximum performance. So basically then we put a different efforts and we put different things and we try different things in order to make sure that this whole line is green to make sure that we are utilizing our tensor cores a lot. Okay. Now before we go ahead and learn more about the different ambarriers and the different synchronization primitives, I want to discuss a bit more about something called proxies. And uh proxies is very closely related to asynchronicity. And uh let's try to understand what are proxies. So in the context of Nvidia H100 proxies are basically something which are used to distinguish between who is performing the memory operations and uh this becomes really important in the context of H100 because with H100 we have this whole new concept of asynchronous units. they were before but with H100 we have these two new asynchronous units that are your TMA and your tensor cores and uh these units perform operations in a very different way than other units which we have and that's why we introduced the whole concept of H of proxies and all that. So if there are two operations which are happening at the same proxy then the hardware would guarantee that these operations are safe and ordered. So this is the general way we do everything you know all your general computation the hardware would almost every time guarantee that these operations are safe and ordered unless the compiler is doing some kind of fancy optimization for latency hiding. But if an operation A is in proxy 1 and operation B is in proxy 2, then the hardware would stop checking. The hardware would assume that these operations are completely unrelated and it would just let them run wild out of order and there would be no communication or no synchronization between these two operations and this would all happen in parallel. And now you can probably understand a bit more about why this concept of fence is important and how it would actually link to these asynchronous units. So in the context of H100 again we have two major types of proxies. You have generic proxies and you have asynchronous proxies. And when you are writing a general CUDA kernel, your thread executes all the instructions sequentially. And when a thread would read or write memory, it would act as a general generic proxy. And uh the whole operation would happen as generic proxy and the hardware would guarantee that these operations happen in the order in which you wrote them again for most of the cases and within that specific thread. But when you talk about asynchronous proxies then we are talking about the hardware mechanism which are performing bulk copies or which are performing the tensor core math and when you issue an instruction like cp.async.bulk or if you issue an instruction like wgmma and I don't want you guys to get confused about these. So cp.async.bulk bulk here is basically the instruction which we use for our asynchronous bulk copies. Those asynchronous copies which are uh which are done using the TMA and WGMMA are basically your really fast matrix multiply accumulate operations which are done on the tensor course. So when you issue instructions like these um in your CUDA kernel then the generic proxy generally kicks off these instruction and it launches these instructions and then it immediately moves to the next line of code and then these instructions which are performed by your asynchronous proxy they run completely independently of generic proxy. They don't know what the thread is currently doing and the thread does not know that you know your proxy is finished or whatever is happening with your other proxy and this is why you need a fence and u we learn a bit more about fence and you know how we synchronize these two proxies but before learning that let's see what are the different ways in which this dynamic where you know the generic proxies never know what asynchronous proxy is doing and The asynchronous proxy never knows what generic proxy is doing can go really wrong. In the first case of it is your read after write hazard or raw hazard. And what is happening is that because generic proxy and asynchronous proxies have different paths to memory, that's how these operations are are isolated because they have really different ways of execution and they have really different paths to the memory. The generic proxy generally operates through the SM's L1 cache. Whereas the asynchronous proxy bypasses the L1 cache and it interacts directly with your L2 cache or HBM and that's how asynchronous proxies get a lot of their their get a lot of their speedups because they bypass the L1 cache and whenever you are doing some kind of stores using generic proxies then these stores are initially held into your local store buffers or your L1 data cache L1 instruction cache and uh these operations are not immediately visible to the rest of the system. So these are not visible to the rest of the different blocks and these are not visible to the TMA or the tensor coursees too. And uh if generic proxies are writing something to the memory address and uh once they are done writing it then they immediately trigger the asynchronous proxy like the TMA to read the same address then asynchronous proxy would likely read some kind of stale data from L2 cache or DM because the new data which was supposed to be written to L2 cache or DM is still struck in generic proxies buffer or generic proxy L1 data cache and because it is still stuck there and the generic proxies have not yet written the whole thing into L2 cache or S or HBM what would happen is called a read after write hazard where even before the write actually executed by the generic proxy you write the whole value. So this is one way in which things can go wrong. The other way that things can go wrong is a right after read uh hazard. And uh basically the way this works is that a thread or a CTA would you know would do a generic proxy read of address X and the way you read it is that you pull the data from X into the L1 cache so that you can you know you can go back and forth and you read it again or something and uh once it has done that then what would happen is that asynchronous proxy would update this value X via L2 cache or HBM. and it would write this this updated value into L2 cache or HBM or you know or your shared memory and uh once it does that then you know it won't be able to edit the view in the generic proxies L1 cache because the asynchronous proxies are generally not interacting with L1 cache. So what would happen is that generic proxy whenever needs that data you know if you have some case where you first read some data and then um you know some kind of asynchronous proxy rewrites it and then the generic proxies reads that data again then instead of you know again going to the L2 cache HBM or shared memory what would happen is that the generic proxy would just go to the go to its own buffer of L1 cache because it does not know that the data is overwritten and what it would do is that it would just read that old data from its own L1 cache and it would assume that this is the data which is is still there in the in the L2 cache or in the HBM or in the shared memory. And what would end up happening is that these generic proxies would keep returning the old value even though X was updated via your asynchronous proxy. And this is called a write after read hazard because you already read some value from the L2 cache or you read some value from the shared memory and you moved it into your L1 cache and now you are just reading it from them and you know you you wrote something and even though you wrote something in the L1 cache or in the L2 cache or in the HBM your your generic proxy is still reading the old value from the from the L1 one data cache and this is called a read after write hazard or war hazard. So the way we prevent these read after write or write after read hazards is by using something called fence and as the name suggest fence is an explicit ordering or visibility point. So it is basically something which we use to make sure that operations are ordered and they are not reordered by the compiler to hide some kind of latency. And uh it constrains how memory effects become observable. And it does that by making sure that you know every single operation till now is actually made can actually make its effects and is you know is fully completed. And uh we use it when the hardware might otherwise let things pass each other especially with asynchronous operations which are not automatically ordered with normal load stores in the way you might assume. So there are two main types of fence. There are there is ordinary fence and then there is crossroxy fence. And those are the two main you know two main times when you actually want to make sure that your your your operations are ordered and they are visible right because what would happen is that your operations which are done by cross which are done by an asynchronous proxy might not be visible to a synchronous proxy or what might happen is that you know there are two operations which are happening in synchronous proxy but one of the operations you know it modifies memory in a way that is not explicit and it's not obvious and then you know the other operation uses that part of memory and then it just the compiler reorders it and everything gets messed up. So in order to prevent that we use fences to make sure that every single um operation is you know is reordered is not reordered and and every single operation behind that fence is perfectly executed before we get ahead and you can have same thing for every single operation ahead of the fence. So that's what fences are for and Nvidia explicitly calls out that a crossroxy fence is required to synchronize across proxies for proper ordering and we learn about crossroxy fence and basically the thing is that as we learned that you know generic proxy and cross proxy and asynchronous proxies are totally independent they don't know what is happening in either you know in the opposite proxy. So in order to just make sure that whatever changes to memory they did we need to use a crossroxy fence and what is cross cross proxy fence we'll learn about it and uh the next point is that all the fences are escoped so you can enforce a fence for the scope of your CTA you can do it for a cluster you can do it for your whole GPU or you can do it for the whole system where it also enforces it for the for the CPU and stuff and the scope determines who must see the ordering and it is tied to a point of coherency in the hierarchy. So basically the thing is that um if you are doing some kind of write in the global memory then the thing is that you know even if every thread in the block knows that this right has been made other blocks don't know it and maybe other blocks might end up using you know some some other outdated data. So in order to prevent that we need to set a scope and then that scope is based on this point of coherency. this point where where you know where it is visible to every single unit or every single agent which we want it to be visible to. So we set it we set this in a way that you know it is visible to everyone which we want it to be visible to and that's how we use these fences and there are two types of fences. There is ordinary fence. So these are the fences which are enforced within a single proxy. And then there is crossroxy fence which is established between two proxies. So in ordinary fence we have two types of fence. We have release fence and we have acquired fence. And both are part of ordinary fence and we have crossroxy fence which is the crossroy fence. And uh in the case of release fence, it enforces this rule where every single memory operations especially writes which appear before the release in the program become visible before anything that happens after release to other threads which synchronize with it. So basically the thing is that if you have an instruction which actually uh which which has this instruction fence.release.cta CTA then it would enforce that every single line of code which came before it must be executed before this instruction and uh you can have it for you know you can have it for your own CTA you can have it for cluster and you know it its basic job is to make sure that everything happens and this is specifically important for rightes because you know later you might use that data so it makes sure that every single right which which happens must appear before release in the program order and that's how it the whole thing works. The next type of fence which we have is acquire fence and acquire fence basically enforces the opposite oneway rule where no memory operation which appears after the acquire instruction in the program order and this especially include reads is allowed to be observed happening before it. So if there is any instruction which is you know which is being launched after this instruction then it cannot be reordered to be launched before it and this is done to make sure that every single right which have been made till now can be read with the current data. So if you have written some kind of data here and then you are going going to read it then this instruction make sure that you are safe to read it. This reading instruction, this operation which was which would read this data cannot just be reordered to happen before the right operation. So that's what acquire fence is for and after the acquire thread is guaranteed to observe the rights which were made visible by a matching release and this matching release must be uh must be made within the chosen scope or proxy. So that is obvious right? So these are the two main types of fences which we have in in in the ordinary fences and um we'll see how these fences are are launched with the M barriers and how these are integrated with barriers later in this lesson. And finally we have our crossroxy lens or crossroxy fence. And um if you access the same location across multiple proxies then you need a crossroxy fence. And for async proxy we use fence.proxy.async to synchronize the memory between the generic and the async proxy. And this is one of the most important instructions which we have because it makes sure that all the changes which are made by your async proxy it can be your uh it can be your tensor course it can be your your TMA for most of the cases it's TMA because the output of tensor course get gets written to registers but we don't have to talk about that but the job of fence.pro proxy.async Async is to synchronize memory between the generic and async proxy and uh it drains and orders the generic proxy and uh yeah let's get back to our presentation and let's now review the whole MARE pipeline. This is the pipeline which is basically powering the whole operations which we saw there and uh the reason for having mriar and the reason for having this whole complicated pipeline is to make sure that the tensor cores don't use the results before the loading is finished. So the way the whole thing works is that first of all we initialize a barrier object in shared memory. So in this object all the information about you know the barrier and the operation will be stored and what we do is that we use marrier.init to create a barrier with something called an expected thread arrival count. And thread arrival count basically makes sure that a specific number of threads have issued the copy instructions. Now this might seem a little vague to you and you might be thinking how we will discuss this in a lot more detail. So in the next step the what we do is that we set the expected transaction count for the async operations which we are going to launch. Now as some of you have might guessed this expected transaction count is the size of the whole data which we are going to transfer in this whole async copy operation and we do this using marrier. Ptx instruction and every time you launch an marrier. Ptx instruction you are reducing the thread arrival count. So now you guys understand what we are doing when we are setting up an expected thread arrival count. What we are doing is that we are basically telling the barrier that these many times I'm going to launch a copy instruction and whenever you are launching the copy instruction you have to set a expected transaction count and whenever you do that your expected thread arrival count goes down. This is basically just to make sure that you are launching the number of copy operations which is predetermined so that once the whole copy operation for the whole asynchronous pipeline for all the computation and all the data loading is done then the expected thread arrival count is zero and u the expected transaction count is also zero. So then we launch an asynchronous operation and u we attach the barrier with this asynchronous operation so that while this operation is being progressed the expected transaction count decreases. So once the whole operation is ended because we launched the asynchronous operation using marrier. the expected uh thread arrival count will be zero. And because we launched the asynchronous operation with the barrier attached, the hardware will automatically decrease the expected transaction count. And uh in the end, the thread arrival count and the expected transaction count will be zero. And once that happen the fl the phase of the barrier will be flipped and the arrival count will be reseted and then we do we'll do the next operation. Now after this you might have a lot of questions. You might have questions regarding thread arrival count about what is it? You might have question about expected transaction count. And you might wonder what is phase and uh what is the meaning of flipping the phase. So let's discuss these in a bit more detail. So a barrier's phase is basically its current reusable state or cycle. So it is a single bit in the barrier object. uh and uh basically this single bit flips whenever the barrier completes a cycle. Then what is transaction count? Transaction count is basically a count which represents the size of your asynchronous operation. And expected transaction count is the amount of work which is yet to be done which is basically the amount of data which is yet to be moved. And because these are asynchronous operation, the hardware automatically decrements the transaction count as the async operation progresses as we discussed. And u we attach the barrier to the async operation to make sure of this behavior. We will look into that in a lot more detail when we will be discussing the whole asynchronous copy operation. And we can reuse a barrier by flipping its face at the end of the operation. And this is one of the most important parts. And uh basically what we will do is that we will flip the s flip this single bit. And uh that will basically uh reset the transaction count. And uh basically the thread arrival count will uh will be all automatically reseted to the value which was set during the initialization. Okay. Now that was a lot to comprehend and uh frankly it's fine. We'll discuss these concepts in a lot more detail in the coming lessons. So it's okay to be a little bit uh tense but uh let's now talk about the first step. Let's now talk about initializing the barrier. And uh this is one of the instruction where a lot of bugs originate from. And if you get this instruction wrong then nothing else matter. So I want your attention right at your screens now. So let's talk about this PTX instruction which we use in initializing this barrier. So basically in this PTX instruction we are setting up the space in which the barrier is being created. Try to make sure that this space is always shared memory because it's really high latency operation to create an object in global memory and then to have your consumer W groups try to reach out to it to uh track the progression of the operation and uh to basically uh do something called a weight operation on this barrier. So it's better to have this to be set to shared and then B64 is basically the size of your barrier object. So it actually exists in shared memory and it's B64 basically means that it's a 64bit object. Then we have here the address of the barrier object and this is the expected thread arrival count. So yeah, this is the value to which the barrier will reset when you reuse the barrier by flipping the phase. And uh the second thing which I want to do is that I want to take a look at one of these uh very important and very popular uh repository for uh CUDA which is Thunder Kittens. And uh what we will do here is that uh we will try to look at some of the source code and we will see how these primitives are implemented. So in this uh repository we can see that there is a function named init sema 4. This is basically the function for initializing the barrier. And as you can see there is the pointer to the barrier object and we have the thread count and transaction count. So it's generally set to zero. We can set the transaction count with the barrier object. And uh basically what we are doing is that we are firstly count uh casting the barrier pointer into the pointer in shared memory and then we are launching the instruction using our inline PTX which we discussed earlier and uh this is how it is launched. Okay. So let's get back to it and uh now let's talk about a few things which you need to keep in mind. So generally your scope is again shared CTA so that the threads can uh threads in the same block can see the barrier. This is really important because the consumer warp groups which are trying to do computation will need to access this barrier in order to basically try to initiate some kind of computation on this barrier. And uh the second thing is that the address and shared memory pointer is creating using CVTA which can which we can also see in this repository. We are using CVTA generic to shared to basically convert the generic pointer in uh into the share the pointer in the shared memory and uh we can use shared cluster to make the barrier visible to all threads in a cluster. This has its own use case which we will discuss further in the lecture. And u you must use map a ptx instruction. I'm really sorry this is map a so yeah we must use the map a ptx instruction to get the address in cluster. So when we are launching the uh the barrier in cluster, we will then need to make sure that instead of this that this pointer is the pointer in uh in the cluster and we will use the map aptx instruction which we discussed earlier in order to convert the pointer into the pointer in the cluster and uh mriar objects must be 64-bit align. This is a strict requirement because unaligned accesses might cause silent corruption. So keep this in mind. This is one of the really important points which everyone should keep in mind. So yeah. Okay. So now let's talk about setting the expected transaction count. So the expected transaction count is set using this code this PTX instruction. So basically marrier. And uh every time you see marrier. You have to assume that your expected thread count will decrease by one as a default behavior. Then the next part of this ptx instruction is expect. This makes sure that what we are trying to do is that we are trying to decrease the expected transaction. uh we are trying to set the expected transaction count and uh then we have to specify the scope of our barrier object the uh the barrier objects type and then we have a placeholder which we'll discuss about and then we have our barrier pointer and then we have the expected transaction count which we are trying to set. So keep this in mind. What this does is that it decreases the expected arrival thread count by one and it adds the expected transaction byes with tx count and uh this is added into the barrier pending transaction count. So basically what happens is that if you are launching multiple uh copy instructions which we often do in uh in double buffering or in triple buffering then uh basically this transaction count is added to barrier spending transaction count. And this is one of the most important points to understand because what might end up happening is that you might end up using the same barrier for these two buffers. And what will end up happening is that once the loading in this buffer is completed, the uh the consumer W group will not start uh doing computation because the expected transaction count was the transaction count for this buffer as well as this buffer. So please make sure that you don't fall in that pitfall. And uh basically you are adding this value into the previously set expected transaction count. If you are doing this for the first time like if you are just creating the barrier and then using this then the transaction count initially will be zero and this much amount of transaction count will be added. So if you call this thing twice with 4096 as a transaction count then your barrier will expect 8192 bytes total before it basically opens up. So now let's talk about the placeholder which we have for the instruction and basically what we are trying to do here is that we are discarding uh an operand an output operand in this whole instruction and that output operand is phase out and it basically represents the current phase and transaction count of the barrier and by current phase I mean that it uh it is the phase and transaction count before this arrive instruction. So it encodes that in a 64-bit encoded token and uh the reason why we are discarding it is because we don't need it. Now the reason why it exists is because of some complicated synchronization issues related to you know thread divergence or something but basically for most of the cases we don't need this encoded token for any kind of synchronization. So what we do is that we discard this whole thing because you know it is expensive to write stuff into our registers and uh it is much more cheaper to just you know use a simple uh one bit integer for you know for doing most of the synchronization as we will discussed later in this in this lesson and how the whole synchronization process is done. But basically uh the way we synchronize it is a lot more cheaper than writing 64-bit encoded tokens everywhere and uh messing things up. So for most of the cases we don't use a place for we don't use a 64-bit encoded token for anything. So we don't put it in the expect transaction count instruction and uh that is the whole thing. So after that let's just take a look at the whole instruction and how it is implemented in the thunder kittens library. So basically we are in the thunder kittens TMA. file in include ops threads util/tmma. And um what we are trying to do here at first is that we are just launching this whole thing for uh for the thread block clusters. And uh we are using the placeholder here. Then u importantly here what we are trying to do is that we are uh putting in the memory address of the barrier. So if you are putting in the whole you know 64bit barrier object then we would probably use uh simple uh uh simple zero here. But what we are trying to do here is that we are putting in the memory address. So we are using these square brackets to represent that and the rest of the stuff is same. We are putting in the marrier address and the number of uh and the transaction count which we want to put in here. So yeah there is that and in the second case we are doing a lot more stuff. So what we are trying to do here is that we are doing it for all the uh all the thread blocks in a thread block cluster. So the way m barrier works with a thread block cluster is that first you need to uh basically create the barrier for every single block in the thread block cluster. Then you have to uh basically initialize the barrier and uh set the expected transaction count of amount of data which is going to you know which is going to arrive at the at the at the barrier. And uh to do that what we do is that we basically first of all do the same thing. We convert the pointer into a pointer in the shared memory state space. And uh then what we do is that we create a holder for holding the address in the distributed shared memory of every single thread block. And then what we do is that we launch the map a instruction to get the address in the distributed shared memory. And uh basically the way it works is that we pass in our holder, we pass in our memory address in the shared memory and we pass in the destination CTA. So this can be different for every single thread block in the thread block cluster. So basically we are able to launch the arrive instruction for every single thread block in the thread block cluster and that is the way we use it for most of our cases. And uh as you can see the instruction here is pretty straightforward. We are passing in the address in the thread block cluster. So we are passing in the distributed shared memory address here. We are passing in the expected bytes of you know this whole operation. And uh we are having the the phase out as placeholder here as well. So this is basically all about how um how embarrier.exect do.exe_tx operation works and uh the next thing which we are going to learn is the synchronization. So till now we have launched the asynchronous copy instruction. So once we get all this done, once we create the barrier, once we set the amount of data which we expect to be moved by this full operation, then we just launch the asynchronous bul copy instruction and we attach the barrier in the asynchronous copy instruction as the with the address. So there is that as well. But basically what we are doing is that once we set the transaction count once we initialize the barrier then we just launch the asynchronous copy instruction. But now the thing is that the consumer WS the Ws which are going to compute on the data which we are trying to move in this case from the the global memory to shared memory or in this case from the thread block to thread block cluster. What we have to do is that we have to make sure that all this data which we are trying to move is you know is computable and at some point we need to make sure that the consumer threads can access this data to compute it. So how does the consumer threads know when they can compute when the whole data is arrived and all that? Well, the way the consumer threads know that now it is the time to compute. Now all the data have arrived and uh now we have to compute is by using something called marrier.try weight parity instruction. And basically what it does is that it um it basically provides a true or false value if the operation is completed. So the the consumer threads query the barrier and ask if this whole operation is completed and um the barrier basically provides a true or a false thing. So if it is false then you know the barrier the consumer threads can ask the barrier again the same question and if it is true then they can just go ahead and compute on the data which you know which uh was just um moved by the tensor memory accelerator. So what is the way through which you know we synchronize? We synchronize using this instruction. So this instruction is marrier.try weight parity instruction. And what we do here is that we pass a phase parity bit into the instruction which basically represents the special execution phase which you must verify is fully completed. So basically as we discussed earlier that the phase can be either zero or one and it tells about the current progression of the operation and uh basically what we do is that we pass a phase parity which represents the phase which you have to verify is fully completed and the hardware compares your current input parity bit against the marrier's current internal parity state to determine if that phase is still active. active. So if that operation is still active then we would be sent back and it would basically return false because what we are trying to query here is if the operation is completed. If the operation which is represented by current phase is completed and um if that whole operation if the phase of the barrier is same as the phase of the the phase which we provided then that means that that whole operation is you know is still in in work and what happens is that we get basically basically what we get is uh a false as answer and uh we we have to keep waiting. That false thing means that we have to still keep waiting. And uh if the whole operation is completed and if the phase parity is different than the internal parity state of the barrier then it will return true and then we can basically work on the data. Then the consumer threads can work on the data. So this instruction is basically called by the consumer threads to query if the whole copy operation is done or not. And uh yeah, basically that's how this whole thing works. And uh yeah, a suction a successful return implies that the barrier has flipped meaning that your current parity now refers to the immediately preceding and therefore completed phase. So that's how the whole thing works. This is the whole instruction. So embarrier.try weight.parity. So this is the try weight instruction. We have another instruction as well which is called test weight. But with tryate what we are trying to do is that we are saying that first of all check if the operation is completed. If it is completed then you know the barrier will open and then we can move ahead. But if it is not completed and this is the important part if it is not completed then basically put the thread which launched this instruction into sleep for a certain number of nanconds. So the general way we synchronize is that we just keep asking this question every single time. So the threads keep launching this instruction all the time until that whole operation is completed and until you know we are um given the signal to go ahead and uh compute on this on the data which is moved. But the thing is that this eats a lot of cycle and uh this just you know overwhelms the hardware a lot. So what we do is that instead of uh instead of basically running this loop a lot of times what we can do is that we can make the the make the thread sleep for a certain number of times so that you know so that the hardware can relax a bit in you know in nice sense. So that's how this whole thing works and um we have the weight complete operand which is basically the boolean result. So it gives us one if the barrier has reached the target phase. It returns a zero if the barrier is not done and uh the thread has just woken up. So this is how the whole thing works. And uh yeah the phase parity is the integer which represents the phase generation which you are waiting for to complete. So if the phase parity is zero and the internal phase of the barrier is zero then that basically means that you are waiting for phase 0 to complete and currently the m the whole operation which is happening and which is basically obstructed by m barrier has also the phase zero. So the whole operation is not completed and once that internal phase of the barrier becomes one then that means that the operation is completed and now we can go ahead and then there is the test weight parity instruction which is basically um which does not have the sleep thing. So it does not have the suspend time hint uh thing in here. the suspend time hint operand in here. And uh this basically just tracks the barrier's progress using a phase parity. And u yeah that's how it basically works. And um if you know if the whole thing is done then it would uh it would give us one and if the barrier is is still busy then it would give us zero. and the address is the pointer to the to the marrier object and phase parity is the parity of the phase which you are waiting for. So this is all that is and the next thing which we have with this same instruction is marrier weight with acquire. So we also have a dotacquire qualifier here and uh it does basically everything which there is in marrier.try Try wait with a really critical distinction and that critical distinction is that if the current synchronization is fully finished then acquire semantic would create a strict memory fence where basically it would ensure that once this whole thing is done once this whole operation is done once we get one as the output in wait complete then what would happen is that basically ally we would basically what would happen is that we would enforce a acquire semantic and acquire fence and uh this fence would obviously be a crossroxy fence where basically the results which are written by the TMA which as we discussed is an asynchronous fence which is an asynchronous proxy. So all the results which are written by the asynchronous proxy that is the TMM are visible to the tensor cores which are a synchronous proxy or you know any other threads which are also synchronous proxy. So if you are loading the data from shared memory into the registers to feed a tensor core then the acquire basically ensures that every single LD instruction you know if you are doing it manually then it ensures that every single LD instruction does not fire until the data is valid and until you know every single thread knows that this new value has arrived and uh every single thread can read the new value and if they fire early then your registers can get garbage from L1 cache. So this acquire is important thing to add in the marrier weight and uh even though WGMMA for most of the cases reads directly from the shared memory especially if we are using SS layout where we are loading a from shared memory and B from shared memory for you know the WGMMA operation the instruction itself requires descriptors and memory state to be consistent and dot acquire basically ensures that every single thing in the dependency chain is respected. Every single thread can view the new data, the data which has just been copied into the shared memory. So there is that and uh now what we are going to do is that we are going to learn about all the complicated ways to implement it. But before that, let's just take a look at this part of the PTX in the PTX documentation. And as you can see, all the instructions are here. And uh as it is written here that the test weight and triate operations test for the completion of the current or immediately preceding phase of the MAR object at the location is specified by ADDR. And we can see the square brackets in the ADR which represent their the location of the barrier and not the barrier object itself. And as it is shown here that test weight is a non-blocking instruction. So once you launch the test weight then it is just launched and you immediately get the control back of the thread and it tests for the completion of the phase as we discussed and try is a potentially blocking instruction and uh it also tests for the completion of phase. But if phase is not complete then the executing thread might be suspended and the suspended thread resumes its execution when the specify if a specified phase completes or before the specified phase completes following a system dependent time limit. So we can specify the time limit in here according to our will and uh the optional 32-bit signed integer operand specifies this whole time limit in nanoseconds that may be used for the time limit instead of the system dependent limit. So basically we have the option to wake up the thread when the barrier face flips or we can just use the our own custom time limit for this whole thing. And uh that is basically how this whole thing works. And uh yeah so there is all that you can read the PDX instructions if you want and uh you know know more about it. But now we'll go to the actual instructions which we have here. So okay so let's now try to understand how the trie instruction is implemented in thunder kittens and uh because this instruction is generally used by the consumer wops to basically get information about the barrier and to know if the whole asynchronous bulk copy instruction is completed. That is the reason why generally this whole thing is launched by the consumer wops and uh they do launch this using a loop. Now how do I know that this is a loop? Well, whenever we are working with loops, we are working with something called labels or bookmarks or things like that. So what you see here is this lab weight instruction and this lab weight basically u bookmark. You can think of lab weight as a bookmark and what it does is that it creates like a time stamp in here. So it says that okay we are working with this line and now this is uh a bookmark or a uh basically a label for which for the thing which we are working with and uh then what we do is that we launch the instruction and once the instruction is launched we get our you know our weight complete uh our weight complete predicate which tells us if you know the the phase which we passed through ph kph phase bit is passed or not. And uh if it is passed then what we do is that we move ahead and the consumer ws are now free to use the data. So they you know they work on the next instruction which they have and basically what we do is that we create another we create another basically bookmark or a label and they are branched to that label. So they skip the next line and they go directly to done and done is also a label and they just go here and after this because there is no instruction so they go and launch the other instructions which they have. But if this predicate says that the operation is not completed then this instruction is false. So basically what would happen is that this would be ignored the thread the thread would launch the next instruction which is branch to lab weight. So it basically says that branch to this instruction. So what it does is that instead of you know launching the other instructions which it has it just directly goes back to lab weight and it again you know it again launches all these instruction and repeats the whole process. So what you can understand from this whole thing is that this is sort of a while loop where we have a condition that P1 must be true and until P1 is true this just whole thing keeps looping keeps looping and keeps looping. Okay, but what about the rest of the thing? So just let me try to explain how the rest of it works. So what we are doing here is that we are creating the predicate and this is the predicate which would in the future contain the value true or false about you know the phase which is represented by kphase bit passing or not. So that's what P1 contains and uh then basically we launch the instruction and uh what we are passing here is the memory address of the barrier. So the memory address of the barrier object we are passing in our kphase bit and we are passing in P1 which would basically tell us if you know if kphase bit phase has passed. Notice that even though we are using try weight here we are not passing in the the sleep operand. So if we go back here then we have this suspend time hint operand which would basically put the thread to sleep until you know for some nanconds but we are not doing here. So basically if we don't have this then what would happen is that you know if we don't pass the the suspend time hint operand then what would happen is that the hardware would assume it to be zero and the hardware would just keep looping over the the whole thing. The next instruction which we have is the test fate instruction and uh what it basically does is that it does the same thing but it first of all is a non-blocking instruction. So it just immediately provides the value to P1 and uh it just follows the same logic. So if P1 is true, if the whole operation is completed, if the phase represented by Kphase bit is passed, then what we do is that we branch to done and we basically launch the next instruction. But if that is not the case, then what we do? So then what we do is that we put the thread to sleep for 5 seconds. So 5 nanconds. So what we are doing here is that we are implementing the same thing as suspend time hint but we are doing it manually. So we are manually putting the thread to sleep for 5 ncond. If this condition is not true and once the thread is you know awakened from sleep after 5 nconds then it basically goes back to lab weight and does the whole thing again. So that is what is happening in this instruction. Rest of the stuff is same. We are passing in the barrier pointer address and we are point passing in the kphase bit which represents the phase which we want to pass. And uh yeah that's basically how this whole thing works. And uh the next instruction which we have is a um an instruction which basically checks if the barrier is passed. So basically what we are doing here is that we are first of all putting in the same thing. So we are creating the predicate which would give us the value if the whole thing passes and if the kphase bit is if the phase represented by kphase bit is done and uh it passes in the the barrier pointer address and uh it also passes in the the kphase bit thing but once it gets p1 then it does something else. So it basically launches an instruction and what this instruction does is that it passes the result which tells if the barrier phase is over or not. So how this thing works? So what this instruction does, what this cell p instruction does is that it takes up a 32bit unsigned integer which we created uh which we basically pass in or which we create. So which we create here using success. So we created this success unsigned integer and we pass this into the instruction and uh what we say here is that look at P1 and uh if P1 says true then write one to the value which is represented by this. So it would write one into the success unsigned integer if P1 is true. But if P1 is false then write zero into this success operand which we have here. So that's how this whole thing works. Rest of the stuff is same there is no loop. So we just once this whole thing is done we just output success and it tells us if the whole thing is you know if the whole thing is true or not if the operation is completed or not and then we just cast it into a boolean. So that's what this whole thing represents and then we have like this huge you know scary a bit scary instruction but it is not that scary and let's see what would happen. So basically what we are trying to do here is that we are creating two registers of 64bit integers and uh they are there to uh to capture the value of start clock and current clock. How this will be used? Let's see. So basically then what we do is that we capture a variable which uh in and we put it into the start block. So we created the registers to hold the start clock and current clock and then what we do is that we inscribe the current clock cycle into start lock. So this is what we are doing through this whole instruction. We are taking the value of current clock uh and we are putting in it putting it in into the start lock variable which we have a start clock register which we created and uh how this will be used we'll talk about that but then we create a predicate for P clock how this will be used we'll learn about that as well and then we create the P1 for you know for outputting if the the phase is done or not and then we just launch our loop and uh we do the same thing. If P1 is true, then we branch to done. But and here is the difference. If P1 is not true, then after this whole instruction is launched, we measure the current clock cycle. So we measure the value of whatever is the current clock cycle. And what we do is that we subtract it from the start clock which we measured before we launched this whole thing. So before the loop was launched, we created a whole um a whole uh register for capturing the starting clock cycle. So the clock cycle before this whole loop was began. And uh if this loop basically fails a lot of time then every time it fails we would be recording the current clock cycle the clock cycle at the moment at which it failed. And then what we are doing is that we are subtracting the current clock cycle from the start clock cycle and we are storing this value into the current clock cycle. So we are doing current clock uh equals start current clock minus start clock. So that is the whole thing which we are doing here. We are subtracting two unsigned 64-bit integer and getting the current value of the clock cycle. And uh we are basically calculating the difference between the clock at which we started and the current clock which we have. And once we get that then what we do is that we have a statement. So the predicate which we created here, this was created to basically know if the if this whole difference if the difference between the current clock cycle at the and the clock cycle at which this whole loop is started is greater than or equal to this big number. So what is this number? Let's see. So 1 2 3 4 5 6. So basically what it does is that it checks if if the difference between the current clock cycle and the clock cycle at which we started is greater than or equal to this. And if it is greater than or equal to this then it sets this to true. If it is not then it sets it to false. And if it is greater than or equal to this whole thing then it is it then that whole thing means that these numbers of clock cycles have passed after you know after we went through this whole loop and what we do is that we crash the whole thing. So we crash the program if this whole value is greater than this big number and uh once we do that then and if that is not true. So if the program is crashed then everything is crashed. We don't move ahead. But if that is not the case, if peak clock is false then we go back and we launch this instruction again. We measure the current clock cycle and if now this whole thing is you know is true then we crash the program or else we just keep doing it. So that is what is done by this whole large scary function. See it is not scary at all if we take a look at it. And uh the second instruction which we have is the simple you know test weight instruction which does the same thing as as we saw. So yeah there is that and uh apart from that no there are okay. So then we let's see another one which is in group util sync.c. So there are some mrier instructions which are implemented in here as well. You see the initiation instruction which is shown here and uh this is to invalidate the the mrier. So basically to kill the barrier and uh what is this? So this is for arrival. We'll learn about this and uh okay so here are the other instructions which we have. So this is also implemented the same thing which we saw there and uh this one is also the same. So that one was try weight and this one is test weight and uh this one is again doing the same uh you know this one is also returning if the whole barrier state is true. And uh yeah same thing we are doing the same thing we are launching the same stuff. So yeah and uh then we have the TMA cluster and uh this also have the same instruction but this is for a thread block cluster and the difference is that this also uses the acquire semantics which we learned about in here. So it also uses the acquire semantics where basically we are creating the strict memory fence which makes sure that every single data which is written is now visible. So yeah the same thing is happening here. We are creating the predicate. We are starting the loop. We are launching this instruction which basically uh does the same thing for a whole cluster and uh it does the same thing for you know for the distribute the barrier which are which is created in the distributed shared memory and uh we are using the pointer which is a pointer in the distributed shared memory for the barrier and uh we are using the acquire semantics which basically make sure that every single write which is done by the operation would be visible and uh it is a crossroxy uh fence and uh rest of the stuff is same. So if the predicate is true then we go to done and we you know keep working with our stuff and if it is false then we basically uh continue looping over lab weight. So there is that and uh yeah and uh what we are doing here. Uhhuh. Okay. So here also we are doing the same you know thing where we return the the current phase and uh if our whole thing is done or not. And uh here we are also using the acquire u semantic. So there is that and we are doing it for a thread block cluster. So there is that as well and we are using memory clobbers and uh yeah the rest of the stuff is pretty much similar. So this is how we implement the same thing for thread block clusters. And uh now let's go to our first file. And uh here as well we are doing the same thing for our thread block clusters. We are using acquire semantics and uh yeah the rest of the stuff is same. We create the predicate. We start the loop. We see if the current phase which is you know which is represented by kphase bit is over. If it is over the predicate returns true we branch to done and we keep launching other stuff in the consumer threads. If it is false then we go back and we do the same thing again. And uh yeah basically that's it for this whole instruction. Okay. So now I want to make a really important point and the point is that in all the instructions which we saw which are launched by the producer VO group whether it is initializing the barrier or setting up the transaction count or actually launching the asynchronous copy instruction itself. All of these operations are launched by a single thread. So a single thread which is also called the leader thread. All other threads in the producer work group are called the followers. And the leader thread is the thread which is launching all these instructions. And this is a really important point because in the consumer groups it's not like that. In the consumer group every single thread is launching the weight operation. Every single thread is launching this operation because every single thread wants is waiting for the TMA copy to complete so that they can work on data. So with that let's talk about how we reuse M barriers. Now the first question which come to mind is why do we have to reuse a barrier? And the reason for that is that because whenever we are creating a barrier, we have to spend a lot of time you know creating a block in shared memory and then putting in all the information about the barrier. It takes a lot of time and whenever we are working with things like loading tiles or loading any kind of you know bulk data from global memory to shared memory and then doing a lot of computation on that the act of creating a barrier takes a lot of time and it eats a lot of time and this is the reason why we basically reuse a barrier instead of creating it because creating a barrier is a high latency operation and and we cannot afford to do that every time we are working with a tile. So this is how we actually reuse a barrier. So we set a phase bit before you know getting into everything. We set an expected phase bit before starting the whole operation. Then we create a variable. Then we launch a loop for iterating over all the tiles. And uh we set the expected transaction count. We launch the async operation. And then comes the part where we use this set phase bit. So we wait for the completion and the completion happens. And now we are here. Now the phase bit of the barrier which we created is same as the phase bit we set here. So now instead of creating another barrier what we do is that we flip the phase. So now what will happen is that the the phase of the barrier is the phase which was set earlier but now we have flipped the phase and now we can use the barrier again. And because all the expected transaction count is zero, we can use it again. But then the question arises what about thread arrival count? So Nvidia barriers have the M barriers have this really cool feature where whenever you flip the phase, what happens is that the expected thread arrival count resets into the value which you set when you were actually setting up the initial thread count. So this is how it works. So basically you get you reuse the whole barrier and you get the same estate as when you actually created the barrier for the first time. So this is really cool feature of Marrier and of Hopper architecture itself and this is the way you basically reuse a barrier. Okay. So till now we have discussed about M barriers and asynchronous bulk copies and I know what you are you might be thinking by now that we only talked about asynchronous copy operations and it really makes sense to use M barriers there and there are a lot of things like you know setting up an expected transaction count and then the internal hardware automatically decreasing that expected transaction count while the operation progresses and the question comes to mind that is MB barrier only for asynchronous bulk copies and the answer is no asynchronous bulk copies are a major use of embarriers but it's not the only place where we are using embarriers so I want you to consider this situation So we use embarriers, we use the producers to write data to shared memory and then what happens is that the consumer starts reading data from shared memory using WGMMA. And uh what happens is that the consu the producer now wants to override the the same part of shared memory with the next style. But because we are reading the data for from the shared memory for WGMMA, we cannot be sure that you know we have already read all the tiles from the shared memory so that the producer can write it. So if producer write overwrites the shared memory before you know the WGMMA operation is has completed reading it then what will happen is that your math will be garbage and u the operation won't be actually the operation won't reap the correct outputs and this is basically the right after read or war hazard and uh It's really bad and we always want to prevent war. So the question comes how do we prevent wars. So the simplest way to prevent war is to create a barrier with an expected thread arrival count. But this time we are not creating a barrier for consumers. We are creating a barrier for producers. So while consumer is busy reading all the data, producers are spinning on the barriers. So producers will try to you know try to override the data in the shared memory but they will be confronted by the barriers and they will start spinning on the barrier and once the consumer finish finishes its whole copy then the consumer will execute the instruction to decrease the thread arrival count and the barrier state flips and the producer can now write the data. So this is how it works. And now the question is that what is the instruction to decrease the arrival count? Do we just you know do we just launch this instruction with transaction count zero or something? And the answer is no. We have a separate instruction for this whole thing and that instruction is this. So this is the instruction simple marrier. scope this is your state space in which the barrier is and then the count and this count is the value by which we are decreasing the barrier thread arrival count. So expected arrivals how much we have to decrease the expected arrivals. This is the number and this is the instruction which the consumer B groups will launch after they are finished reading the data for WGMMA and this will just decrease the pending arrival thread count with this number. Okay. So now let's see how this is implemented in thunder kittens. So I found a single implementation of this instruction in thunder kittens and uh here it is. So it is basically implemented for the state state space of cluster and uh as you can see initially there is the standard practice of converting the pointer so that we can use it uh with use it in uh distributed shared memory and uh then basically we are just launching the instruction and uh as you can see here that uh we are passing in the phase out placeholder and uh we are passing in uh the address of the marrier and the count with which we want to decrease the expected thread arrival count. So basically this is the whole implementation and uh yeah so now let's talk about uh another really important topic which is about compilers and reordering and um basically what happens is that when we are you know compiling our code then what often happens is that compilers reorder the operations so that they can get some kind kind of efficiency gains. So basically if it is if it is much more performant or if even if we get a little bit of performance gain if we are launching this instruction before this instruction even though we are operating on these operations sequentially. What will happen is that the compiler will run this and then it will run this. And because the whole final operation does not depends on the order in which these are executed, it does not it does not really affect things a lot. But with GPU programming and whenever we are working with multi-threaded code, what ends up happening is that this reordering might break the correctness of the code. So if our thread A is executing this and if my thread B is executing this and uh if we reorder the operations in different ways then we might end up getting really different results and uh to prevent this behavior to prevent the compiler from reordering the operations we use this instruction. So basically in marrier dot release suffix basically provides release semantics and what this means is that we we implement a one-way hardware fence where the memory writes above this line cannot reorder below it. So basically what it means is that we create a fence and all the operations before this instruction all the operation before we launch this instructions are basically we cannot they cannot reorder after this instruction. So everything which was executing before this instruction will execute that way and ne neither the compiler nor the hardware can reorder it. And this is another important point that sometimes the compiler does not reorders it. Sometimes the hardware itself in order to achieve some latency hiding or stuff the hardware itself basically reorders a few operations and to prevent that we use this instruction. So after a thread finishes producing data that other thread will consume we have to use this instruction. So let's see how this is also implemented in in our repository and uh I have this function and uh basically we are doing nothing we are just u having an if condition so that the thread zero executes this instruction and then we convert the pointer to a pointer in shared memory and then we just launch the instruction. Pretty straightforward. We know how this works and uh yeah basically only all the threads in CTA can view this and uh basically the this is for telling us that the pointer is in the CTA state space and uh yeah basically that's it. So now let's talk about destroying a barrier. So mrier.inwell Inwell is the is the instruction which invalidates the barrier object. It basically wipes the hardware's tracking of the barrier state and it frees up the shared memory address which is uh in which the barrier is located so that we can use it for some other purposes and uh it basically converts the generic pointer to a 32-bit shared memory offset so that the GPU hardware can address the correct local memory bank. So this is basically the code for destroying the barrier and with that we have learned the whole barrier pipeline how things work. We use marrier.init instruction to create a barrier and we set an expected thread arrival count and we start with phase zero and then we have expect_tx embarer.exect_tx do.exe_tx and uh this basically tells the barrier to expect this number of uh bytes from an upcoming async copy operation and uh then we launch our asynchronous copy operation and uh we basically the other consumer threads launch this instruction and u this spins until the phase zero is And uh once all the bytes and threads arrive the barrier flips and uh the computation begins and then once the computation ends with we manually flip the phase and then we start working on the next style. So this is basically how things work when we are working with marrier and uh and uh asynchronous bul copies. So let's now move on to the a few more you know a few more amber instructions. Uh these are not very standard instructions. I couldn't find most of them in in thunder kittens or in a lot of other popular repositories. So you can think of these as some kind of a special mention. I found them in the PTX stock. So I thought that I can just you know I can just talk about them. So the first instruction is marrier. Urop and there are a lot of ways to do this and uh it is basically the standard arrive operation. But if you remember that uh when we were creating a barrier uh we were setting up an expected transaction count and uh the way it works is that whenever we reuse the barrier and whenever we flip the phase the transaction count is set back to this count. So what uh what this instruction does is that it uh permanently decrements the expected thread arrival count for the barrier. So even if you are flipping the phase and trying to reuse the barrier, the expected thread arrival count will be the previous value minus the number which we put in here. So let's look at it. So basically we have this instruction and then sam.expect_tx and nocomplete. So basically the first in uh this instruction what it does is that it works pretty normally. we just decrement the uh we set the expected transaction count but we also permanently decrement the expected thread arrival count by one. So these two instructions are here for this and uh that's what they do and this is just the whole phase state thing. So for one instruction we have it here and for other we have it as a placeholder. Okay. Then then we have dots. So it can be basically dotrelease and relaxed. Is basically the normal way of doing things. You can think of it as the difference between mrier.cope and uh mriarrier.re.escope. So basically what it does is that it uh it basically just you know specifies that no memory ordering is enforced and u if we use release then it ensures that all memory rights are performed by the thread before the arrival are visible and uh yeah and then we have dot no complete here and uh what this does is that it instructs the hardware to perform the hard the arrival without triggering the phase completion. So even if we reach you know expected transaction count zero and expected thread arrival count zero the barrier phase will not flip. So this instruction is for that we you might use it in some occasions or another but yeah not very widely used and uh not that important for most of the use cases. So now let's talk about another really important PTX instruction and that is barrier.cluster and the reason why it is so important is that we don't have a way to synchronize the threads within a cluster. We have a way to synchronize the threads within a single block and it is bear sync or sync threads but we don't have a way to do it across a thread block cluster and barrier.cluster is the ptx instruction which does that so it also guarantees that any rights to distributed shared memory before the barrier are visible to all the blocks in the cluster after the barrier. So it basically helps us getting uh the release and acquire semantics. And uh the key us use use case for that is TMA multiccast. And the way it basically works is that in the initial step thread block zero might be running but thread block one might not be running. And uh basically if block zero tries to write into block one then the program will crash. So we have to make sure that every single block in the cluster is running and by launching this operation we can do that. So now let's look at the two important instructions which we have in barrier.cluster and uh these are those two instructions. First is barrier.cluster.arive and uh what this does is that it signals the arrival to the cluster. So uh it it signals the arrival to the barrier and uh it is a non-blocking operation. So it continues executing independent instructions which are next in the queue. And uh then we have barrier.cluster.weight. And uh this is basically a blocking instruction. So it basically blocks every single thread until every single thread in the cluster has signaled arrive. So once every single uh thread in the cluster have launched this instruction then we can then this barrier will be opened and uh basically what it does is that it guarantees that all the data which is written by all uh which is written by other blocks and this block is now safe to read. So basically every single uh thread had to launch some instruction and then it had to launch this instruction so that and then we wait for it so that you know all the previous instructions before this have been executed and this is really important because you know block A can write into block B's memory. So we need a barrier to make sure that block B does not prematurely read the data before block A has finished writing. So this is basically the purpose for it and uh yeah so now let's also discuss the way it is implemented in thunder kittens and uh basically this is the way so we don't have any operands we just launch these instructions like that and uh as you can see we are launching barrier.clusterive cluster. And uh we have also these two things here. So release basically gives us the release semantics which basically makes sure that all the operations which are uh which are execute which are there before this instructions are not reordered with all the uh all the all the instructions which which are there after the instructions. So there's no reordering between the previous and the next instructions and uh align basically means that uh all the threads in the cluster are uh you know simultaneously launching it. So there is nothing all the threads in the uh I think in the are simultaneously launching it. So that is that and uh we have the next instruction which is barrier.cluster.weight weight.quire dot align and uh what this acquire means is that every single instruction which comes after it is not reordered to be before it. So that's what acquired means and aligns aligned against means that every single thread in the WOB have uh have executed this instruction. So yeah that's about it with barrier.cluster cluster and uh now let's move to asynchronous group and uh this is basically another way of uh launching uh an asynchronous copy instruction. So generally the asynchronous copy instruction is something like cp.ac async.bulk.tensor tensor and uh then we are using barrier here M barrier here and u asynchronous bul groups are the other way to synchronize and uh basically this is more lightweight than m barriers. So the way it works is that instead of you know having having a way where you know we are tracking the whole progress of this copy instruction what we have is that we launch a bunch of copy instructions and then what we do is that we batch them in a group. So I will launch a few instructions with my thread and then I'll batch them in a group and uh basically we make sure that um a bunch of these operations are done and uh until a few until a bunch of groups are uh group operations are done we don't launch the next few operations. And the way we synchronize it with everything is that once all these instructions have been released and once we have batched them in a group and once we start waiting on these group what we do is that we synchronize all the threads using sync threads and that basically means that okay all the instructions till now have been done and we'll see this later in a lot more detail. So basically this is the first instruction and what it does is that once we launch a bunch of uh asynchronous bulk copy operations using cp.async.bulk with bulk group now we just want to commit these uh these these instructions. So initially when we just launch our program for the first time we have not launched any we have not launched this instruction any time. So what it does is that it takes all the instructions which have launched till now and it patches them into a single group. So it basically packs them in a group and uh now what we can do is that we can make the unit wait until execution of n units are finished and that is the next instruction. But basically what we can see what what we can take from this instruction is that any instruction which is before which is launched before it and which is not committed into another group is now committed within this group and every single instruction which is launched after this instruction is part of a separate group and it is is still uncommitted. So every single instruction which is launched after after this is just uncommitted independent instruction and uh what we can have is that we can have multiple instructions within uh within a single group and we can have multiple groups with multiple cp.asing.pbulk operations. Okay. So next we have cp.acing.weight weight group. And this is basically the next thing which goes hand inhand with commit group and what it is is that it is a blocking instruction which make sure that the threads which are launching this wait until at most nomed groups remain pending. So basically if we have weight group zero then we'll wait for all the groups to compete complete and if we have weight group two then we'll keep uh we'll keep waiting uh waiting until the last two instructions uh the last two groups are still working in the background. So if we have let's say five groups and uh we then uh basically we then launch weight group two then this means that until the three groups are done operating then we can you know get the control to back to the our thread. So basically this count refers to the groups which are is still pending and not groups which have completed. So wait group two means that wait until the two most recently committed groups are is still pending and all the older groups must have been completed. So this is this and uh we can have this read uh read suffix in here and what it does is that it stalls the execution of the thread and it also enforces an acquire fence. So this is really important to again um to basically avoid the read after write scenarios because you risk reading the older data before you know you actually get the newer data. So this is this and uh the way it is implemented in thunder kittens is also pretty straightforward. So we set some count and then what we do is that we just launch the instruction with uh the memory clober and uh basically yeah this is for uh weight group and uh I think this is for weight group with uh with uh release semantics and uh yeah so basically this is it and uh we are basically done. Finally, we have our named barriers and these are another bunch of really important barriers and we would see them in a lot of other repositories. So the way it works is that they are kind of like sync threads but they are just for a few threads which are defined. So basically what we do is that we launch this instruction and we launch them for a particular barrier ID which basically tells us that okay this is this specific barrier and we set it with a expected thread count and this expected thread count is very much similar to the ex embar's expected thread count and uh what this does is that it blocks the execution or it non-blocks the execution but reduces the expected thread arrival count of the barrier until B number of threads arrive. So that's what is done by a barrier or named barriers. And uh as it is said here that named barriers let you create multiple independent synchronization points inside a single block. So different subsets of WOPS can coordinate without forcing unrelated WOBs to stop. And these are generally used for your for your producer consumer pipeline where your producer which are doing the copy and you know they don't have to stop for all the other threads to to complete. So what you might do is that you might have your whole producer WOB or whole producer VO group launching this instruction and uh basically launching the non non-blocking version of this instruction and just signaling at the barrier and then you can have your consumer WOS which are launching this instruction on the barrier with same ID so that they can be blocked until the whole data arrives. So this is what named barriers do and as it is said here that PTX explicitly allows different VOs to use the same named barrier with different operations such as mixing arrive and sync to build consumer producer pipelines. So that's what this is for and uh let's just go ahead and let's talk about these instructions again. So we have barrier.cta.sync.aligned aligned. And what this does is that it basically it basically creates a barrier with ID A where B number of threads have to signal arrive for this barrier to open and for them to proceed with the next instruction. Then we have barrier.ct.aligned or barrier barrier. And what this does is that it creates a non-blocking barrier which in IDA and it just reduces the arrival count of the barrier until B you know until B threads arrive. So that's what this does. And then we have two other special operations. And what these do is that if we have a condition where you know every single thread is arriving at the barrier then every single thread would arrive with a predicate zero or one and uh you know the ID is a and uh and the the the number of threads are B. And what happens is that the barrier would collect all of these predicates and it would reduce them together and reduce them in a way that it it would produce an output and that output would basically be that these are the number of threads which provided one as a predicate and these are the number of threads which provided zero as a predicate. And basically we can use this for all kinds of bookkeeping where you know there is a certain operation which was launched by n number of threads and certain operation which was launched by other number of threads and we can use it to track. We also have this this exclamation and what this is saying is that uh what this is saying is that uh give me as output the number of threads which are signaling zero and if we don't have it then it says that give me as output the number of threads which are signaling one. So if there are 128 threads and we are launching it without this bank then what it would mean is that okay if 64 threads are launching are launching uh are launching with one and 64 threads are launching with zero then it would just return 64. If there are 63 threads which launch with with with one and there are you know 65 threads which launch with zero then it would give us 63 as an output because we are not using bang here. If we use it then we would get 65 as an output if 65 threads are launching it with zero. So that's what this does and then we have this final instruction and this basically tells us to apply and or or operation. So basically if if we are using end here and uh if you are using it without this bank then it would only basically it would only give one as an output here at C until every single thread provides one as an input. So for end, every single thread must provide one as an input and as you can probably guess for or just a single thread can provide you know whatever value which we might expect and uh one or zero according to this bang and then we would just give the output. So that's how this whole thing works and that's the logic of named barriers and with that this lesson ends. Thanks for watching this lesson and yeah let's go. Okay so now let's talk about coupens map. In the previous lessons, we talked about the TMA and how we have this new unit which basically orchestrates the whole asynchronous bulk copies. And we also talked about M barriers and other synchronization ways which are really good tracking mechanism for completion of asynchronous bulk copies. But in this lesson, we're going to talk about the object which is which basically is the most important part of asynchronous pul copies and without this object asynchronous pul copies won't be possible and uh that my friends is tensorm but before starting let's talk about TMA and just review everything which we learned in the first lesson. So what is TMA? TM is the new unit which enables fully asynchronous data movement within the GPU with minimal thread involvement and basically a single thread launches DMA instruction and immediately continues execution and the whole operation is handled by the hardware in the background and uh instead of address calculation you get something called a descriptor which basically gives you a good description of what your tensor which you are moving looks like or what any kind of data which you are moving looks like and uh TMA can transfer data to shared memory of multiple SMS simultaneously. This is what we call the multiccast feature of asynchronous bul copies and the TMA can handle 1D to 5D tensips. Great. So now let's talk about how H100 gets the perfect asynchronicity and just review everything which we looked at you know a lesson before. So we need all the units to be busy all the time and H100 gets it right by having multiple buffers loading using TMA and having tensor cores do the opa operation simultaneously. So if you remember the previous lesson, we we reviewed a basically a producer consumer pipeline where we saw that we had two buffers and we were loading data using the tensor memory accelerator and then tensor coursees were doing the operation simultaneously. So basically this is the way we get you know all the units to be busy all the time in H100. And uh the more important point is that you don't need a lot of complex engineering which you previously needed. So if you remember you know writing code in the previous generations of the GPU you would remember that you had to index everything and you had to you had to manually manage everything and then there would come sizzling and all this different stuff which you had to handcode and you had to index and then SMS will be basically calculating all the indexes and then going to the global memory fetching data. That was a lot of hard work and complex engineering. But now we have descriptors and these descriptors handle everything from data to layout to optimization and just they just leave a lot of heavy lifting to the hardware so that you guys can focus on writing more and more kernels and work on pipelines. So now let's talk about the need for descriptors. Why do we need descriptors? And uh let's just see why this whole asynchronous pipeline is powered by descriptors. So descriptors are the piece of puzzle which enable asynchronous nature of H100. We talked about this in earlier times. The whole fetching of data would happen using indexing where threads would need to compute the address of the data which they wanted and then fetch and you know that was really complicated and u this was fixed a bit in a where you know the copy instruction was non-blocking but thread is still had to compute the addresses for every 16 bytes so it was not really fixed and SMS were thred to the copy engine So now with H100s the descriptors encapsulate the entire transfer state meaning that every single bit of in information which is required to complete the job is contained in a 128 byt object in the memory. And uh you can basically just use this 128 byt object and you can give it to the tensor memory accelerator and it will make sure that the data is transferred asynchronously and in the background without any SMS involvement. So this is great. So this was the we talked about the way we used to do it in the previous generations but now we just use the CUDA API to create a tensor map. We encode it with the information required and then we just launch the operation and provide this tensor map to you know or testate the whole operation. So now let's talk about the the the internals of the tensor map and what are the information which we instill into a tensor map. So first of all, tensome tensor map is a descriptor a 128 bit descriptor which stores information about the memory based pointer tensor shape stride data type alignment and swizzling memory space basically where the tensor exists and order and rank and uh optionally it can also handle tiling or interle layouts. So these are a lot of things which uh tensor map can handle. So now let's get into it and now let's learn how to create a tensor map object. So we use this to create a tensor map object. It is created in host memory as a local variable and uh it has a size of 128 bytes and uh it must be 128 bytes a line. So that's pretty important and uh it's not a regular pointer but a specified data structure which is created by Nvidia. So you can actually look at uh some of uh this in Nvidia's documentation where they talk about how they created it and all that stuff and u then we use coupens map and code tile this function to encode the values in tensor map. So now let's see how this information is encoded. So here we can see that we are basically first of all pointing the pointing uh the to the tensor map. Then we are telling it the data type of the tensor which we are going to transfer. Then we are talking about the rank of the tensor the global address of the tensor its dimension. It's a stride then the dimension of the tile and the element stride and then the interle pattern swizzling pattern L2 promotion and out of bounds fill and uh we are going to learn all of these in this lesson. So let's get on with it. So first of all pretty simple this is the tensor map object which we create using this this which we learned here and uh basically this is just a pointer to the object and uh we need to pass this because you know the the the we need to know where the tensor map exists so that we can access it and then use it to enable the whole asynchronous data movement. Then we have this and what this does is that it determines the actual data type of the data and uh which you are going to copy from HBM or from shared memory and uh TM engine uses this to automatically get the memory alignment and the transfer size. So this is important and we have a lot of data types which we can use. So just review it. See if you know everything. If you don't then just learn about it. And uh yeah, but uh one of the things which I want to discuss a bit more about is this flush to zero mode. What is flush to zero? What does it mean? So flush to zero mode is just a floating point arithmetic operation which basically handles the out of bound numbers. So basically if you your tensor has some out of bound numbers then it will replace them with zeros and instead of you know computing them normally or crashing the whole program. So that's really important. Then let's talk about the tensor rank and global memory address which we are passing in here. So okay. So now let's talk about tensor rank and global memory address. So tensor rank is the number of dimensions of tensors. So basically if I have a 1D tensor then uh it will be one. If I have a 2D tensor then it will be two and so on. It is not the matrix rank from linear algebra. So that is one thing to take care of. And uh basically the tensor which we are talking about is the tensor in global memory from which we are copying data. So it is not the the the tensor which is you know which we are copying itself but it's the tensor from which we are copying the data. And then global memory address here is the memory address of the tensor in HPM. And uh yeah basically that's it. Then we have global dimension. So global dimension is an array which specifies the size of each dimensions of the tensor in terms of number of elements. So basically just the size of each dimension and uh if we have an rd dimensional array then we can arrange data in the following way where the first element of the global dimension tensor is the innermost dimension. basically it represents the size of the innermost dimension and the last element the earth element is basically r minus one because i