7. Accelerators and Heterogeneous Computing

Introduction and Context

High-performance computing (HPC) increasingly depends on systems that combine different kinds of processors rather than relying on identical general-purpose CPU cores. A contemporary computing node may contain multicore CPUs, graphics processing units (GPUs), field-programmable gate arrays (FPGAs), tensor-processing units, and other domain-specific accelerators. Such a system is heterogeneous because its computational resources differ in architecture, programming model, memory organization, performance characteristics, and energy efficiency.

The central opportunity is specialization. CPUs are flexible and effective for control-intensive, irregular, and sequential work. GPUs contain many simpler arithmetic units and are highly effective for large collections of similar operations. FPGAs can implement deeply customized pipelines, while tensor accelerators are designed for matrix and tensor operations common in machine learning. The central difficulty is that performance does not arise merely from placing an accelerator in a system. Data must be moved to the right device, organized appropriately, processed by an efficient kernel, and returned or consumed without creating synchronization and communication bottlenecks.

For data scientists, heterogeneous computing matters because many workloads—deep learning, simulation, image processing, graph analytics, optimization, scientific statistics, and large-scale linear algebra—contain substantial parallelism. This chapter develops the architectural concepts, programming methods, performance models, and decision-making principles required to use accelerators effectively.

Learning Objectives

After studying this chapter, a student should be able to:

1. Explain the architectural differences between CPUs, GPUs, FPGAs, and domain-specific accelerators.

2. Identify data-parallel, task-parallel, and pipeline-parallel opportunities in a workload.

3. Describe GPU execution, memory hierarchies, kernels, threads, blocks, and occupancy.

4. Develop the conceptual structure of CUDA, OpenCL, SYCL, and directive-based programs.

5. Optimize memory transfers, coalescing, shared-memory use, vectorization, and kernel configuration.

6. Select an appropriate accelerator programming model for a practical workload.

7. Evaluate speedup, throughput, latency, energy efficiency, portability, and total cost.

8. Diagnose common performance and correctness problems in heterogeneous applications.

The Architecture of Specialization

A CPU typically contains a small number of sophisticated cores with large caches, advanced branch prediction, out-of-order execution, and strong support for irregular control flow. It is therefore well suited to operating-system activity, orchestration, preprocessing, serial algorithms, and workloads whose parallelism is limited.

A GPU uses a large number of arithmetic units organized for throughput. Rather than optimizing primarily for the rapid completion of one instruction stream, it executes many threads in groups. The same instruction is commonly applied to different data elements, a style associated with single instruction, multiple data (SIMD), or its GPU-oriented variant, single instruction, multiple threads (SIMT). GPUs perform particularly well when thousands of operations follow similar control paths and access memory in regular patterns.

FPGAs are reconfigurable hardware devices. An application can be represented as a circuit or pipeline, allowing operations to be customized and executed concurrently. They can provide excellent latency and energy efficiency, but development, compilation, and debugging are generally more demanding than GPU programming.

Application-specific integrated circuits and domain-specific accelerators sacrifice generality for efficiency. Tensor processing units, for example, often include matrix-multiplication hardware and reduced-precision arithmetic. Their value is greatest when applications match the operations for which they were designed.

The choice of accelerator should not be reduced to peak floating-point operations per second. Relevant criteria include:

• arithmetic intensity and memory bandwidth requirements;
• regularity of control flow and data access;
• data-transfer volume between host and device;
• latency and throughput objectives;
• precision, numerical reproducibility, and supported data types;
• software portability and developer expertise;
• acquisition, power, and operational cost.

A useful principle is that an accelerator is effective only when the work assigned to it is sufficiently large and parallel to compensate for launch, synchronization, and transfer overheads.

Forms of Parallelism in Heterogeneous Workloads

Data parallelism occurs when the same operation is independently applied to many data elements. Vector addition, matrix multiplication, convolution, and pixel transformations are common examples. Data parallelism is the most direct target for GPU kernels.

Task parallelism divides an application into different activities that may execute concurrently. For example, one device may perform feature extraction while a CPU prepares the next batch. Pipeline parallelism organizes stages so that different data items occupy different stages simultaneously. A streaming analytics system may read, filter, transform, and classify records in an overlapped pipeline.

Irregular workloads require more care. Graph traversal, sparse matrix algorithms, and dynamic programming may contain uneven workloads, indirect memory accesses, and branch divergence. These applications can still benefit from accelerators, but their performance depends on data restructuring, load balancing, specialized libraries, or hybrid CPU–GPU execution.

The decomposition of an application should distinguish three roles: the host, which controls execution; the device, which performs accelerated computation; and the data path, which determines how information moves among host memory, device memory, caches, and on-chip storage. A poor data path can eliminate the benefit of a fast kernel.

GPU Execution and Memory Hierarchy

A GPU program generally consists of host code and one or more device kernels. The host allocates memory, transfers data, launches kernels, and coordinates results. A kernel is a function executed by many device threads. Threads are grouped into blocks, and blocks form a grid. Threads within a block can cooperate through fast shared memory and synchronization barriers. Blocks are normally scheduled independently, which permits execution across multiple streaming multiprocessors.

The exact terminology differs across vendors, but the underlying ideas are broadly similar. A thread has private registers and may access global device memory. Shared memory is on-chip storage visible to threads in the same block. Constant and read-only caches may accelerate repeated access to invariant data. Global memory has high capacity but substantially greater latency than registers and shared memory.

A typical one-dimensional kernel maps a thread to an element using an expression conceptually equivalent to:

`index = block_index × block_size + thread_index`

Threads whose indices exceed the data range must avoid memory access. This boundary check is essential because the grid is often rounded up to a convenient block size.

Memory coalescing is a major GPU optimization. When adjacent threads access adjacent memory locations, hardware can combine requests into efficient transactions. Strided, random, or misaligned access may require more transactions and reduce effective bandwidth. A structure-of-arrays layout can therefore outperform an array-of-structures layout when threads process the same field across many records.

Shared memory can reduce repeated global-memory access. A block first loads a tile of data cooperatively, synchronizes, and then reuses the tile. This technique is central to optimized matrix multiplication and stencil computation. It must be used carefully: excessive shared-memory allocation can reduce the number of simultaneously resident blocks, and incorrect synchronization can produce race conditions or deadlock-like behavior.

Occupancy, Latency Hiding, and Kernel Efficiency

Occupancy is commonly described as the proportion of the hardware's maximum possible active warps or wavefronts that are resident on a multiprocessor. Higher occupancy can help hide memory latency because the processor switches to another ready group while one group waits. It is not, however, a direct measure of performance. A kernel with moderate occupancy may be faster if it uses registers efficiently, has good cache behavior, and performs substantial useful work.

Occupancy is constrained by threads per block, registers per thread, shared memory per block, and hardware limits. Increasing block size may improve scheduling efficiency until register or shared-memory pressure limits active blocks. Conversely, reducing register use through compiler settings may increase occupancy but cause spills to local memory, making the kernel slower.

Performance analysis should therefore examine achieved occupancy alongside memory throughput, instruction throughput, branch efficiency, cache behavior, and execution time. Profilers are more reliable than intuition for identifying the limiting resource.

Arithmetic intensity is the ratio of arithmetic operations to bytes transferred. The roofline model uses arithmetic intensity to compare an application's demand with the machine's compute and bandwidth ceilings. A low-intensity kernel is generally memory-bound; optimizing arithmetic instructions will have little effect unless memory movement improves. A high-intensity kernel may be compute-bound and benefit from vectorization, instruction-level parallelism, or specialized math units.

Programming Models: CUDA, OpenCL, SYCL, and Directives

CUDA is a mature platform for NVIDIA GPUs. It provides explicit kernel syntax, device-memory management, streams, events, unified-memory options, libraries, and profiling tools. CUDA offers detailed control and a rich ecosystem, but source-level portability is limited when code depends on NVIDIA-specific features.

OpenCL defines a portable framework for heterogeneous devices. A host program discovers platforms and devices, creates contexts and command queues, manages buffers, and launches kernels. Kernels can be compiled for different vendors. This portability is valuable, although development may involve more explicit setup and vendor-dependent performance tuning.

SYCL provides modern C++ abstractions for heterogeneous computing. It allows kernels to be expressed using familiar C++ constructs, queues, buffers or unified shared memory, and command dependencies. SYCL can target multiple back ends, but portability still depends on compiler maturity, library support, and whether the application uses vendor-specific extensions.

Directive-based models such as OpenMP target offloading and OpenACC allow programmers to annotate loops or regions for execution on an accelerator. Compilers generate much of the device code. These approaches are attractive for incrementally porting established scientific applications, particularly when the computational structure is already expressed in regular loops. They provide less low-level control than CUDA or explicit kernel programming, but they can reduce development effort and preserve a largely unified source base.

The programming model should match the purpose of the project. Explicit APIs are appropriate when fine-grained control and maximum device-specific performance are essential. SYCL or OpenCL may be preferred when multi-vendor portability is important. Directives are useful when an existing CPU application must be accelerated rapidly and its parallel regions are well understood.

A Kernel Development Workflow

Effective accelerator development is iterative rather than purely syntactic. First, profile the CPU application and identify a region that consumes significant time and has exploitable parallelism. Next, establish a correct baseline and define numerical tolerances. Then model data movement and decide whether the selected region is large enough to amortize transfer and launch costs.

The computational region should be expressed as a kernel or offloaded loop. Data layout, index mapping, boundary conditions, and synchronization must be designed before tuning. After correctness testing on small and adversarial inputs, benchmark the implementation with representative data sizes. Finally, use a profiler to determine whether the kernel is limited by memory bandwidth, arithmetic throughput, occupancy, divergence, synchronization, or host–device communication.

A simple conceptual CUDA-style vector addition is shown below:

```cpp

__global__ void add(const float a, const float b, float* c, int n) {

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

if (i < n) c[i] = a[i] + b[i];

}

// Launch with enough blocks to cover n elements

add<<<(n + 255) / 256, 256>>>(a_d, b_d, c_d, n);

```

The example illustrates a one-thread-per-element mapping and a boundary guard. It is not necessarily an optimized implementation: for small arrays, transfer and launch overhead may dominate; for larger arrays, alignment, allocation strategy, stream usage, and memory bandwidth become relevant.

Managing Transfers and Overlapping Work

In discrete accelerator systems, host memory and device memory are physically separate. Copying data across the interconnect can be costly. A basic optimization is to transfer data in large batches rather than many small fragments. Another is to retain intermediate data on the device so that several kernels can execute without repeated round trips.

Pinned, or page-locked, host memory can support faster asynchronous transfers, although excessive use reduces flexibility for the operating system. Streams allow independent sequences of operations to overlap when hardware and dependencies permit. A pipeline may copy batch n+1 while the device computes batch n and the host processes results from batch n−1. Double or triple buffering can support this pattern.

Unified or shared memory simplifies programming by presenting a common address space, but it does not eliminate the cost of data movement. Page migration, access faults, and poor locality may result in unpredictable performance. Explicit prefetching, access advice, or deliberate data placement may be necessary for demanding applications.

Synchronization should be minimized but not removed where correctness requires it. Excessive device-wide synchronization can serialize work. At the same time, eliminating a required barrier can cause a race in which one thread reads data before another thread has written it.

Vectorization and Numerical Considerations

Vectorization applies one instruction to multiple data values. CPUs may use SIMD instruction sets such as AVX, while GPUs use groups of threads and specialized matrix units. Vectorization is most effective when data are contiguous, loop iterations are independent, and operations have compatible types.

Padding arrays, aligning allocations, simplifying loop dependencies, and using compiler-supported vector types can improve vector execution. However, vectorization may be hindered by branches, pointer aliasing, reductions, irregular indexing, and function calls. Compiler reports and generated-code inspection are useful when performance matters.

Accelerators often support lower-precision formats, including half precision or specialized numerical formats. Lower precision can increase throughput and reduce memory traffic, especially in machine learning, but it can also increase rounding error, overflow, underflow, and instability. Mixed-precision algorithms commonly perform multiplication in lower precision while accumulating or correcting results in higher precision. Numerical validation must be part of optimization rather than an afterthought.

Floating-point operations are not generally associative. A parallel reduction may produce a different order of additions from a CPU implementation and therefore a slightly different result. Applications involving financial calculations, scientific reproducibility, or regulatory reporting should define acceptable error bounds and, where necessary, use deterministic reduction strategies.

Heterogeneous Application Patterns

In machine learning, CPUs commonly handle input decoding, data preparation, scheduling, and communication, while GPUs or tensor accelerators execute matrix multiplication, convolution, and attention operations. Performance depends not only on kernel speed but also on batch size, data-loader efficiency, memory capacity, and movement between training stages.

In scientific computing, a CPU may manage sparse control flow and boundary conditions while a GPU performs dense linear algebra or stencil updates. In data analytics, filtering and aggregation may be fused into device kernels to reduce intermediate transfers. Image and video systems use pipelines in which decoding, transformation, inference, and encoding are assigned to the devices best suited to each stage.

Financial risk analysis, molecular simulation, computational fluid dynamics, genomics, and recommendation systems all illustrate the same broader principle: performance is an end-to-end property of the workflow. A highly optimized kernel cannot compensate for an inefficient data pipeline or an unsuitable partition of responsibilities.

Illustrative Case Study: Accelerating a Batch Feature Pipeline

Background

Consider an illustrative data-science team processing millions of customer records. Each record contains numerical attributes used to generate standardized features for a predictive model. The original CPU implementation performs normalization, nonlinear transformation, and a small reduction for every record.

Situation and Problem

The pipeline meets correctness requirements but occupies most of the batch-processing window. Initial inspection shows that records are independent, operations are mostly identical, and the input is stored in a row-oriented structure. However, the application repeatedly copies intermediate arrays and invokes several short kernels.

Analysis

The workload has substantial data parallelism and moderate arithmetic intensity. The row-oriented layout causes threads processing the same feature to access memory with a stride, reducing coalescing. The short kernels also create launch overhead. Profiling indicates that memory transactions and host–device transfers, rather than arithmetic operations, are the main constraints.

Approach

The team converts the input to a structure-of-arrays representation and transfers a complete batch to the GPU. Normalization and nonlinear transformation are fused into one kernel, while the reduction uses a hierarchical block-level method. Device memory is reused across batches, and two streams overlap transfer of the next batch with computation on the current batch. A CPU implementation remains available for small batches and for validation.

The team tests several block sizes, measures achieved bandwidth and occupancy, and compares results within a predefined numerical tolerance. They also monitor whether asynchronous operations actually overlap rather than assuming that stream use guarantees concurrency.

Outcome

The optimized design improves end-to-end throughput because it reduces transfers, improves memory access, and removes intermediate kernel launches. The improvement is smaller for small batches, where fixed overhead dominates. The hybrid implementation therefore selects the CPU path below a measured size threshold.

Learning Points

The case demonstrates that acceleration requires workload analysis, not simply code translation. Data layout, kernel fusion, batching, transfer overlap, numerical validation, and an informed CPU fallback all influence the result. It also illustrates why kernel-level speedup and application-level speedup must be reported separately.

Challenges, Risks, and Common Mistakes

A frequent mistake is to offload a loop without measuring its parallelism or transfer cost. Another is to maximize occupancy mechanically while ignoring memory behavior. Uncoalesced accesses, excessive branching, bank conflicts in shared memory, register spills, and unnecessary synchronization are common sources of lost performance.

Correctness risks include out-of-bounds accesses, data races, incorrect reduction logic, uninitialized device memory, and assumptions about thread ordering. Debugging is more difficult when errors are asynchronous; explicit error checks after launches and synchronization points are essential during development.

Portability is also a technical risk. Code that uses a vendor-specific library, precision mode, or synchronization primitive may not run on another accelerator. Portability should be evaluated at the level of performance as well as correctness: a program may compile on several devices but achieve useful throughput on only one.

Operational risks include thermal limits, power constraints, driver incompatibility, memory exhaustion, and scarce accelerator availability in shared clusters. Production deployments need monitoring, resource quotas, reproducible environments, and clear fallback behavior.

Governance, Professional Practice, and Sustainability

Accelerator-based data processing can reduce energy per computation, but greater efficiency may encourage larger workloads and erase absolute savings. Energy, cooling, embodied hardware impacts, and utilization should be considered in procurement and architectural decisions.

Professional practice requires transparent benchmarking. Reports should state data size, precision, hardware, software versions, compiler flags, warm-up policy, transfer inclusion, and statistical treatment of repeated measurements. Claims based only on peak theoretical throughput are misleading.

Where accelerated analytics processes personal or sensitive information, the deployment must preserve access controls, encryption, retention policies, and applicable data-protection requirements. Debugging data copied to device memory, profiler traces, and crash dumps should be treated as potential information exposures. Reproducibility also matters: nondeterministic execution and changing driver versions can complicate audit and scientific review.

