ECE 408 MT1 Study Guide

(T/F) Host code can transfer data from host to device directly, and device code can transfer data from device to host directly.

False, host code can transfer data between host and device (using cudaMemcpy) but device code generally cannot perform data transfers to the host directly. Device code mostly can r/w per thread registers, and r/w per-grid global memory.

(T/F) Starting around 2004, computer architects began incorporating sophisticated branch prediction logic and increasing clock frequencies more rapidly with each process generation—as transistor sizes shrank—in order to reduce chip power consumption.

False, Dennard scaling, which allowed clock speeds to increase with each process generation while maintaining power consumption entered around 2005-2006, after that focus shifted to emphasizing parallelism and specialization.

(T/F) The sign function (sign(x) = 1 for x > 0, sign(0) = 0, and sign(x) = −1 for x < 0) can be used as a great activation function for DNNs.

False, since sign is not a smooth function, its zero derivative away from zero

(T/F) Consecutive threads within the same block are assigned to the same or contiguous warps, and consecutive blocks within the same grid are allocated to the same or neighboring SMs.

False, while it is true that consecutive threads within the same block are assigned to the same/contiguous warps, it is not guaranteed that consecutive blocks within the same grid to be allocated to the same or neighboring SM’s.

(T/F) Control divergence can lead to inefficiency because if threads from the same warp take different paths, all threads must wait until the longest path is completed, potentially leading to idle cycles for some threads if paths are different lengths.

True, the hardware executes these paths in multiple passes, where only the threads on the active path execute while the others are inactive.

(T/F) A constant memory read is not always faster than a global memory read, especially when only a few threads read from constant memory.

True, when only a few threads read from constant memory the benefit of caching and broadcasting is reduced, so the access may not be faster than global memory.

(T/F) The CUDA compiler generates an intermediate representation called PTX, which is not specific to any particular GPU. At runtime, the GPU driver just-in-time compiles the PTX into GPU-specific assembly.

True, the cuda compiler generates an intermediate representation called PTX (Parallel Thread Execution), which is a virtual ISA code not specific to any particular GPU architecture

(T/F) Each individual thread has its own copy of the local variables stored in registers.

True, each thread has its own set of local variables, typically stored in registers

(T/F) We can learn NAND with a perceptron.

True, since NAND is the negation of AND, and AND is linearly separable and can be learned by a single layer perceptron NAND is able to be learned as well.

If we want to allocate an array of n floating point numbers in the GPU global memory and have a pointer variable called device_array point to this array, what is the correct call to cudaMalloc()

on the host side? Circle the correct answer.

1. cudaMalloc(device_array, n*sizeof(float));

2. cudaMalloc((void )device_array, nsizeof(float));

3. cudaMalloc((void)&device_array, n*sizeof(float));

4. cudaMalloc((void **)&device_array, n*sizeof(float));

5. None of these answers are correct

The standard is cudaMalloc(void **devPtr, size_t size), so it would be 4. cudaMalloc((void **)&device_array, n*sizeof(float));

If we want to copy an array, host, of 5 floating-point numbers from the host memory to the GPU constant memory and have a pointer variable called device_array point to this constant memory array, what is the correct call on the host side? Circle the correct answer.

1. cudaMemcpy(device_array, host, 5*sizeof(float), cudaMemcpyH2D);

2. cudaMemcpyToSymbol(device_array, host, 5*sizeof(float));

3. cudaMemcpy(host, device_array, 5*sizeof(float));

4. cudaMemcpyToSymbol(host, device_array, 5*sizeof(float));

5. None of these answers are correct.

Since we need to copy to constant memory, we use cudaMemcpyToSymbol(filter_c, filter, FILTER_DIM * sizeof(float)); So the correct answer is 2. cudaMemcpyToSymbol(device_array, host, 5*sizeof(float));

3. How many CUDA threads are in each block as the result of the following kernel call?

#define VECTOR_N 1024

#define ELEMENT_N 256

...

scalarProd<<VECTOR_N, ELEMENT_N>>>(d_C, d_A, d_B, ELEMENT_N);

1. 1024*256

2. 1024

3. 256

4. 1024+256

Kernel launch syntax is:

kernel<<<gridDim, blockDim, sharedMemSize, stream>>>(kernel_arguments);

Where gridDim specifiies the number of thread blocks in the grid (can be 1d,2d,3d)

Where blockDim specifies the number of threads per block

sharedMemSize (optional) specifies the amount of shared memory in bytes to allocate per block

stream (optional) specifies the cuda stream in which the kernel is launched, defaults to 0.

so the answer is 3. 256

4. How do we declare a 5x5 constant memory floating point array Mc?

1. constant float[5][5] Mc;

2. constant float Mc[5][5];

3. __constant__ float[5][5] Mc;

4. __constant__ float Mc[5][5];

5. None of these answers are correct.

4. __constant__ float Mc[5][5];

The __constant__ keyword is used to declare constant memory variables

