ハンズオン
Lab 2

ハンズオンについて• この講義は以下の教材をもとにしています
https://github.com/openhackathonsorg/gpubootcamp/blob/58e1329572bebc508ba7489a9f9415d7e0592ab8/hpc/nways/nways_labs/nways_MD/English/Python/jupyter_notebook/cupy/cupy_guide.ipynb• Lab 1

• CuPy による GPU コンピューティングの講義• ハンズオン• Exercise 1-4 を実施• Lab 2
• ハンズオンで使うコードの概要説明• Nsight Systems の概要説明• ハンズオン• Lab Task を実施

The Serial Codedcdreadhead and dcdreadframe• 以下の Cell 1 を参照• https://github.com/openhackathonsorg/gpubootcamp/blob/58e1329572bebc508ba7489a9f9415d7e0592ab8/hpc/nways/nways_labs/nways_M
D/English/Python/jupyter_notebook/cupy/serial_RDF.ipynb• または、/work/EDU5/ユーザ名/hands-on/Lab2/nways_serial.py を参照• dcdreadhead : DCDFile からフレーム数、原⼦数を読み込み• dcdreadframe : MDAnalysis ライブラリを使って、10 (6720 atoms / frame) フレームを読み込み• 上記 2 つの関数はホストで実⾏

The Serial Codepair_gpu

The Serial Codepair_gpu

NSIGHT SYSTEMS
System profiler
Key Features:
System-wide application algorithm tuning•••

Locate optimization opportunities•••

Visualize millions of events on a very fast GUI timeline
Or gaps of unused CPU and GPU time

Balance your workload across multiple CPUs and GPUs••

Multi-process tree support

CPU algorithms, utilization and thread state
GPU streams, kernels, memory transfers, etc

Command Line, Standalone, IDE Integration

OS: Linux (x86, Power, Arm SBSA, Tegra), Windows, MacOSX(host)
GPUs: Pascal+
Docs/product: https://developer.nvidia.com/nsight-systems

CPUutilization
Processes &threads
OS runtime
APIs
NVTX
annotations

CUDA &cuBLAS APIs

GPU CUDA
Kernels & memorytransfers

NVTX projected onGPU CUDA streams

CPU IP &backtracesample dataNsight Systems 101
Quick start• Nsight Systems CLI によるプロファイリング
$ nsys profile [options] <application> [application-arguments]• -t <parameters>: トレースする API を指定。デフォルトは、cuda, opengl nvtx, osrt

• --stats <true|false>: true でプログラム実⾏時の統計情報を標準出⼒に表⽰• -o <filename>: 出⼒ファイル名を指定• --force-overwrite <true|false>: true で出⼒ファイルの上書きを許可。デフォルトは falseなど... 詳細は、nsys --help or nsys [specific command] --help で確認可能• <filename>.nsys-rep が出⼒される• ローカルに <filename>.nsys-rep を転送し、Nsight Systems UI で可視化

Nsight Systems user guide:
https://docs.nvidia.com/nsight-systems/UserGuide/index.html

Nsight Systems 101
Example• 実⾏コマンド例$ nsys profile--stats true --force-overwrite true -o my_report python solution.py
NVTX Range Statistics:• 統計情報が標準出⼒に出⼒されるTime (%)-------87.9
9.7
2.4

Total Time (ns)--------------1,015,559,769
112,401,039
27,280,924

Instances--------111

Avg (ns)--------------1,015,559,769.0
112,401,039.0
27,280,924.0

Med (ns)--------------1,015,559,769.0
112,401,039.0
27,280,924.0

Min (ns)------------1,015,559,769
112,401,039
27,280,924

Max (ns)------------1,015,559,769
112,401,039
27,280,924

StdDev (ns)----------0.0
0.0
0.0

Style
------PushPop
PushPop
PushPopRange------------------------Read_File
CuPy_Pair_gpu_Circulation
Entropy_Calculation...
[5/8] Executing 'cudaapisum' stats report
CUDA API Statistics:
Time (%)-------94.3
5.10.40.00.0...

Total Time (ns)--------------219,106,094
11,788,845
958,943
106,326
67,470Num Calls--------4
12
11Avg (ns)-----------54,776,523.5
11,788,845.0
479,471.5
106,326.0
67,470.0

Med (ns)-----------8,807.5
11,788,845.0
479,471.5
106,326.0
67,470.0

Min (ns)---------6,97811,788,845
6,294
106,326
67,470Max (ns)----------219,081,501
11,788,845
952,649
106,326
67,470

StdDev (ns)------------109,536,651.7
0.0
669,174.0
0.00.0Med (ns)-----------11,817,865.0

Min (ns)---------11,817,865

Max (ns)---------11,817,865

StdDev (ns)----------0.0

Name---------------------------cudaMalloc
cudaDeviceSynchronize
cudaHostAlloc
cuModuleLoadData
cuModuleUnload[6/8] Executing 'gpukernsum' stats report
CUDA Kernel Statistics:
Time (%)-------100.0

Total Time (ns)--------------11,817,865

Instances--------1

Avg (ns)-----------11,817,865.0

Name------------cupy_pair_gpu[7/8] Executing 'gpumemtimesum' stats report
CUDA Memory Operation Statistics (by time):
Time (%)-------98.1
1.9

Total Time (ns)--------------143,295
2,784Count----41

Avg (ns)-------35,823.8
2,784.0

Med (ns)-------46,207.5
2,784.0

Min (ns)-------3,232
2,784

Max (ns)-------47,648
2,784

StdDev (ns)----------21,738.5
0.0Operation-----------------[CUDA memcpy HtoD][CUDA memcpy DtoH][8/8] Executing 'gpumemsizesum' stats report
CUDA Memory Operation Statistics (by size):
Total (MB)---------1.629
0.016

Count----41Avg (MB)-------0.407
0.016

Med (MB)-------0.538
0.016

Min (MB)-------0.016
0.016

Max (MB)-------0.538
0.016

StdDev (MB)----------0.261
0.000Operation-----------------[CUDA memcpy HtoD][CUDA memcpy DtoH]

Nsight Systems 101
Example• *.nsys-rep を⼿元の端末に転送• ⼿元の端末で Nsight Systems を起動し、File->Open で *nsys-rep ファイルを選択
マウスオーバーすることで、ブロックサイズ、ブロック数などを確認可能

GPU カーネル
データ転送ハンズオンについて• この講義は以下の教材をもとにしています
https://github.com/openhackathonsorg/gpubootcamp/blob/58e1329572bebc508ba7489a9f9415d7e0592ab8/hpc/nways/nways_labs/nways_MD/English/Python/jupyter_notebook/cupy/cupy_guide.ipynb• Lab 1

• CuPy による GPU コンピューティングの講義• ハンズオン• Exercise 1-4 を実施• Lab 2
• ハンズオンで使うコードの概要説明• Nsight Systems の概要説明• ハンズオン• Lab Task を実施

ハンズオン• Lab Task

• nways_serial.py を CPU 実⾏し時間計測 (標準出⼒の kernel compute time に注⽬)
• nways_serial.py を CuPy Raw Kernel を使って GPU 対応、実⾏、時間計測• 正しく GPU 化できたかを確認しながら進める (標準出⼒の s2 と s2bound が以下の値であれば OK)
s2 value is -2.43191
s2bond value is -3.87014• 必要に応じて、前記講義資料で復習• (時間があれば、Nsight Systems を使ったプロファイリングにもチャレンジしてみましょう)回答例
raw_kernel = cp.RawKernel(r'''extern "C" __global__
void cupy_pair_gpu(const double* d_x, const double* d_y, const double* d_z,
unsigned long long int *d_g2, int numatm, int nconf,
double xbox, double ybox, double zbox, int d_bin){double r, cut, dx, dy, dz;
int ig2;
double box;
box = min(xbox, ybox);
box = min(box, zbox);
double del = box / (2.0 * d_bin);cut = box * 0.5;
int id1 = blockIdx.y * blockDim.y + threadIdx.y;
int id2 = blockIdx.x * blockDim.x + threadIdx.x;
if (id1 >= numatm || id2 >= numatm) return;if (id1 > id2) return;for (int frame = 0; frame < nconf; ++frame) {dx = d_x[frame * numatm + id1] - d_x[frame * numatm + id2];
dy = d_y[frame * numatm + id1] - d_y[frame * numatm + id2];
dz = d_z[frame * numatm + id1] - d_z[frame * numatm + id2];
dx = dx - xbox * (round(dx / xbox));dy = dy - ybox * (round(dy / ybox));dz = dz - zbox * (round(dz / zbox));}r = sqrtf(dx * dx + dy * dy + dz * dz);
if (r < cut) {ig2 = (int)(r / del);atomicAdd(&d_g2[ig2], 2);}}''', 'cupy_pair_gpu')

NVIDIA Modulus
NVIDIA Modulus

Platform for developing Physics ML surrogate model• A training and inference pipeline - using
Physics (governing equations) and Data(simulation/observations)• Customizable and scalable platform• Higher level of abstraction for domain
experts• Build generalized AI surrogates forparameterized domain• Near real-time high-fidelity simulation

Developing digital twins for weather, climate, and energy [S41823]

Modulus EULA – It’s free.

Physics-ML categorization
Fully datadriven

https://developer.nvidia.com/modulus#Data

Inductivebias

Physicsconstrained

Fully physics
drivenPhysics

12