• Home
• Descriptor
• Announcements
• Schedule and Notes
• Assignments
• Feedback
• Sponsors

• View
• Print
• Logout ()
  

• PmWiki
• Installed Recipes

Execution Core

There are 32 threads per warp, and they are executed together on one SM. All 32 threads execute the same instruction because they share a sequencer, and all 32 move together. Furthermore, for each warp, its instructions are executed sequentially. The SM can execute 8 (G80) or 2 independent half warps, 2×16 (Fermi) of these threads per cycle, and will take 4 (G80) or 2 (Fermi) cycles to execute all 32 threads.

A SM is given one, but potentially up to eight, thread block(s) to execute at a time. Each block consists of up to 32 warps (1024 threads). These warps can be executing different instructions, even though within the warp, all the threads execute the same instruction. Every 4 (G80) or 2 (Fermi) cycles, the SM chooses a warp that is ready to run, and runs it for the next 4 (G80) or 2 (Fermi) cycles. There is no delay when switching to a different warp at the end of the 4 or 2 cycles. (Note, for G80, there is a maximum limit of 768 threads, or 24 warps, per SM. For Fermi, 1536 threads, or 48 warps). Note that each SM in Fermi has two parallel warp schedulers and sequencer. One scheduler handles even warps and the other the odd warps.

All the thread blocks that go to the SMs execute the same program, called a kernel. (Note, this is a different notion from executing the same instruction in the program.) The GPU will execute all of the thread blocks for a given kernel, and then move on to the next kernel. Thread blocks within a kernel are arranged in a grid, and are addressed by two-dimensional indices.

Scheduling Hierarchy

• Kernel (also known as grid):
  • Composed out of thread blocks.
  • You can have arbitrarily large number of thread blocks with in a kernel.
• Thread Block (Also called CTA: cooperative thread array):
  • It will usually have either one or two dimensional indices.
  • You can have up to 32 warps within a thread block.
• Warp
  • Consist of 32 threads
• Thread:
  • It also has either one or two dimensional indices.
  • Logical concept.
    • Threads within a warp are not independent.
    • They should execute the same instruction.
    • If they need to diverge, then mask operations are required: this is managed by hardware)
• This units of execution asking their names (indices) can be viewed as SPMD model.
  • SPMD: Every one executes same code (kernel) and it can do different things by asking “Who am I?”

Overview of scheduling hierarchy

• The CPU tells GPU what to do in the granularity of kernel.
• GPU thread scheduler (giga thread engine) tells cores what to do in the granularity of thread block.
• Core tells ALUs what to do; this happens at the granularity of warp (it can be seen as vector instruction).
• Finally, there are operations within a lane. these are called threads.

Reason for introducing concept of threads instead of programming using vector instructions

• This is meant to present abstraction where users do not need to think about vector instructions.
• They let you program on SIMD machine without having to think exactly what is happening in every cycle, who is communicating with whom, etc.

Scheduling Thread Blocks for Execution

• No need to think dependencies but need to think resources that is available.
  • Need to think about available registers and available local shared memory. Basically the scheduler needs to think about software manage resources.
  • How many thread blocks can you put in one core? Min(Total scratch pad / CTA Scratch pad, Total registers / CTA registers, 8)
    • The number 8 is artificial number. Minimum number of hardware management registers are required and this limits maximum number of blocks you can have for a core concurrently. The number 8 is actually a large number and you will rarely put 8 blocks to a SM.
• It could look like out of order scheduler.
• The scheduler is called GigaThread Engine in Fermi architecture.
• Each core can handle multiple thread blocks concurrently.

Scheduling Warps for Execution

Each SM implements a hardware scoreboard scheduler to determine which warp to run next. (Note that because the scheduler chooses between warps, and not threads, it has much less work to do.) Warps which have all of the operands ready for its next instruction are eligible for execution. The scheduler uses a round-robin policy to select an eligible warp to execute.

The scoreboard keeps track of all register operands of all instructions in the instruction buffer to prevent hazards.

Once selected, the same instruction is executed for all 32 threads in the warp. This takes 4 cycles for G80 or 2 cycles for Fermi, because the SM has 8 SPs or 16 SPs respectively, which can process one thread per cycle.

In the case of a long-latency memory operation, the currently-executing warp can still proceed, as long as it doesn’t use the new data as an operand in its instructions.

Overlapping communication and computation

• Can be done by software pipelining.
• Device driver supports this. It figures out what DMA operation can be executed next and which kernels are ready for execution.

Locality vs Parallelism

• How many threads per SM are appropriate?
• How many threads are required to tolerate memory latency?
  • EX) Assuming one forth of instructions are memory operations and they are accessing global memory, we need minimal of 13 warps to fully tolerate 200 core cycle memory latency.
    • 13 warps x 4 cycles per warp x interval between memory operations (4) equals to 208 cycles
• Throwing too many threads to SM will limit locality since it will reduce available memory space for each thread
• Why do we want locality?
  • saves bandwidth
  • saves power
  • gives lower latency
• Do you have reasonable amount of total parallelism per SM?

Communication

• Two shared name spaces: global memory and shared memory
• Threads belonging to the same thread block may communicate with each other through shared memory
• Execution kernels and thread blocks may communicate with each other through device memory (global memory)

Synchronization

• Thread blocks are implicitly synchronized at the end of kernel execution.
• How do threads within a thread block synchronize?
  • __syncthreads() is used
  • Every other thread within the thread block should reach this point to proceed
• Thread blocks must be independent
  • Presumed to run to completion without pre-emption
  • Thread blocks cannot be removed arbitrarily
  • If you are waiting for more number of thread blocks than resource available you face deadlock

Memory

• Registers has one instruction latency
• L1 is not coherent, accessing takes 2 cycles but equivalent to one instruction
• Not all memory operations were able to utilize L1 cache on G80 or G92; only special instructions were able to utilize the cache
• Access to L2 needs to go through gigantic cross bar.
• L2 is shared between all cores, and it is coherent because there are only one
• Off chip memory space has two separate name space; private and global
  • Private is local memory but kept in offchip

Design considerations for building a register file for GPU

• Multiple smaller SRAMs is better than one big SRAM.
• How many banks do you need? What is reasonable number?
  • 3 Read 1 Write * 32 threads (a warp)
  • 3*32 = 96 banks? it is too small; overhead of peripherals will be too high.
  • 4 banks? not enough banks, not enough to do independent accesses.
  • 8 banks? maybe reasonable to support two independent schedulers.
  • Basically, the more banks you have, the fewer bank conflicts you have.
  • You need to think latency, power, bank conflicts, locality, width of access, and aspect ratio of SRAM. These will fight with each other, so it is not a simple trade off.