For a vector addition, assume that the vector length is 8000, each thread calculates 8 output elements, and the thread block size is 512 threads. The programmer configures the kernel launch to have a minimal number of thread blocks to cover all output elements. How many threads will be in the grid?

1. 1000

2. 8196

3. 8192

4. 1024

5. 8200

6. None of these answers are correct

Step 1: calculate # of threads needed: 8000/8 = 1000 threads

Step 2: calculate # of TB needed. blocks = 100/512 = 2 blocks (rounding up)

Step 3: calculate total # of threads in grid: 2 × 512 = 1024 threads

Correct answer: 4. 1024

What are the possible values of dst[0] after this kernel execution?

__global__ void kernel(char *dst) {

dst[0] = blockIdx.x;

}

// dst is a pointer to an array of one char allocated on

// the device and initialized to value of 3, e.g., dst[0]=3

kernel<<<2,1>>>(dst);

1. 0

2. 0 or 1

3. 3

4. 0 or 1 or 3

5. 1

Kernel is launched with 2 blocks each with 1 thread, and it can writes its blockIdx.x value to dst[0]. Since there is no guarantee in the order of block execution, it can be overwritten by 0 or 1.

2. 0 or 1

A particular CUDA device’s streaming multiprocessor (SM) can take up to 1536 threads and up to 4 thread blocks. Which of the following block configurations could guarantee the SM be fully utilized?

1. 256 threads per block

2. 384 threads per block

3. 576 threads per block

4. 1024 threads per block

5. None of these answers are correct

2. 384 threads per block,

since this exactly matches the maximum number of blocks per SM and fully utilized the 1536 threads.

Consider a kernel where for every floating-point operation (FLOP) performed, it accesses 8 bytes from global memory. What can you infer about this kernel's performance characteristics based on its bytes/flops ratio assuming a GPU with global memory bandwidth of 150 GB/s and 1,000 GFLOP/s of compute power?

1. The kernel is likely compute-bound since it performs more computations relative to memory accesses.

2. The kernel is likely memory-bound since it performs fewer computations relative to memory accesses.

3. The kernel optimally balances between memory accesses and computatons, ensuring maximum GPU utilization.

4. The bytes/flops ratio is unrelated to performance botlenecks and provides no useful insight.

The bytes per flop ratio is 8 B/FLOP, so therefore the kernel requires 8 bytes of memory access for every floating point operation.

The max FLOP rate supported by memory bandwidth is: mem bandwidth/B per FLOP = 150 GB/s / 8 B/FLOP = 18.75 GFLOP/s

The GPU’s peak performance is 1,000 GFLOP’s, which is higher than 18.75 GFLOPs.

So therefore, the answer is 2. The kernel is likely memory-bound since it performs fewer computations relative to memory accesses.

Which type of memory will be most suitable (fastest accesses) for their purpose?

A. Suppose your kernel code requires certain threads to read data items writen by other threads in the same thread block.

shared memory, since shared memory is on chip, low latency, and accessible by all threads within a block, enabling fast inter-thread communication.

Which type of memory will be most suitable (fastest accesses) for their purpose?

B. Suppose your kernel code requires all threads to share the data items and those data items remain same throughout the execution.

constant memory, since it is cached, read only memory optimized for broadcast access patters where all threads read the same data

1. Consider a DRAM system with a burst size of 512 bytes and a peak bandwidth of 240 GB/s. Assume a thread block size of 1024 and warp size of 32 and that A is a floating-point array in the global memory.

What is the maximal memory data access throughput we can hope to achieve in the following access to A?

int i = blockIdx.x * blockDim.x + threadIdx.x;

float temp = A[4*i] + A[4*i+1] + A[4*i+2];

1. 240 GB/s

2. 180 GB/s

3. 120 GB/s

4. 60 GB/s

5. None of them are correct.

Given:

DRAM burst size = 512 bytes

Peak bandwidth = 240 GB/s

TB size = 1024 threads

warp size = 32 threads

Each thread accesses 3 elements of A, the elements accessed by consecutive threads are spaced by 4 floats apart because of the 4*i indexing.

T0: A[0], A[1], A[2}

T1: A[4], A[5},… and so on

DRAM burst size is 512 bytes, each float is 4 bytes, so 512 bytes correspond to 128 floats, only some of the data is used by threads though due to the memory access pattern (3 out of 4 floats)

The peak bandwidth is 240 GB/s, since only ¾ of the data fetched is useful, effective throughput is 240 GB/s * ¾ = 180 GB/s

Therefore, the answer is 2. 180 GB/s

For a 4x16 2-dimensional array M, elements are placed into the linearly addressed memory space according to the row major convention. Assume that we are using 4×4 blocks and that the warp size is 4, k is the number of loading iteration. Which of the following array accesses in the kernel code has coalesced memory access?

A: M[k*Width+blockIdx.x*blockDim.x+threadIdx.x]

B: M[(blockIdx.y*blockDim.y+threadIdx.y)*Width+k]

1. A

2. B

3. Both

4. Neither

Block size: 4 x 4 threads
Warp size: 4
k is the number of loading iterations
Width = 16 (number of columns in M)

index = r x Width + c

Access A:
M[kWidth+blockIdx.xblockDim.x+threadIdx.x]
k Width corresponds to row k
blockIdx.xblockDim.x+threadIdx.x corresponds to the column index

We have a fixed row and consecutive col, so accesses are coalesced.

Access B:
B: M[(blockIdx.yblockDim.y+threadIdx.y)Width+k]
blockIdx.yblockDim.y+threadIdx.y)Width corresponds to row
k corresponds to the col

We have a fixed col, and consecutive row, so accesses are not coalesced

Final Answer: 1. A

(T/F) GPUs are commonly categorized as latency-oriented devices.

False, GPU’s are throughput oriented since they contain many small (not powerful) ALU’s

(T/F) CPU execution and GPU execution could overlap with each other.

True, they are free to work concurrently and execute asynchronously.

(T/F) Shared memory is private per thread block.

True, only threads in the same block can access

(T/F) Global memory is a separate hardware unit from the GPU core.

True, global memory is usually implemented as off-chip DRAM, which is physically separate from the GPU cores.

(T/F) The "grid-block-thread" hierarchy is a software interface and does not describe the GPU’s physical architecture.

True, this hierarchy is a software execution model used to organize and manage parallel execution of threads.

(T/F) In recent NVIDIA GPU architecture, a warp consists of 32 threads that execute the same instruction at a time.

True, a warp consists of 32 threads following the SIMT (single instruction, multiple threads) model.

After doing Lab 4, Ze has mastered the usage of constant memory. While reviewing his Lab 1 for his midterm preparation, he suggests copying two input vectors to constant memory. Evaluate whether this strategy is beneficial by briefly stating one benefit and one drawback (each in a phrase for no more than SIX words) of using constant memory for this purpose.

Benefit: memory cached, efficient for broadcasting

Drawback: limited size, not scalable for vectors

Which factor(s) is/are important for achieving optimal global memory coalescing in CUDA?

a) Making each thread within a warp accesses consecutive memory address.

b) Maximizing the number of global memory accesses per kernel launch.

c) Using constant memory for all data accesses.

d) Aligning kernel execution to the size of the L2 cache

a, since global memory coalescing occurs when threads in the same warp access consecutive global memory locations, allowing the hardware to combine these accesses into a single memory transaction

Which of the following is a correct consideration about the function call __syncthreads()?

a) It ensures global synchronization across all threads in all blocks executing a kernel.

b) It must be used at the beginning of a kernel to synchronize thread execution.

c) It synchronizes all threads within a block but using it excessively can lead to deadlocks if not all threads reach the synchronization point.

d) It is optional and can be omitted for kernels that do not access shared memory.

c

What is the grid size and block size, respectively, that will be produced as a result of the following kernel call?

scalarProd<<<1024, 256>>>(d_C, d_A, d_B, ELEMENT_N);

a) 256, 1024

b) 1024, 256

c) 256, 1024+256

d) 1024, 1024+256

The kernel call arguments are <<< numBlocks, numThreadsPerBlock >>>

numBlocks is also known as the grid size (num blocks in the grid)

numThreadsPerBlock is also known as the block size (num threads per block)

So the correct answer is b) 1024, 256

Which line(s) below don’t have a memory-coalesced pattern? Assume that array g is in global memory and BLOCK_WIDTH = 32.

a) int a = g[threadIdx.x];

b) int a = g[threadIdx.x * BLOCK_WIDTH + blockIdx.x];

c) int a = g[BLOCK_WIDTH / 2 + threadIdx.x];

d) int a = g[blockIdx.x * BLOCK_WIDTH / 2 + threadIdx.x];

a. Threads access consecutive elements, t_0 accesses g[0], t_1 accesses g[1], so therefore this is a classic coalesced access

b. For thread t, the index is t * 32 + blockIdx.x, so threads access memory locations spaced by 32 elements, so not coaleced

c. Threads access consecutive elements starting at offset 16, therefore coalesced access

d. Threads access g[16 * blockIdx.x + t] for thread t. Within a block, threads access consecutive elements starting at 16 * blockIdx.x, therefore coalesced access

Final answer: B

For a Tiled-matrix multiplication kernel, if we use a 16 x 16 tile, what is the reduction of memory bandwidth usage for input matrices A and B?

a) 1/8 of the original usage

b) 1/16 of the original usage

c) 1/32 of the original usage

d) 1/64 of the original usage

The global memory accesses are reduced by a factor equal to the tile width

b) 1/16 of the original usage

Consider the following code:

kernel<<VECTOR_N, ELEMENT_N>>>(d_C, d_A, d_B, ELEMENT_N);

• Q: How many CUDA threads are in each block as the result of the

following kernel call?

A: ELEMENT_N

Consider the following code:

kernel<<VECTOR_N, ELEMENT_N>>>(d_C, d_A, d_B, ELEMENT_N);

Q: How many CUDA threads will be created as the result of the

following kernel call?

A: VECTOR_N * ELEMENT_N

Q: For a vector addition, assume that the vector length is 16000,

each thread calculates 8 output elements, and the thread block

size is 256 threads. The programmer configures the kernel launch

to have a minimal number of thread blocks to cover all output

elements. How many threads will be in the grid?

A:

• How many threads do we need? 16000/8 = 2000

• How many blocks of threads do we need to run 2000 threads?

ceil(2000/256) = 8

• Thus, how many threads will be running? 8 * 256 = 2048

Q: A CUDA kernel is launched with 512 thread blocks each of which

has 256 threads. If a variable is declared as a local variable in the

kernel, how many versions of the variable will be created through

the lifetime of the execution of the kernel?

A:

• How many threads will be created? 512 * 256 = 131072

• So, there will be as many copies of the local variable, one in each thread.

Q: A CUDA kernel is launched with 32 thread blocks each of which

has 32×32 threads. How many threads are in a block and how

many threads are launched altogether?

A:

• How many threads in a block will be created? 32 * 32 = 1024

• How many threads across all blocks? 1024 * 32 = 32,768

Q: Suppose you have a 2D grid of blocks where each block is

configured as a 16×16 thread block. If the grid dimensions are

32×16 blocks, calculate:

• The total number of threads per block.

• The total number of threads in the entire grid.

A:

• How many threads in a block will be created? 16 * 16 = 256

• How many thread blocks are created? 32 * 16 = 512

• How many threads across all blocks? 256 * 512 = 131,072

Q: You have a 3D grid of blocks with dimensions 8×4×2, and each

block is configured as a 4×8×8 thread block. Calculate:

• The total number of threads per block.

• The total number of threads in the entire grid.

A:

• Total number of threads per block: 4 * 8 * 8 = 256

• Total number of blocks in the grid: 8 * 4 * 2 = 64

• Total number of threads in the entire grid: 256 * 64 = 16,384

Q: A particular CUDA device’s streaming multiprocessor (SM) can

take up to 1536 threads and up to 4 thread blocks. Which of the

following block configurations could guarantee the SM be fully

utilized?

• 256 threads per block, 384 threads per block, 576 threads per block

A:

• 1536 / 256 = 6 thread blocks – too many for SM

• 1536 / 384 = 4 thread blocks per SM – just the right number

• 1536 / 576 = 2 thread blocks per SM – not enough to fully utilize the SM

Q: A 1D array of N floating point elements is to be processed in a

one-element-per-thread fashion by a GPU. The target GPU has 8

SMs, each with 16 SPs. What is the best execution configuration for

this kernel?

A:

• We do not know the max number of threads the SM can support.

• But even SM count and SP/core count alone are insufficient; one needs device

residency limits plus the kernel’s register/shared-memory footprint and a

performance measurement.

Consider the following CUDA kernel:

global void do_work(int q, int *A) {

int result = 0;

if (q < 5) {

	result = threadIdx.x;
}

A[threadIdx.x] = result;

}

Q: Is there is a control divergence in this code?

A: No since the value of q is the same for ALL threads in the thread

block. For a thread to diverge, its execution path must depend on

something unique to the thread, e.g., on threadIdx. More generally,

divergence requires a predicate that differs among active lanes

within the same warp, not necessarily a value unique to the thread.

Q: Consider a 2D input matrix of 256 rows and 4 columns. We use

16 x 16 block of threads to perform operations on the input matrix

so that each thread processes exactly one input element. How

many blocks need to be launched?

A:

• 256 rows * 4 columns = 1024 elements to compute

• 16 * 16 = 256 threads per block

• 1024 / 256 = 4 blocks of threads

• Note that here we ignore the 2D geometry and assume our own index

scheme (explicitly flattened/remapped), which is perhaps more that one

typically wants to do 

Q: Suppose your kernel code requires certain threads to read data

items written by other threads in the same thread block. Which

type of memory will be most suitable (fastest accesses) for this

purpose?

• Register

• Shared memory

• Global memory

A:

• Shared memory

Q: What are the possible values of *dst after this kernel execution?

__global__ void kernel(char *dst) {

dst[0] = blockIdx.x;

}

// dst is a pointer to an array of one char allocated on

// the device and initialized to value of 3, e.g., dst[0]=3

kernel<<<2,1>>>(dst);

• Either 0 or 1

• This is a data race. CUDA does not guarantee block execution order, so the final

writer is nondeterministic.

Q: Which expression produces the more efficient global-memory

access pattern for the specified launch?

int tx = blockIdx.x * blockDim.x + threadIdx.x;

int ty = blockIdx.y * blockDim.y + threadIdx.y;

A) B[ty * Width + tx] = 2 * A[ty * Width + tx];

B) B[tx * Width + ty] = 2 * A[tx * Width + ty];

A: Cannot tell. We do not know how the kernel was launched.

Q: Assume four-byte floats, 512-byte-aligned input, streaming data

initially absent from cache, each required 512-byte region fetched

exactly once for the two loads, and enough concurrency to sustain

