Questiоn 1 CUDA [18 pоints] [15 minutes] (1.1) [12 Pоints] With а 32‐wide wаrp аnd 64B cache lines (no subblocking/sectoring), how many memory transactions will be generated per warp, for each of the following half‐precision float array (16 bits) accesses? Assume a threadblock size of (256,1,1) and no inter‐warp merging. (1.1.1) A[threadIdx.x*2i] (where i >= 0) Your answer must cover all values of i. (1.1.2) A[threadIdx.x+2i] (where i >= 0) Your answer must cover all values of i. (1.2) [6 Points] With a warp size of 64, and a 1D threadblock of size [x,y,z]=[256,1,1], what is the SIMD efficiency of line 7? Assume that 2i is a constant where i >= 0. Your answer must cover all values of i. Question 2 CUDA Optimizations [20 Points] [25 minutes] (May be better to attempt at the end) We want to compute the prefix sum of a vector of length n, which is defined as Output[i]=∑j=0iInput[j]{"version":"1.1","math":"∑j=0iInput[j]"}for all i from 0 to n‐1. For example, if the input is [5,0,8,6], the output is [5,5,13,19]. Note that unlike a single plain summation, here every prefix’s sum is calculated (as many prefixes as the input vector length). Prefix sum has numerous use cases including sparse computations, parallel task assignment. Assume the vector length is a large power of two and a multiple of 32, the warp size. (2.1) (6 points) What is your basic strategy to parallelize this computation? Hint: Assuming the vector is drawn left to right with index 0 at the left, compute in parallel the sums (plural) of increasingly longer sub‐sequences. This strategy is similar to that of parallel summation of a large number of numbers. Obviously, the computation is an iterative process. Assuming a thread per vector element (a la classic CUDA) where tid = vector index. There is no iteration 0 (believe me ‐ it is easier this way). Show below by drawing lines from the cells in the row above that are added to the cell in the row below which is updated to show how the summation would proceed. (2.2) [4 points] Write an expression for the cells (indices) whose values are ready after iteration i (i >= 1) using the variable “tid” and “i”. (2.3) [4 points] Write an expression for the thread ids that are active in iteration i, using the variable “tid” and “i”, keeping in mind some cells are already ready (hint: use binary representation of thread ids and bit masking). (2.4) [6 points] Using the variables “i”, “tid” and “vector[]” which is updated in place, write an expression for the cells (indices) being added and the cell (index) being updated in iteration i, only for the active threads from Q(2.3). State the expression as an assignment statement x = y+z. Question 3 GPU Pipeline [22 Points] [25 minutes] Consider the design of a modern GPU covered in class. Pictured here is a design with a single warp scheduler (issue) and only two parallel pipelines (ALU and memory). (3.1) [3 points] If more than one instruction per warp were allowed to be in the operand collector stage of the pipeline, what problem would emerge? (3.2) [3 points] How would you fix the problem in Q(3.1) in a light-weight manner while allowing more than one instruction per warp to be in the operand collector stage? Be specific in your answer. (3.3) Control-flow divergence [16 points] (3.3.1) [4 points] In control flow-divergent code, what two other workload characteristics, one for code and the other for data, make such divergence severely hurt performance? The question already states that the code is divergent so answers saying divergence hurts performance will receive zero credit. (3.3.2) [3 points] If nested branches are rare and loop branches do not diverge (there is other substantial branch divergence), how would you handle branch divergence better than modern GPUs? (3.3.3) [5 points] How would you change the rest of the GPU microarchitecture to address the problems caused by your above branch strategy in Q(3.3.2)? Address at least three substantial components. (3.3.4) [4 points] What code characteristics would keep the above changes effective and reasonable (hint: recall your answer to Q(3.3.1))? What (possibly unrealistic) code characteristics would make the changes the absolute minimum? Question 4 Performance [20 Points] [25 minutes] In untiled, uncached FP16 MM of two N x N matrices A and B to produce an N x N matrix C, assuming an inner product dataflow, each row of A and each column of B are brought from memory to produce a cell of C – i.e., 2 N * 2 bytes fetched for N multiply-accumulates to produce a cell of C which is written to memory. Repeating this process N2 times gives operations (Ops) per byte (counting multiply and accumulate separately) as Thus, untiled MM is ridiculously memory-bound. Now consider a t x t tiled FP16 MM (caching implicit). The cache can hold only one tile of each of A, B, and C. Here is a brief description of standard tiled MM. Assume the A tile is fetched from memory and held, while the B and partially-updated C tiles are fetched from memory to compute the MM for the entire tile, and update the partial C tile which is written to memory. Then the B tile moves horizontally and is fetched from memory (while the A tile is held), one step at a time, to produce C tiles until the horizontal end of the B matrix. Then the A tile moves down one step and is fetched from memory, and the B and C tiles are brought from/written to memory again. (4.1) [4 points] For a given A tile position, how many operations occur (in terms of N and t) counting multiply and accumulate separately? (4.2) [6 points] For a given A tile position, how many bytes are fetched from/written to memory (in terms of N and t)? (4.3) [2 points] How many A tile positions are there? (4.4) [4 points] What is the operations per byte for tiled MM? Show your work. (4.5) [4 points] If the GPU computes at 64 T Flops/s and the memory bandwidth is 1.5 TB/s, when does tiled MM stop being memory-bound for large matrices? Question 5 Synthesis [20 Points] [25 minutes] Assume a GPU workload has a two-phase behavior. Phase 1 has significant spatial locality across the warps of a threadblock well beyond what the cache can exploit, and Phase 2 has little spatial locality well below what the cache exploits but significantly higher temporal locality than usual. In both cases, the cache (and in general on-chip) capacity is highly limited. The workload switches dynamically and unpredictably between these two phases, which are long. There is little control-flow divergence in either phase. Assume a GPU microarchitecture similar to that of Q3, a GTO warp scheduler, and a high-performance GPU+memory architecture in general. Merely repeating the schemes from the papers will receive zero credit. You may not change the cache configuration or add more on-chip SRAM. Please read ALL the questions before answering, so you don’t answer earlier than where it is asked. The questions may seem open-ended and generic, but are crafted carefully so that there is only one correct answer. Be concise and specific in your answer. Please do not “dump core”! (5.1) [3 points] How would you detect phase 1? (5.2) [4 points] How would you exploit phase 1 for better performance? (5.3) [2 points] What strategy would not work for phase 1? (5.4) [3 points] How would you detect phase 2? (5.5) [4 points] What is the impact of phase 2 on the cache? (5.6) [4 points] How would you handle phase 2 for better performance? Remember: No discussion on Piazza or Disqus after the exam until online exam window closes. Congratulations, you are almost done with the Midterm Exam. DO NOT end the Honorlock session until you have submitted your work. When you have answered all questions: Use your smartphone to scan your answer sheet and save the scan as a PDF. Make sure your scan is clear and legible. Submit your PDF as follows: Email your PDF to yourself or save it to the cloud (Google Drive, etc.). Click this link to submit your work: Midterm Exam Return to this window and click the button below to agree to the honor statement. Click Submit Quiz to end the exam. End the Honorlock session.
Eаch оf the fоllоwing is аn importаnt part of personalized medicine EXCEPT: