第12回配信講義 計算科学技術特論B（2024）

CUDAやOpenACCによる GPUコンピューティング Akira Naruse, 5th July 2024

AGENDA  GPU Computing  OpenACC  CUDA 2

GPU Computing 3

GPU Computing Low latency + High throughput CPU GPU 4

アプリケーション実行 アプリケーション・コード GPU Do i=1,N CPU 計算の 重い部分 End do 逐次部分は CPU上で実行 並列部分は GPUで実行 5

GPUの構造 (NVIDIA H100) 16,896 CUDA cores /chip この全てを活用できるかが鍵 128 FP32 CUDA cores /SM 132 Stream Multiprocessor (SM) /chip 6

GPU APPLICATIONS 数千のアプリケーションがGPUに対応 www.nvidia.com/en-us/accelerated-applications 7

LEADING APPLICATIONS (HPC) https://developer.nvidia.com/hpc-application-performance 8

• https://developer.nvidia.com/hpc-application-performance

GPU Programming 1 9

1 NVIDIA Math Libraries 10

GPU Programming 2 1 11

2 Accelerated Standard Languages 12

Applications using STDPAR 13

GPU Programming 2 3 1 14

3

Incremental Optimizations
OpenMP

OpenACC

void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma omp parallel for
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma acc parallel copy(y[:n]) copyin(x[:n])
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

...
saxpy(N, 3.0, x, y);
...

...
saxpy(N, 3.0, x, y);
...
15

GPU Programming 2 3 4 1 16

Platform Specialization

4
CPU

void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

...
saxpy(N, 3.0, x, y);
...

CUDA
__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)
y[i] += a*x[i];
}

...
size_t size = sizeof(float) * N;
cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);
...

17

OpenACC 18

GPU Programming 19

OPENACC GPU CPU 簡単: 既存のコードに コンパイラへのヒントを追加 Program myscience ... serial code ... !$acc kernels do k = 1,n1 do i = 1,n2 ... parallel code ... enddo enddo !$acc end kernels ... serial code … End Program myscience 既存のC/Fortranコード 強力: 少ない労力で、コンパイラがコー ドを自動で並列化 ヒントの 追加 オープン: 複数コンパイが、様々なプロ セッサをサポート NVIDIA/AMD GPU, X86/ARM CPU 20

GPUコンピューティング アプリケーション・コード Do i=1,N CPU OpenACC GPU 計算の 重い部分 End do 逐次部分は CPU上で実行 並列部分は GPUで実行 21

SAXPY (y=a*x+y)
OpenMP

OpenACC

void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma omp parallel for
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

void saxpy(int n,
float a,
float *x,
float *restrict y)
{
#pragma acc parallel copy(y[:n]) copyin(x[:n])
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

...
saxpy(N, 3.0, x, y);
...

...
saxpy(N, 3.0, x, y);
...

 omp  acc

 データの移動
22

SAXPY (y=a*x+y, FORTRAN) OpenMP OpenACC subroutine saxpy(n, a, X, Y) real :: a, X(:), Y(:) integer :: n, i subroutine saxpy(n, a, X, Y) real :: a, Y(:), Y(:) integer :: n, i !$omp parallel do do i=1,n Y(i) = a*X(i)+Y(i) enddo !$omp end parallel do end subroutine saxpy !$acc parallel copy(Y(:)) copyin(X(:)) do i=1,n Y(i) = a*X(i)+Y(i) enddo !$acc end parallel end subroutine saxpy ... call saxpy(N, 3.0, x, y) ... ... call saxpy(N, 3.0, x, y) ...  FORTRANも同様 23

OPENACC: PARALLEL DIRECTIVE
void saxpy(int n, float a,
float *x,
float *restrict y)
{
#pragma acc parallel copy(y[:n]) copyin(x[:n])
for (int i = 0; i < n; ++i) {
y[i] += a*x[i];
}
}

並列化領域を
指示する

...
saxpy(N, 3.0, x, y);
...

24

OPENACC: KERNELS DIRECTIVE
void saxpy(int n, float a,
float *x,
float *restrict y)
{
#pragma acc kernels copy(y[:n]) copyin(x[:n])
for (int i = 0; i < n; ++i) {
y[i] += a*x[i];
}
}
...
saxpy(N, 3.0, x, y);
...

並列化領域を
指示する

Parallel: ユーザ主導
Kernels: コンパイラ主導

25

OPENACC: LOOP CLAUSE
void saxpy(int n, float a,
float *x,
float *restrict y)
{
#pragma acc kernels copy(y[:n]) copyin(x[:n])
#pragma acc loop independent
for (int i = 0; i < n; ++i) {
y[i] += a*x[i];
}
}
...
saxpy(N, 3.0, x, y);
...

ループの並列化方法を
指示する

independent: 並列化可能
seq: 逐次実行
collapse: ループ融合
gang/vector: 並列化粒度
...
26

OPENACC: DATA CLAUSE
void saxpy(int n, float a,
float *x,
float *restrict y)
{
#pragma acc kernels copy(y[:n]) copyin(x[:n])
for (int i = 0; i < n; ++i) {
y[i] += a*x[i];
}
}
...
saxpy(N, 3.0, x, y);
...

データの移動方法を
指示する

copyin: CPU  GPU
copyout: CPU  GPU
copy: (両方)
create: (メモリ確保)
...
27

OPENACC DATA DIRECTIVE
void saxpy(int n, float a,
float *x,
float *restrict y)
{
#pragma acc kernels present(x, y)
for (int i = 0; i < n; ++i) {
y[i] += a*x[i];
}
}

データ領域を
指示する

#pragma acc data copy(y[:N]) create(x[:N])
{
produce_x();
saxpy(N, 3.0, x, y);
consume_y();
}
28

簡単にコンパイル
OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
{
$ nvc –acc=gpu
–Minfo=acc –gpu=... saxpy.c
#pragma acc parallel copy(y[:n]) copyin(x[:n])
saxpy:
#pragma omp parallel for
16, Generating
for present_or_copy(y[:n])
(int i = 0; i < n; ++i)
Generatingy[i]
present_or_copyin(x[:n])
+= a*x[i];
Generating
Tesla code
}
19, Loop is parallelizable
Accelerator
kernel generated
...
19, #pragma
acc3.0,
loopx,gang,
saxpy(N,
y); vector(128) /* blockIdx.x threadIdx.x */
...

29

簡単に実行
OpenACC
void saxpy(int n, float a,
float *x,
float *restrict y)
{
$ nvc –acc=gpu
–Minfo=acc –gpu=... saxpy.c
#pragma
acc parallel copy(y[:n]) copyin(x[:n])
$ nsys nvprof ./a.out
saxpy:
#pragma omp parallel for
...16, Generating
present_or_copy(y[:n])
for
(int Calls
i = 0; i < n;
Time(%)Generating
Time present_or_copyin(x[:n])
Avg++i)
Min
Max Name
y[i]
+=
a*x[i];
62.95%Generating
3.0358ms Tesla code
2 1.5179ms 1.5172ms 1.5186ms [CUDA memcpy HtoD]
}
31.48%
1.5181ms
1 1.5181ms 1.5181ms 1.5181ms [CUDA memcpy DtoH]
19, Loop
is parallelizable
5.56%Accelerator
268.31us kernel generated
1 268.31us 268.31us 268.31us saxpy_19_gpu
...
19, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
saxpy(N, 3.0, x, y);
...
30

OPENACCが加速する科学技術計算 www.openacc.org/success-stories 31

NUMECA FINE/TURBO Commercial CFD Application CHALLENGE • Accelerate 20 year old highly optimized code on GPUs SOLUTION “ “ OpenACC enabled us to target routines for GPU acceleration without rewriting code, allowing us to maintain portability on a code that is 20-years old David Gutzwiller, Head of HPC NUMECA International • Accelerated computationally intensive routines with OpenACC RESULTS • Achieved 10x or higher speed-up on key routines • Full app speedup of up to 2x on the Oak Ridge Titan supercomputer • Total time spent on optimizing various routines with OpenACC was just five person-months www.openacc.org/success-stories 32

CLOVERLEAF Performance Portability for a Hydrodynamics Application Speedup vs 1 CPU Core CHALLENGE • Application code that runs across architectures without compromising on performance SOLUTION • Use OpenACC to port the CloverLeaf mini app to GPUs and then recompile and run the same source code on multi-core CPUs RESULTS We were extremely impressed that we can run OpenACC on a CPU with no code change and get equivalent performance to our OpenMP/MPI implementation. “ “ Wayne Gaudin and Oliver Perks Atomic Weapons Establishment, UK • Same performance as optimized OpenMP version on x86 CPU • 4x faster performance using the same code on a GPU Benchmarked Intel(R) Xeon(R) CPU E5-2690 v2 @ 3.00GHz, Accelerator: Tesla K80 www.openacc.org/success-stories 33

LSDALTON Minimal Effort Large-scale application for calculating highaccuracy molecular energies Lines of Code Modified # of Weeks Required # of Codes to Maintain <100 Lines 1 Week 1 Source Big Performance LS-DALTON CCSD(T) Module Benchmarked on Titan Supercomputer (AMD CPU vs Tesla K20X) “ Janus Juul Eriksen, PhD Fellow qLEAP Center for Theoretical Chemistry, Aarhus University “ OpenACC makes GPU computing approachable for domain scientists. Initial OpenACC implementation required only minor effort, and more importantly, no modifications of our existing CPU implementation. Speedup vs CPU 11.7x 7.9x ALANINE-1 13 ATOMS www.openacc.org/success-stories 8.9x ALANINE-2 23 ATOMS ALANINE-3 33 ATOMS 34

POWERGRID Advanced MRI Reconstruction Model CHALLENGE • Produce detailed and accurate brain images by applying computationally intensive algorithms to MRI data • Reduce reconstruction time to make diagnostic use possible SOLUTION Now that we’ve seen how easy it is to program the GPU using OpenACC and the PGI compiler, we’re looking forward to translating more of our projects Brad Sutton, Associate Professor of Bioengineering and Technical Director of the Biomedical Imaging Center University of Illinois at Urbana-Champaign “ “ • Accelerated MRI reconstruction application with OpenACC using NVIDIA GPUs RESULTS • Reduced reconstruction time for a single high-resolution MRI scan from 40 days to a couple of hours • Scaled on Blue Waters at NCSA to reconstruct 3000 images in under 24 hours www.openacc.org/success-stories 35

NEKCEM Computational Electromagnetics Application CHALLENGE • Enable NekCEM to deliver strong scaling on next-generation architectures while maintaining portability SOLUTION The most significant result from our performance studies is faster computation with less energy consumption compared with our CPU-only runs. “ “ Dr. Misun Min, Computation Scientist Argonne National Laboratory • Use OpenACC to port the entire program to GPUs RESULTS • 2.5x speedup over a highly tuned CPU-only version • GPU used only 39 percent of the energy needed by 16 CPUs to do the same computation in the same amount of time www.openacc.org/success-stories 36

MAESTRO & CASTRO Astrophysics 3D Simulation 2 weeks to learn OpenACC CHALLENGE • Pursuing strong scaling and portability to run code on GPU-powered supercomputers SOLUTION 2 weeks to modify code On the reactions side, accelerated calculations allow us to model larger networks of nuclear reactions for similar computational costs as the simple networks we model now “ “ • OpenACC compiler on OLCF’s Titan supercomputer RESULTS • 4.4x faster reactions than on a multi-core computer with 16 cores • Accelerated calculations allow modeling larger networks of nuclear reactions for similar computational costs as simple networks Adam Jacobs, PhD candidate in the Department of Physics and Astronomy at Stony Brook University www.openacc.org/success-stories 37

INCOMP3D 3D Fully Implicit CFD Solver OpenACC is a highly effective tool for programming fully implicit CFD solvers on GPU to achieve true 4X speedup “ “ Lixiang Luo, Researcher Aerospace Engineering Computational Fluid Dynamics Laboratory North Carolina State University (NCSU) CHALLENGE • Accelerate a complex implicit CFD solver on GPU SOLUTION • Used OpenACC to run solver on GPU with minimal code changes. RESULTS • Achieved up to 4X speedups on Tesla GPU over parallel MPI implementation on x86 multicore processor • Structured approach to parallelism provided by OpenACC allowed better algorithm design without worrying about GPU architecture * CPU Speedup on 6 cores of Xeon E5645, with additional cores performance reduces due to partitioning and MPI overheads www.openacc.org/success-stories 38

CUDA 39

GPU Programming 40

CUDAプログラミング  プログラミングモデル  アーキテクチャ  性能Tips 41

GPUコンピューティング CPU CPU Memory GPU Memory GPU PCI  高スループット指向のプロセッサ  分離されたメモリ空間 42

GPUプログラム
CPU
void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

...
saxpy(N, 3.0, x, y);
...

GPU
__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)
y[i] += a*x[i];
}

...
size_t size = sizeof(float) * N;
cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);
cudaDeviceSynchronize();
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);
...
43

GPU実行の基本的な流れ CPU GPU  GPUは、CPUからの制御 で動作 入力データ 転送 GPUカーネル 投入  GPU上で 演算 同期   入力データ: CPUからGPUに 転送 (H2D) GPUカーネル: CPUから投入 出力データ: GPUからCPUに 転送 (D2H) 出力データ 転送 44

GPUプログラム
GPU

CPU
void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)
y[i] += a*x[i];
}

入力データ転送
カーネル起動
...
saxpy(N, 3.0, x, y);
...

同期

出力データ転送

...
size_t size = sizeof(float) * N;
cudaMemcpy(d_x, x, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, y, size, cudaMemcpyHostToDevice);
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);
cudaDeviceSynchronize();
cudaMemcpy(y, d_y, size, cudaMemcpyDeviceToHost);
...
45

GPUプログラム (Unified Memory)
GPU

CPU
void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)
y[i] += a*x[i];
}

...

カーネル起動
...
saxpy(N, 3.0, x, y);
...

同期

saxpy<<< N/128, 128 >>>(N, 3.0, x, y);
cudaDeviceSynchronize();
...
46

GPUプログラム (Unified Memory)
GPU

CPU

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx;
if (i < n)
y[i] += a*x[i];
}

void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];
}

プリフェッチ

...
saxpy(N, 3.0, x, y);
...

...
int dev_id = 0;
size_t size = sizeof(float) * N;
cudaMemPrefetchAsync(x, size, dev_id);
cudaMemPrefetchAsync(y, size, dev_id);
saxpy<<< N/128, 128 >>>(N, 3.0, x, y);
cudaDeviceSynchronize();
...
47

GPUカーネル
GPU

CPU

}

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx.x;
if (i < n)
y[i] += a*x[i];
Global スレッドID
}

...
saxpy(N, 3.0, x, y);
...

...
saxpy<<< N/128, 128 >>>(N, 3.0, d_x, d_y);
...

void saxpy(int n, float a,
float *x, float *y)
{
for (int i = 0; i < n; ++i)
y[i] += a*x[i];

 GPUカーネル: 1つのGPUスレッドの処理内容を記述
 基本: 1つのGPUスレッドが、1つの配列要素を担当
48

Execution Configuration
(ブロック数とブロックサイズ)
スレッドID

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx.x;
if (i < n)
y[i] += a*x[i];
}
ブロックサイズ

ブロックID

...
saxpy<<< (N+127)/128, 128 >>>(N, 3.0, d_x, d_y);
...