Emerging Directions

Several developments are reshaping heterogeneous computing. Unified programming environments increasingly seek one source model across CPUs, GPUs, and specialized accelerators. Compilers are improving automatic kernel generation, loop transformation, graph optimization, and heterogeneous scheduling. Machine-learning systems increasingly use mixed precision, sparsity, operator fusion, and tensor-specific hardware.

Near-memory and processing-in-memory approaches aim to reduce the cost of moving data. Chiplet-based designs integrate specialized components within a package, while high-bandwidth memory increases the bandwidth available to suitable workloads. Distributed heterogeneous systems also combine accelerators across nodes, making collective communication and topology awareness important.

These developments do not remove the need for algorithmic understanding. Automatic tools can optimize a well-structured computation, but they cannot reliably compensate for unnecessary data movement, poor partitioning, or an algorithm with insufficient parallelism.

Best-Practice Framework for Accelerator Projects

A disciplined project should begin with a measured baseline and a clearly defined objective, such as latency, throughput, energy per record, or cost per training run. The team should select kernels using profiling evidence, choose a programming model consistent with portability and control requirements, and design data placement before writing detailed device code.

Correctness tests, numerical tolerances, race detection, and boundary-case validation should precede aggressive optimization. Tuning should proceed from high-impact changes—batching, data reuse, layout, algorithm choice, and transfer reduction—to lower-level parameters such as block size and register limits. Each change should be benchmarked on representative workloads.

Finally, the optimized implementation should be documented, monitored in production, and periodically reevaluated as hardware, compilers, and data distributions change. The best accelerator solution is not necessarily the fastest isolated kernel; it is the most effective, maintainable, reliable, and economically justified end-to-end system.

Chapter Summary

Heterogeneous computing combines CPUs with GPUs, FPGAs, tensor units, and other specialized processors. CPUs provide flexibility, while accelerators exploit data, task, or pipeline parallelism through specialized execution resources. GPU performance depends on kernel design, memory coalescing, shared-memory reuse, occupancy, arithmetic intensity, vectorization, and synchronization.

CUDA, OpenCL, SYCL, and directive-based approaches provide different balances of control, portability, and development effort. Successful applications minimize transfers, reuse device data, overlap communication with computation, and validate numerical correctness. Performance must be assessed end to end using meaningful metrics rather than peak hardware specifications. Professional deployment also requires attention to portability, security, reproducibility, energy, governance, and maintainability.

Key Terms

Accelerator: A processor or hardware component designed to execute particular computations more efficiently than a general-purpose CPU.

Heterogeneous computing: Computing that uses different processor architectures within one application or system.

Kernel: A function launched for execution by many accelerator threads or work-items.

SIMD: Single instruction, multiple data execution in which one instruction operates on several data values.

SIMT: Single instruction, multiple threads execution, commonly used to describe GPU thread organization.

Host: The CPU and associated system software that controls accelerator execution.

Device: The accelerator on which offloaded kernels execute.

Thread block: A cooperating group of GPU threads that can share on-chip memory and synchronize.

Warp or wavefront: A hardware execution group of threads processed together.

Occupancy: The proportion of available hardware execution capacity occupied by active thread groups.

Memory coalescing: Combining neighboring thread memory accesses into efficient memory transactions.

Shared memory: Fast, programmer-managed on-chip memory available to threads in a block.

Arithmetic intensity: The number of arithmetic operations performed per byte of data moved.

Roofline model: A performance model relating arithmetic intensity to compute and memory-bandwidth ceilings.

Kernel fusion: Combining multiple computational stages into one kernel to reduce intermediate storage and launch overhead.

Pinned memory: Host memory locked in physical memory to support efficient data transfer.

Stream: An ordered sequence of accelerator operations that may execute concurrently with other streams.

Vectorization: Executing one operation on multiple data elements using vector or parallel hardware.

Directive-based offloading: Using compiler directives to identify regions for accelerator execution.

Mixed precision: Using more than one numerical precision within an algorithm to improve performance while controlling error.

Review Questions

Short Answer Questions

1. Why are GPUs generally more suitable for regular data-parallel workloads than for highly irregular control flow?

2. Define memory coalescing and explain its effect on GPU performance.

3. What constraints influence occupancy?

4. Why can unified memory still incur performance costs?

5. Distinguish between a kernel and the host code that launches it.