240 GB/s of physical read traffic. Ignore other traffic. What is the

ideal useful-input throughput?

int i = blockIdx.x * blockDim.x + threadIdx.x;

float temp = A[4*i] + A[4*i+1];

A:

• From a burst of 512 bytes, we will only use every other 8 bytes due to

reading from [4*i] and [4*i+1] memory locations.

• So, we can expect 120 GB/s.

Q: Consider 2D tiled convolution with mask width of 5x5 and

output tile width of 16x16 applied to an input image of size 32x32.

Assume we use Strategy 1. The use of shared memory reduces the

number of global memory accesses. What is the reduction in

global memory accesses for thread block (0,0)?

A:

• Count # of global memory accesses:

row 0: 9+12+15*14

row 1: 12+16+20*14

other 14 rows: 15+20+25*14

• Count # of input elements to be

used in the calculation: 18*18

• Reduction: 5929/324≈18.30×

Q: For a tiled 2D convolution kernel with 30x30 output tiles and

3x3 mask (and thus 32x32 input tile), how many warps in each

thread block have control divergence? (Assume strategy 2: block

size covers input tile.)

A:

• thread block size is 32x32

• All 32x32 threads participate in data load

• What does each warp do?

• Warps 0-29: 30 threads compute, and 2 threads do not

• Warps 30-31: do not compute

• Thus, only 2 warps do not have control divergence

Q: The A40 GPU has a peak FP32 performance of 37.4 TFLOPS and

48 GB of GDDR6 memory with a memory bandwidth of 696 GB/s.

These GPUs do not support FP64. What should be the byte-to-

FLOP ratio for this GPU to fully utilize its compute capability? How

much of data reuse is needed to make full use of the compute

resources?

A:

• 696 GB/s / 37.4 TFLOPS = 0.019 B/FLOP

• Another way to look at it: for each byte loaded from global memory, we

need to use it in 53.7 floating point operations

What are the 7 steps of a cuda host function?

1. compute size

2. cudaMalloc the device buffers

3. cudaMemcpy input host→device

4. define dim3 block and grid

5. launch <<<grid, block>>>

6. cudaMemcpy output device → host

7. cudaFree the device buffers

How do you compute grid size when N may not be a multiple of TILE_WIDTH?

dim3 dimGrid, Ceiling division: (N + TILE_WIDTH - 1) / TILE_WIDTH in each dimension. Just plain N/TILE_WIDTH truncates and leaves edge elements uncovered

Suppose the input vectors have N elements and each block has 32 threads. What is/are the correct expression(s) for the grid dimension? N has type 'int'. ceil(f) returns the smallest integer greater than or equal to the given float value. floor(f) returns the largest integer less than or equal to a given number. You may assume N is positive and no overflow happens during the operation.

(a) dim3 gridDim((N-1)/32 + 1);

(b) dim3 gridDim(floor(float(N)/float(32)));

(c) dim3 gridDim((N+31)/32);

(d) dim3 gridDim(N/32 + 1);

(e) dim3 gridDim(N/32);

(f) dim3 gridDim(ceil(float(N)/float(32)));

a, c, and f

Imagine that we want to use each thread in a vector addition kernel to compute the sums for 8 adjacent elements. Which of the following expressions correctly maps the thread and block indices to idx, the array index of the first element to be processed by a thread?

idx = threadIdx.x * blockDim.x + blockIdx.x + 16;

idx = threadIdx.x * blockDim.x + blockIdx.x + 8;

idx = (blockIdx.x * blockDim.x + threadIdx.x) * 8;

idx = 16 * (blockIdx.x * blockDim.x + threadIdx.x) * 8;

idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

c) idx = (blockIdx.x * blockDim.x + threadIdx.x) * 8;

In terms of the vector length N, how many Bytes are read from global memory by the vector addition kernel?

(a) 2N Bytes

(b) 4N Bytes

(c) 8N Bytes

(d) 12N Bytes

(e) 8N² Bytes

(c) 8N Bytes

Which of the following statements are wrong?

(a) The "__global__" function qualifier indicates this function can be called from the host or from the device.

(b) By default, any traditional C program is a CUDA program that contains only host code.

(c) In general, the CPU execution and the GPU execution do not overlap.

(d) "threadIdx" and "blockIdx" are two built-in variables that are read-only.

(c) In general, the CPU execution and the GPU execution do not overlap.

Imagine that we want to use each thread in a vector addition kernel to compute the sums for 9 elements as follows.

Each thread block processes 9*blockDim.x consecutive elements that are broken into 9 contiguous sections. All threads in a block process the first section together, with each thread processing one element. The threads then move on to the next section, in which each thread again processes one element. In each section, consecutive threads process consecutive elements.

Which of the following expressions correctly computes the array index of the second element to be processed by a thread into the variable idx?

idx = blockIdx.x * blockDim.x * 9 + blockDim.x + threadIdx.x;

idx = (blockIdx.x * blockDim.x + blockDim.x + threadIdx.x) * 9;