ブロック数

ブロックサイズ

ブロック数 x ブロックサイズ ≧ 配列要素数

49

スレッド階層 y[i] = a*x[i] + y[i] (スレッド、ブロック、グリッド) x[] 0 127 128 255 256 383 384 y[] 0 127 128 255 256 383 384 0 127 128 0 255 127 256 0 383 127 384 0 スレッド (global) (local) ブロック0 ブロック2 ブロック1 グリッド  ブロックサイズ(スレッド数/ブロック)は、カーネル毎に設定可能  推奨: 128 or 256 スレッド 50

Execution Configuration
(ブロック数とブロックサイズ)
スレッドID

__global__ void saxpy(int n, float a,
float *x, float *y)
{
int i = threadIdx.x + blodkDim.x * blockIdx.x;
if (i < n)
y[i] += a*x[i];
}
ブロックサイズ

ブロックID

...
(N+255)/256,
256
saxpy<<< (N+
(N+127)/128,
63)/ 64, 128
64 >>>(N, 3.0, d_x, d_y);
...

ブロック数

ブロックサイズ

ブロック数 x ブロックサイズ ≧ 配列要素数

51

2D配列のGPUカーネル例
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
Globalスレッド ID (x)
int i = threadIdx.x + blodkDim.x * blockIdx.x;
int j = threadIdx.y + blodkDim.y * blockIdy.y;
if ( i < N && j < N )
GlobalスレッドID (y)
C[i][i] = A[i][j] + B[i][j];
}