Long Answer Questions

1. Compare CPUs, GPUs, FPGAs, and domain-specific accelerators in terms of architecture, strengths, limitations, and suitable workloads.

2. Explain the GPU memory hierarchy and discuss how global memory, shared memory, registers, and caches affect kernel performance.

3. Compare CUDA, OpenCL, SYCL, and directive-based programming models with respect to control, portability, and development effort.

4. Explain the roofline model and show how arithmetic intensity guides optimization decisions.

5. Describe a complete workflow for profiling, implementing, validating, and tuning an accelerator-based application.

Analytical / Application Questions

1. A GPU kernel has high occupancy but low throughput. Accesses are irregular and profiler results show poor memory efficiency. Diagnose the likely problem and propose two changes.

2. An application reports a tenfold kernel speedup but only a 1.4-fold end-to-end speedup. Analyze possible causes and recommend measurements or design changes.

3. A research group needs an application to run on CPUs, GPUs from multiple vendors, and future accelerators, while minimizing changes to an existing loop-based codebase. Which programming approach would you initially evaluate, and what trade-offs would you explain?

22. Multiple Choice Questions with Answers

Question 1

Which workload is most naturally suited to GPU acceleration?

A. A sequential parser with unpredictable branches

B. The same arithmetic operation applied independently to millions of elements

C. A small configuration routine executed once

D. An algorithm whose steps depend on the immediately preceding result

Correct Answer: B. The same arithmetic operation applied independently to millions of elements

Question 2

What is the primary benefit of coalesced global-memory access?

A. It eliminates the need for synchronization

B. It combines neighboring thread requests into more efficient memory transactions

C. It guarantees maximum occupancy

D. It converts floating-point operations to integer operations

Correct Answer: B. It combines neighboring thread requests into more efficient memory transactions

Question 3

Which factor can directly limit the number of active blocks on a multiprocessor?

A. The number of comments in the kernel

B. Register or shared-memory use per block

C. The file name of the source program

D. The number of CPU cores in the host

Correct Answer: B. Register or shared-memory use per block

Question 4

A kernel has low arithmetic intensity and is memory-bound. Which optimization is most likely to help first?

A. Increasing the number of transcendental operations

B. Reducing unnecessary data movement and improving access locality

C. Adding more conditional branches

D. Increasing numerical precision without changing the algorithm

Correct Answer: B. Reducing unnecessary data movement and improving access locality

Question 5

What is a principal purpose of using streams in accelerator programming?

A. To guarantee that every operation executes serially

B. To permit independent transfers and computations to overlap when dependencies allow

C. To increase the size of device registers

D. To replace all error checking

Correct Answer: B. To permit independent transfers and computations to overlap when dependencies allow

Question 6

Which programming approach generally provides the greatest vendor-specific low-level control on supported NVIDIA GPUs?

A. CUDA

B. OpenCL only

C. A spreadsheet formula

D. A database query language

Correct Answer: A. CUDA

Question 7

Why might kernel fusion improve end-to-end performance?

A. It necessarily increases the number of host-device transfers

B. It can reduce intermediate memory traffic and kernel-launch overhead

C. It guarantees that all branches disappear

D. It removes the need for numerical validation

Correct Answer: B. It can reduce intermediate memory traffic and kernel-launch overhead

Question 8

A parallel reduction produces a slightly different floating-point result from the CPU version. What is the most likely explanation?

A. Floating-point addition is not generally associative, so the operation order changed

B. The GPU cannot perform addition

C. Memory coalescing changes the mathematical value by definition

D. Occupancy always causes arithmetic overflow

Correct Answer: A. Floating-point addition is not generally associative, so the operation order changed

Question 9

Which consideration is most important when selecting between an explicit kernel API and a directive-based approach?

A. Whether the source code has a short file name

B. The required balance of low-level control, portability, and development effort

C. Whether the application contains any variables

D. The number of comments required by the compiler

Correct Answer: B. The required balance of low-level control, portability, and development effort

Question 10

In the illustrative feature-processing case, why was a CPU fallback retained for small batches?

A. Small batches may not amortize accelerator transfer and launch overheads

B. GPUs cannot process numerical features

C. CPUs always have greater memory bandwidth

D. The fallback guarantees identical instruction ordering

Correct Answer: A. Small batches may not amortize accelerator transfer and launch overheads