idx = threadIdx.x * blockDim.x + blockDim.x + blockIdx.x + 9;

idx = 4 * blockIdx.x * blockDim.x * 9 + threadIdx.x;

idx = blockIdx.x * blockDim.x + threadIdx.x + blockDim.x + 9;

a) idx = blockIdx.x * blockDim.x * 9 + blockDim.x + threadIdx.x;

Which of the CUDA calls below copies an array of 160250 doubles from host array h_A to device array d_A, recording the returned value in err?

cudaError_t err = cudaMemcpy (d_A, h_A, 160250 * sizeof (double), cudaMemcpyFromHost);

cudaEvent_t err = cudaMemcpy (d_A, h_A, 160250 * sizeof (double), cudaMemcpyHostToDevice);

cudaError_t err = cudaMemcpy (d_A, h_A, 160250 * sizeof (double), cudaMemcpyHostToDevice);

cudaEvent_t err = cudaMemcpy (160250 * sizeof (double), h_A, d_A, cudaMemcpyHostToDevice);

cudaEvent_t err = cudaMemcpy (d_A, h_A, 160250, cudaMemcpyHostToDevice);

c) cudaError_t err = cudaMemcpy (d_A, h_A, 160250 * sizeof (double), cudaMemcpyHostToDevice);

Which of the following CUDA kernels could possibly have control divergence? Assuming warp size is 32. And exclude all overflow/underflow scenarios.

Has control divergence, since if statement is not a multiple of 32, it will split the warp.

Which of the following CUDA kernels could possibly have control divergence? Assuming warp size is 32. And exclude all overflow/underflow scenarios.

No control divergence, since dependent on blockIdx

Which of the following CUDA kernels could possibly have control divergence? Assuming warp size is 32. And exclude all overflow/underflow scenarios.

No control divergence

Which of the following CUDA kernels could possibly have control divergence? Assuming warp size is 32. And exclude all overflow/underflow scenarios.

Has control divergence, since it could be a multidimentional block

Which of the following CUDA kernels could possibly have control divergence? Assuming warp size is 32. And exclude all overflow/underflow scenarios.

No control divergence, since the if statement will never be executed

Each streaming multiprocessor (SM) in an NVIDIA GPU with compute capability 8.6 supports up to 1536 threads and up to 16 thread blocks. In each block, how many warps should be used to maximize the use of both of these resources?

3

Assume that the following matrix size specifications are passed to your matrix multiplication kernel in MP2:
numARows=26
numAColumns=38
numBRows=38
numBColumns=23
numCRows=26
numCColumns=23
Remember that the matrices contain floats.

How many Bytes are read from global memory by the kernel?

181792 Bytes

Assume that the following matrix size specifications are passed to your matrix multiplication kernel in MP2:
numARows=26
numAColumns=38
numBRows=38
numBColumns=23
numCRows=26
numCColumns=23
Remember that the matrices contain floats.

How many Bytes are written to global memory by the kernel?

2392 Bytes

Assume that the following matrix size specifications are passed to your matrix multiplication kernel in MP2:
numARows=26
numAColumns=38
numBRows=38
numBColumns=23
numCRows=26
numCColumns=23
Remember that the matrices contain floats.

How many floating-point operations are performed by the kernel?

45448 floating-point operations

Which of the following statements on control divergence in SIMT(Single Instruction, Multiple Threads) architecture are correct?

(a) Control divergence happens when threads within the same block are executing different branches

(b) Control divergence can lead to inefficiency because if any thread takes a different branch, all threads must wait until the longest path is completed, potentially leading to idle cycles for some threads

(c) Control divergence increases the computational throughput by utilizing more cores

(d) None of these are correct

(d) None of these are correct

You need to process a 240 x 238 image (238 pixels in the horizontal direction, 240 pixels in the vertical direction) with the kernel called DiagonalKernel(). More specifically, numCols is 238 and numRows is 240.

You decide to use a grid of 2D blocks, where each block has a dimension 16x16 threads.

How many warps will be generated during the execution of the kernel?

How many warps will have control divergence?

1800

232

A kernel is launched with 224 thread blocks, each consisting of 527 threads. The kernel contains a shared memory variable, X. How many versions of X are created during execution of the kernel?

224 versions

The kernel code fragment below is intended to have a subset of threads do some work, synchronize with each other, and then do some more work.

// some work for all threads if (threadIdx.x < totalSize) { // do some work __sync_threads (); // do some more work } else { __sync_threads (); } // some more work for all threads

Which of the following answers best identifies the problem with the code and explains how the code should be fixed?

a) Threads in a block must execute the same static syncthreads calls--in other words, the same lines in the code--not just the same number of syncthreads calls. Two conditionals with the same test must be used, and the __syncthreads calls executed outside of the conditionals.

b) Only some of the threads in a block execute the first __syncthreads, but that function synchronizes all threads, and is thus not necessary. Remove both instances to correct the code.