ブロックサイズ (x,y)
...
4
32,
8 );
dim3 sizeBlock( 64,
16, 16
dim3 numBlocks( N/sizeBlock.x, N/sizeBlock.y );
MatAdd<<< numBlocks, sizeBlock >>>(A, B, C);
...

ブロック数 (x,y)

 ブロックサイズ(ブロック形状)は、1D～3Dで表現可能
52

ブロック・マッピング、スレッド・マッピング dim3 sizeBlock(16,16) (0,0) (0,0) (1,0) (0,0) dim3 sizeBlock(32,8) (0,2) (1,1) (1,2) (1,0) (0,1) (1,1) (0,2) (1,2) (0,3) (1,3) (0,4) (1,4) (2,0) (15,15) (0,1) (0,0) (31,7) (2,1) (2,2) ブロックID(blockIdx) スレッドID(threadIdx) 53

CUDAプログラミング  プログラミングモデル  アーキテクチャ  性能Tips 54

GPUアーキテクチャ概要 Streaming Multiprocessor  「並列」プロセッサ、 H100:132 L2 cache (60 MB)  全SMからアクセス可能な キャッシュ NVIDIA H100 DRAM (HBM3)  全SMからアクセス可能な メモリ 55

Stream Multi-Processor 演算ユニット  INT32: 64個  FP32: 128個  FP64: 64個 その他ユニット  LD/ST, SFU, etc. レジスタ(32bit): 64K個 (256KB) 共有メモリ/L1キャッシュ: 256KB NVIDIA H100 56

GPUカーネル実行の流れ  CPUが、GPUに、グリッドを投入  具体的な投入先は、Giga Thread Engine グリッド ブロック スレッド スレッド ブロック 57

GPUカーネル実行の流れ  Giga Thread Engine(GTE)が、SMに、ブロックを投入  GTEは、ブロックスケジューラ  グリッドをブロックに分解して、ブロックを、空いているSMに割当てる グリッド ブロック スレッド スレッド ブロック 58

ブロックをSMに割り当て  各ブロックは、互いに独立に実行  ブロック間では同期しない、実行順序の保証なし  1つのブロックは複数SMにまたがらない  1つのSMに、複数ブロックが割当てられることはある グリッド ブロック ブロック ブロック ブロック 59

GPUカーネル実行の流れ  SM内のスケジューラが、スレッドをCUDAコアに投入 グリッド ブロック スレッド スレッド ブロック 60

GPUカーネル実行の流れ  SM内のスケジューラが、ワープをCUDAコアに投入  ワープ: 32スレッドの塊  ブロックをワープに分割、実行可能なワープを、空CUDAコアに割当てる グリッド ブロック ワープ ワープ ブロック 61

ワープのCUDAコアへの割り当て  ワープ内の32スレッドは、同じ命令を同期して実行 SIMT (Single Instruction Multiple Threads)  各ワープは、互いに独立して実行  同じブロック内のワープは、明示的に同期可能(__syncthreads()) グリッド ブロック ワープ ワープ ブロック ワープ ワープ 62

GPUアーキの変化を問題としないプログラミングモデル Pascal, CC6.0 64 cores /SM Kepler, CC3.5 192 cores /SM Maxwell, CC5.0 128 cores /SM 63

CUDAプログラミング  プログラミングモデル  アーキテクチャ  性能Tips 64

デバイスメモリへのアクセスは、まとめて  コアレス・アクセス  32スレッド(ワープ)のロード/ストアをまとめて、メモリトランザクションを発行  トランザクションサイズ: 32B, 64B or 128B  トランザクション数は、少ないほど良い 配列 128B x1 Best Good Bad 128B境界 0 1 0 128B境界 28 29 30 31 0 1 2 3 14 15 16 17 1 30 31 64B x2 32B x32 2 65

リソース使用率 (Occupancy) SMの利用効率を上げる SMに割当て可能なスレッド数を、 上限に近づける  レジスタ使用量(/スレッド) CUDAコア数: 64 最大スレッド数: 2048 最大ブロック数: 32 共有メモリ: 160KB レジスタ数(32-bit): 64K個  できる限り減らす リソース量/SM (A100)  DP(64bit)は、2レジスタ消費  レジスタ割り当て単位は8個  レジスタ使用量と、割当て可能なスレッド数の関係  32レジスタ: 2048(100%), 64レジスタ: 1024(50%)  128レジスタ:512(25%), 256レジスタ: 256(12.5%) 66

リソース使用率 (Occupancy) SMの利用効率を上げる SMに割当て可能なスレッド数を、 上限に近づける  スレッド数(/ブロック)  64以上にする  64未満だと最大ブロック数がネックになる CUDAコア数: 64 最大スレッド数: 2048 最大ブロック数: 32 共有メモリ: 160KB レジスタ数(32-bit): 64K個 リソース量/SM (A100)  共有メモリ使用量(/ブロック)  できる限り減らす  共有メモリ使用量と、割当て可能なブロック数の関係  32KB:2ブロック, 16KB:4ブロック, 8KB:8ブロック 67

リソース使用率 (Occupancy) 空き時間を埋める  CUDAストリーム (≒キュー) (*) 操作: GPUカーネル, データ転送  同じCUDAストリームに投入した操作: 投入順に実行  別のCUDAストリームに投入した操作: 非同期に実行 (オーバラップ実行) [CUDAストリームの効果例] GPUカーネルとデータ転送が オーバラップして 同時に実行されている 68

分岐を減らす  ワープ内のスレッドが別パスを選択すると遅くなる  ワープ内のスレッドは、命令を共有 (SIMT)  ワープ内のスレッドが選んだ全パスの命令を実行  あるパスの命令を実行中、そのパスにいないスレッドはinactive状態  Path divergenceを減らす  できる限り、同ワープ内のスレッドは 同じパスを選択させる A[id] > 0 true 0 処理 X false 1 2 3 処理 Y 69

まとめ  GPU Computing  OpenACC  CUDA 70