c) Since only the X dimension is being used to split the thread block, synchronization amongst the threads happens automatically. Such synchronization is only needed when the Y or Z dimension of thread blocks are used. Remove both instances of __syncthreads to correct the code.

d) The threads with thread indices bigger than totalSize have no useful work to do, so the extra code outside of the conditional should simply be moved into the then block to fix the code.

e) The __syncthreads function should not be called by only some of the threads in a block. The two groups of threads should be split into distinct thread blocks (or kernel launches).

a) Threads in a block must execute the same static syncthreads calls--in other words, the same lines in the code--not just the same number of syncthreads calls. Two conditionals with the same test must be used, and the __syncthreads calls executed outside of the conditionals.

Assume that the following matrix size specifications are passed to your tiled matrix multiplication kernel in MP3:
numARows=83
numAColumns=123
numBRows=123
numBColumns=81
numCRows=83
numCColumns=81
Remember that the matrices contain floats.

Also assume that you are using 32×32 tiles.

Recall that floating-point arithmetic operations include mathematical operations such as addition and multiplication.

Consider a matrix multiplication implementation where only threads responsible for an element in the output matrix C are required to perform floating-point operations. How many floating-point operations are executed in this implementation with respect to the parameters above?

1721088

Assume that the following matrix size specifications are passed to your tiled matrix multiplication kernel in MP3:
numARows=83
numAColumns=123
numBRows=123
numBColumns=81
numCRows=83
numCColumns=81
Remember that the matrices contain floats.

Also assume that you are using 32×32 tiles.

Recall that floating-point arithmetic operations include mathematical operations such as addition and multiplication.

Now consider an implementation where all threads launched perform floating-point operations. How many floating-point operations are executed now?

  2359296

In your Lab 3 kernel implementation, what is the purpose of syncthreads()? Select all correct options.

Checkbox options

(a) Make sure the shared memory tile has been loaded in before FP operations begin.

(b) Make sure the shared memory tile has been fully utlized before being re-written by the next iteration in the for loop.

(c) Make sure the threads can visit the correct global memory position and copy the corresponding value from global memory to shared memory.

(d) Make sure the threads that operate on the halo cells can't copy the matrix value into the shared memory.

(e) None of these are correct.

a and b

Assume that the following matrix size specifications are passed to your tiled matrix multiplication kernel in MP3:
numARows=70
numAColumns=55
numBRows=55
numBColumns=61
numCRows=70
numCColumns=61
Remember that the matrices contain floats.

Also assume that you are using 4×4 tiles.

How many Bytes are read from global memory by the kernel?

How many Bytes are written to global memory by the kernel?

487960 Bytes

17080 Bytes

Consider a 2D convolution tiling strategy in which each thread loads one value into shared memory from input N. In an internal tile (not near any boundary), given an output tile of size 21 x 21 and a mask of size 7 x 7, how many warps in each thread block experience control/branch divergence? If the number of threads is not a multiple of 32, the last warp is not full. Do NOT count absent threads as a cause of control divergence.

If necessary, you may assume the threads that produce the output are in the middle area of blocks.

19

Consider a 2D convolution tiling strategy in which each thread loads one value into shared memory from input N. Given an output tile of size 32 x 32 and a mask of size 7 x 21. For any particular input element, how many thread blocks (on average) access it, in the limit as N becomes large. Do not consider accesses to the mask/filter, only input N. Please keep your answer to four significant figures.

1.930

Consider a 2D convolution tiling strategy in which each thread loads one value from the input matrix. Suppose the output tiles have size 16 x 16 and the mask has size 9 x 9.

1. Assume that the GPU supports up to 2048 threads per SM. How many thread blocks can execute simultaneously on each SM?

2. How many bytes of shared memory are used on each SM? Assume that the input is an array of floats.

3. What is the average number of times used for each value loaded to shared memory in an internal tile?

4. Consider the previous setup (same block size and mask size, but different input size and output size) where each thread loads 4 values (i.e. tiles of 2x2) into shared memory. With the number of threads in each thread block unchanged, in an internal tile, what is the new average number of times used for each value loaded into the shared memory?

1. 3 blocks per SM

2. 6912 bytes

3. 36

4. 56.25

For 1D convolution, when MASK_WIDTH = 3, which of the three tiling strategies presented in lectures is the most suitable? (Hint: think about memory)

Strategy 3, since the shared memory tile holds only the output tiles elements and halo cells are read straight from global memory

Consider a 2D convolution tiling strategy in which each thread loads one value from input N. Consider 28 x 28 output tiles and a mask with radius 2. In an internal tile (ie not near any boundary), what is the size of the corresponding input tile?

1024

Describe the 1st tiling strategy

Using an example of a 1D with an output tile of 8 and a mask width of 3, so input tile is 10 wide

Block covers the output tile, load the input tile in multiple steps.

Block size: 8 threads (the output tile size)

Shared memory: the full input tile, 10 floats (halo included)

Loading takes several steps

1. some thread load left halo

2. every thread loads 1 core element

3. some threads load the right halo

Compute: every thread computes one output, and no thread is idle

Describe the 2nd tiling strategy

Using an example of a 1D with an output tile of 8 and a mask width of 3, so input tile is 10 wide

Block size: 10 threads (the input tile size)

Shared memory: the full input tile, 10 floats.

Loading: every thread loads exactly one element in a single step

Compute: only the 8 middle threads produce outputs, the halo threads sit idle during compute, which causes divergence

Describe the 3rd tiling strategy

Using an example of a 1D with an output tile of 8 and a mask width of 3, so input tile is 10 wide

Block size: 8 threads (the output tile size)

Shared memory: only the core 8 floats (least of the three)

Loading: every thread loads one core element

Compute: Each thread checks whether a neighbor it needs is in the shared tile. If it is, it reads from shared. If it is a halo cell, it reads from global memory.

For a 2D output tile of 16 × 16 and a 5 × 5 mask, how many threads per block and how many shared memory floats does strategy 1 use?

Block size: 16×16 threads = 256 threads per block

Shared mem = full input tile = (16+5-1) x (16+5-1) = 20×20=400 floats

For a 2D output tile of 16 × 16 and a 5 × 5 mask, how many threads per block and how many shared memory floats does strategy 2 use?

Block size: full input tile = (16+5-1) x (16+5-1) = 20×20=400 threads per block

Shared mem = full input tile = (16+5-1) x (16+5-1) = 20×20=400 floats

How do I write step 1 of the cuda host function in CUDA? (plus sign kernel)

size_t bytes = N * N * sizeof(float);

How do I write step 2 of the cuda host function in CUDA? (plus sign kernel)

float *d_in, *d_out; 
cudaMalloc((void**)&d_in, bytes); cudaMalloc((void**)&d_out, bytes);

How do I write step 3 of the cuda host function in CUDA? (plus sign kernel)

cudaMemcpy(d_in, h_in, bytes, cudaMemcpyHostToDevice);

How do I write step 4 of the cuda host function in CUDA? (plus sign kernel)

dim3 blockDim(TILE_WIDTH, TILE_WIDTH);
dim3 gridDim((N + TILE_WIDTH - 1) / TILE_WIDTH,   // x -> columns
             (N + TILE_WIDTH - 1) / TILE_WIDTH);  // y -> rows

How do I write step 5 of the cuda host function in CUDA? (plus sign kernel)

plus_sign_sum<<<gridDim, blockDim>>>(d_in, d_out, N);

How do I write step 6 of the cuda host function in CUDA? (plus sign kernel)

cudaMemcpy(h_out, d_out, bytes, cudaMemcpyDeviceToHost);

How do I write step 7 of the cuda host function in CUDA? (plus sign kernel)

cudaFree(d_in); cudaFree(d_out);

What are the 10 steps for a cuda kernel function? (plus sign kernel, tiled matmul, A*B+C)

1. Abbreviate indices

2. Map this thread to its output element(s)

3. Declare shared tile(s), sized from the problem

4. Declare per-thread accumulators (registers)

5. Loop over phases (only if one tile can't cover the whole input needed)

6. LOAD: guarded global -> shared, 0 if out of bounds

7. BARRIER: everything loaded before anyone reads

8. COMPUTE: read only from shared memory

9. BARRIER: everyone done reading before the next phase overwrites

10. WRITE: guarded store to global (the only place you guard on output bounds)

How do I write step 1 of a cuda kernel function? (plus sign kernel, A+B*C kernel)

int bx = blockIdx.x, tx = threadIdx.x;  
int by = blockIdx.y, ty = threadIdx.y;  

How do I write step 2 of a cuda kernel function? (plus sign kernel)

    // 2. Map this thread to its output element(s)
    int row = by*TILE_WIDTH + ty;
    int col = bx*TILE_WIDTH + tx;

How do I write step 3 of a cuda kernel function? (plus sign kernel)

__shared__ float rowTile[TILE_WIDTH][TILE_WIDTH];
__shared__ float colTile[TILE_WIDTH][TILE_WIDTH];

How do I write step 4 of a cuda kernel function? (plus sign kernel)

float rowSum = 0.0f, colSum = 0.0f;

How do I write step 5 of a cuda kernel function? (plus sign kernel)

for (int i = 0; i < ((width + TILE_WIDTH - 1) / TILE_WIDTH)); i++) {

How do I write step 6 of a cuda kernel function? (plus sign kernel, tiled matmul, A*B+C)

// load tile into shared memory
if (row < width && tile * TILE_WIDTH + threadIdx.x < width) {
	rowTile[threadIdx.y][threadIdx.x] = in[row*width + tile*TILE_WIDTH + threadIdx.x];
} else {
	rowTile[threadIdx.y][threadIdx.x] = 0;
}
		
if (tile*TILE_WIDTH+threadIdx.y < width && col < width) {
	colTile[threadIdx.y][threadIdx.x] = in[(tile*TILE_WIDTH + threadIdx.y) * width + col];
} else {
	colTile[threadIdx.y][threadIdx.x] = 0;