NVIDIA cuda programming, open source and AI

NVIDIA cuda programming, open source and AI

1. CUDA & AI
Ph.DTaewookKang
[email protected]
daddynkidsmakers.blogspot.com

2. Maker, Ph.D.
12 books author
TK. Kang
BIM principle and Digital transformation: BIM principle & DX
사이트소개, Profile (dxbim.blogspot.com)

3. Overview

4. Open source service market
Open Source Services Market worth 32.95 Billion
USD by 2022

5. Open source -GithubOpenSlicer

6. Open source
http://guswnsxodlf.github.io/software-license
GNU General Public License(GPL) 2.0
–의무엄격. SW 수정및링크경우소스코드 제공의무
GNU Lesser GPL(LGPL) 2.1
–저작권표시. LPGL 명시. 수정한라이브러리 소스코드 공개
Berkeley Software Distribution(BSD) License
–소스코드 공개의무 없음. 상용SW 무제한사용가능
Apache License
–BSD와유사. 소스코드 공개의무 없음
Mozilla Public License(MPL)
–소스코드 공개의무 없음. 수정코드는MPL에의해배포
MIT License
–라이선스 / 저작권만 명시조건

7. Autonomous Vehicle-CARLA for Urban Driving Simulator

8. Deep Learning -You only look once (YOLO)

9. Deep Learning –Google Tensorflow
CNN
(convolution neural network)
Deep Learning
Feature –classification
Learning
KICT

10. IoT–Embedded computer for prototyping

11. Cloud platform–MQTT, RaspberryPI, Blynk, ITFFF
Packing
Wireless
Sensor
Gateway
IoT
Control
Big data
analysis
Protocol
IoT
connection
service
A BIM ANALYSIS OF HVAC AND RADIANT COOLING SOLUTIONS, ROBERT CUB ICK, 2016
KICT

12. Cesium
http://cesiumjs.org/Seattle/
GIS-Cesium for visualization
KICT

13. BIM –xBIMToolkit
KICT
4Projects –Collaborative BIM Tool

14. Robotics –ROS(robot operation system)

15. Robotics –Robot Operating System(ROS). Open Drone Map. Arduino. Mission planner.
Trimble GPS
카메라
카메라
스캐너
IMU
DMI
KICT

16. Blockchain–Smart ContractPragma solidity ^0.4.0
Contract Lottery {
uintcount = 0;
uintprize_money= 0;
function () payable public {
count++;
prize_money+= msg.value;
if(count % 7 == 0) {
msg.sender.transfer(prize_money);
prize_money= 0;
}
}
}

17. Construction startups

18. IoT-based field management -Pillar
Pillar Technology

19. Construction safety –Spot-r

20. Field Data Mining–Smartvid.io

21. Autonomous Field Monitoring -Doxel

22. Autonomous construction –Built Robotics

23. Startupswith open source
Smart construction business model matrix

24. CUDA

27. CUDA
Ray Tracey's blog: GPU path tracing tutorial 3: GPU-friendly
Acceleration Structures. Now you're cooking with GAS!

28. CUDA

29. CUDA
•CUDA archive download
•cuDNN download
•CUDA, cuDNN 설치, 버전확인및개발(daddynkidsmakers.blogspot.com)
•GPU CUDA 기반데이터병렬처리 지원파이썬Numba 라이브러리 사용법소개
•mac999/cuda: CUDA programming example (github.com)

30. CUDA

31. CUDA

32. CUDA
CUDA C++ Programming Guide (nvidia.com)
CUDA_C_Programming_Guide.pdf (nvidia.com)

33. CUDA
CUDA 아키텍처(Multi-Process Service :: GPU Deployment and
Management Documentation)
CUDA C++ Programming Guide (nvidia.com)
CUDA_C_Programming_Guide.pdf (nvidia.com)

34. CUDA
CUDA 아키텍처(NVIDIA H100 Tensor Core GPU Architecture
Overview)
NVIDIA Hopper 아키텍처 심층분석하기 -NVIDIA Technical Blog
Streaming Multiprocessor

35. CUDA
CUDA 아키텍처 (Multi-Process Service :: GPU Deployment and
Management Documentation)
Tesla Fermi Kepler Maxwell Pascal …

36. CUDA
CUDA 아키텍처(NVIDIA GPU Architecture: from Pascal to Turing to
Ampere -WOLF Advanced Technology)

37. CUDA

38. CUDA

39. CUDA
Grid > Block > Thread
func<<<B, S>>>();

40. CUDA
Grid > Block > Thread
__global__
void add(int n, float *x, float *y)
{
int index = threadIdx.x;
int stride = blockDim.x;
for (int i = index; i < n; i += stride)
y[i] = x[i] + y[i];
}
add<<<1, 256>>>(N, x, y);
Block, Thread number

41. CUDA
N=data size for calculation
x, y=vector array
An Even Easier Introduction to CUDA | NVIDIA Technical Blog

42. CUDA

43. CUDA

44. CUDA
NVIDIA/cuda-samples: Samples for CUDA Developers which demonstrates features in CUDA Toolkit
(github.com)

45. CUDA

46. CUDA
#include <stdio.h>
#include <cuda_runtime.h> // For the CUDA runtime routines (prefixed with
"cuda_")
#include <helper_cuda.h> // CUDA Kernel Device code
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements)
{
C[i] = A[i] + B[i];
}
}
https://docs.nvidia.com/cuda/cuda-c-programming-
guide/index.html#c-language-extensions
__global__ __device____host__ __shared__
char, short, int, long, longlong, float, double

47. CUDA
int main(void)
{
cudaError_t err = cudaSuccess; // Error code to check return values for CUDA calls
// prepare input dataset
int numElements = 50000;
size_t size = numElements * sizeof(float);
printf("[Vector addition of %d elements]\n", numElements);
// allocate host CPU input vectors and init
float *h_A = (float *)malloc(size); // Allocate the host input vector A
float *h_B = (float *)malloc(size); // vector B
float *h_C = (float *)malloc(size); // vector C
if (h_A == NULL || h_B == NULL || h_C == NULL)
{
fprintf(stderr, "Failed to allocate host vectors!\n");
exit(EXIT_FAILURE);
}
for (int i = 0; i < numElements; ++i)
{
h_A[i] = rand()/(float)RAND_MAX;
h_B[i] = rand()/(float)RAND_MAX;
}
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory

48. CUDA
for (int i = 0; i < numElements && i < 10; i++)
{
fprintf(stderr, "A=%.2f, B=%.2f\n", h_A[i], h_B[i]);
}
// allocate device CUDA input vector A, B, C
float *d_A = NULL;
err = cudaMalloc((void **)&d_A, size);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
float *d_B = NULL;
err = cudaMalloc((void **)&d_B, size);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}

49. CUDA
float *d_C = NULL;
err = cudaMalloc((void **)&d_C, size);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Copy the host input vectors A and B in host memory to the device input vectors in device memory
printf("Copy input data from the host memory to the CUDA device\n");
err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector A from host to device (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector B from host to device (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}

50. CUDA
// Launch the Vector Add CUDA Kernel
int threadsPerBlock = 256;
int blocksPerGrid =(numElements + threadsPerBlock -1) / threadsPerBlock;
printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid, threadsPerBlock);
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements);
err = cudaGetLastError();
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}

51. CUDA
// Copy the device result vector in device memory to the host result vector in host memory.
printf("Copy output data from the CUDA device to the host memory\n");
err = cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to copy vector C from device to host (error code %s)!\n",
cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Verify that the result vector is correct
for (int i = 0; i < numElements; ++i)
{
if (i <= 10)
fprintf(stderr, "A=%.2f, B=%.2f, C=%.2f\n", h_A[i], h_B[i], h_C[i]);
if (fabs(h_A[i] + h_B[i] -h_C[i]) > 1e-5)
{
fprintf(stderr, "Result verification failed at element %d!\n", i);
exit(EXIT_FAILURE);
}
}
printf("Test PASSED\n");

52. CUDA
// Free device global memory
err = cudaFree(d_A);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector A (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_B);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector B (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
err = cudaFree(d_C);
if (err != cudaSuccess)
{
fprintf(stderr, "Failed to free device vector C (error code %s)!\n", cudaGetErrorString(err));
exit(EXIT_FAILURE);
}
// Free host memory
free(h_A);
free(h_B);
free(h_C);
printf("Done\n");
return 0;
}

53. CUDA

54. CUDA
numElements = 500,000,000

55. CUDA
numElements = 500,000,000
nvcc -o test_cpu_gpu -I "C:\ProgramData\NVIDIA Corporation\CUDA Samples\v11.3\common\inc"
test_cpu_gpu.cu
Step by Step: CUDA 7.5 on Windows |
Neuralotic (xilodyne.com)

56. CUDA
numElements = 500,000,000

57. CUDA CUDA_Compiler_Driver_NVCC.pdf (nvidia.com)

58. CUDA CUDA_Compiler_Driver_NVCC.pdf (nvidia.com)

59. CUDA

60. CUDA
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}
CUDA C++ Programming
Guide (nvidia.com)

61. CUDA

62. CUDA
// 각정점에대해, 정점데이터를 처리하는 BLOCK, THREAD X, Y 위치를입력받고 ,
// SIN, COS함수를통과해, Y값인W값을계산해, POS에저장함.
// 각위치데이터계산을병렬처리하므로 , 전체실수연산시간이크게줄어듬을 보여줌.
__global__ void simple_vbo_kernel(float4 *pos, unsigned int width, unsigned int height, float time)
{
unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
// calculate uv coordinates
float u = x / (float) width;
float v = y / (float) height;
u = u*2.0f -1.0f;
v = v*2.0f -1.0f;
// calculate simple sine wave pattern
float freq = 4.0f;
float w = sinf(u*freq + time) * cosf(v*freq + time) * 0.5f;
// write output vertex
pos[y*width+x] = make_float4(u, w, v, 1.0f);
}
// 매시간에대한메쉬의정점높이값을 실수계산, 병렬처리.
void launch_kernel(float4 *pos, unsigned int mesh_width,
unsigned int mesh_height, float time)
{
// execute the kernel
dim3 block(8, 8, 1);
dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
simple_vbo_kernel<<< grid, block>>>(pos, mesh_width, mesh_height, time); // 블록과스레드유닛병렬처리
}

63. CUDA
void runCuda(struct cudaGraphicsResource **vbo_resource)
{
// map OpenGL buffer object for writing from CUDA
float4 *dptr;
checkCudaErrors(cudaGraphicsMapResources(1, vbo_resource, 0));
size_t num_bytes;
checkCudaErrors(cudaGraphicsResourceGetMappedPointer((void **)&dptr, &num_bytes,
*vbo_resource));
// execute the kernel
// dim3 block(8, 8, 1);
// dim3 grid(mesh_width / block.x, mesh_height / block.y, 1);
// kernel<<< grid, block>>>(dptr, mesh_width, mesh_height, g_fAnim);
launch_kernel(dptr, mesh_width, mesh_height, g_fAnim); // 병렬처리계산
// unmap buffer object
checkCudaErrors(cudaGraphicsUnmapResources(1, vbo_resource, 0));
}

64. CUDA
// CALL BACK 함수
void display()
{
sdkStartTimer(&timer);
// run CUDA kernel to generate vertex positions
runCuda(&cuda_vbo_resource);
glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
// set view matrix
glMatrixMode(GL_MODELVIEW);
glLoadIdentity();
glTranslatef(0.0, 0.0, translate_z);
glRotatef(rotate_x, 1.0, 0.0, 0.0);
glRotatef(rotate_y, 0.0, 1.0, 0.0);
// render from the vbo
glBindBuffer(GL_ARRAY_BUFFER, vbo);
glVertexPointer(4, GL_FLOAT, 0, 0);
glEnableClientState(GL_VERTEX_ARRAY);
glColor3f(1.0, 0.0, 0.0);
glDrawArrays(GL_POINTS, 0, mesh_width * mesh_height);
glDisableClientState(GL_VERTEX_ARRAY);
glutSwapBuffers();
g_fAnim += 0.01f;
sdkStopTimer(&timer);
computeFPS();
}
void timerEvent(int value)
{
if (glutGetWindow())
{
glutPostRedisplay();
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
}
}
Int main()
{
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
}

65. CUDA

66. CUDA

67. CUDA
CUDA Math API :: CUDA Toolkit
Documentation (nvidia.com)
CUDA Math API :: CUDA Toolkit
Documentation (nvidia.com)

68. CUDA
CUDA C++ Programming Guide (nvidia.com)
The graph can only
contain kernel nodes,
memcpy nodes, memset
nodes, and child graph
nodes.

69. CUDA
CUDA C++ Programming Guide
(nvidia.com)

70. CUDA python
Numba: A High Performance Python
Compiler (pydata.org)
numba/numba: NumPy aware dynamic
Python compiler using LLVM (github.com)

71. CUDA python
•NVIDIA Compute Capability 5.0 이상CUDA 지원GPU 장착PC (2023.12 시점)
•NVIDIA CUDA 11.2 이상(링크참고)
•NVIDIA TX1, TX2, 자비애, 젯슨나노
•GTX 9, 10, 16 시리즈. RTX 20, 30, 40 시리즈. H100 시리즈
CONDA 환경의경우, 다음과같이CUDA 툴킷을자동설치.
conda install cudatoolkit
넘바는cuda python을이용해NVIDIA GPU CUDA와바인딩.
conda install nvidia::cuda-python
conda install numba
Numba: A High Performance Python
Compiler (pydata.org)

72. CUDA python
from numba import cuda
@cuda.jit
def increment_by_one(an_array):
tx = cuda.threadIdx.x# cuda 객체에포함된쓰레드ID 획득
ty = cuda.blockIdx.x# 그리드내 블럭ID 획득
bw = cuda.blockDim.x# 블럭폭(쓰레드갯수) 획득
# Compute flattened index inside the array
pos = tx + ty * bw# 해당스레드접근인덱스계산
if pos < an_array.size:# 전달된배열경계체크
an_array[pos] += 1# 배열값하나씩증가
threadsperblock = 32
blockspergrid = (an_array.size + (threadsperblock -1)) // threadsperblock
increment_by_one[blockspergrid, threadsperblock](an_array)
cpu_mem = np.arrage(10)
cuda_mem = cuda.to_device(cpu_mem)
cpu_mem = cuda_mem.copy_to_host()

73. CUDA python
•커널함수내에서메모리를 동적할당가능(참고-Memory management).
•예외처리 지원안됨. 일반적으로 Zero Divide 에러발생안됨
•list, dic, set 사용안됨
•print 사용가능
•재귀함수 지원
•abs(), bool, complex, enumerate(), float, int, len(), min(), max(), pow(), range(),
round(), zip() 지원
•cmath, math 라이브러리 지원
•operator 연산자add, eq, ge, sub, xor 등함수지원
•NumPy 일부함수지원(예. sin, cos, tan, deg2rad 등). 배열생성등메모리할당관련
함수는지원안됨
•Fastmath, cuda.random,cuda.select_device 함수등지원
•공유메모리락등처리위한atomic,thread, lock 함수지원

74. CUDA python
커널함수디버깅을 위해서는 다음과같은pdb를사용.
from pdb import set_trace; set_trace()
이외, 기존쿠다코드인cu 루틴을호출할수있는기능을지원.
@cuda.jit(link=['functions.cu'])
def multiply_vectors(r, x, y):
i = cuda.grid(1)
if i < len(r):
r[i] = mul(x[i], y[i])

75. CUDA python
import numpy as np
from numba import cuda
from tqdm import tqdm
import time
@cuda.jit
def f(a, b, c):
tid = cuda.grid(1)# threadIdx.x + (blockIdx.x * blockDim.x)
# if tid % 100 == 0:
#print('tid=', tid)
size = len(c)
if tid < size:
c[tid] = a[tid] + b[tid]
N = 10000000
a = cuda.to_device(np.random.random(N))# 천만개난수생성및a 배열값할당
b = cuda.to_device(np.random.random(N))# b 배열값할당
c = cuda.device_array_like(a) # 계산된값전달받을 c 배열생성

76. CUDA python
# CPU 모드커널함수계산성능측정
start_time_cpu = time.time()
f.forall(len(a))(a, b, c)
end_time_cpu = time.time()
cpu_execution_time = end_time_cpu -start_time_cpu
print(f'CPU len={len(a)}, ', c.copy_to_host())
# CUDA 모드커널함수계산성능측정
start_time_cuda = time.time()
nthreads = 256# Enough threads per block for several warps per block
nblocks = (len(a) // nthreads) + 1# Enough blocks to cover the entire vector depending on its length
f[nblocks, nthreads](a, b, c)
end_time_cuda = time.time()
cuda_execution_time = end_time_cuda -start_time_cuda
print(f'CUDA len={len(a)}, ', c.copy_to_host())
print(f"CUDA threads: {nthreads}, blocks: {nblocks}")
print(f"CPU Execution Time: {cpu_execution_time:.6f} seconds")
print(f"CUDA Execution Time: {cuda_execution_time:.6f} seconds")

77. CUDA python
Examples —Numba 0+untagged.2155.g9ce83ef.dirty documentation

78. CUDA -3D shader
Chapter 32. An Introduction to Shader Interfaces | NVIDIA Developer
Shader Basics -The GPU Render Pipeline | GPU Shader Tutorial (shader-tutorial.dev)

79. CUDA -3D shader
The Cg Tutorial -Chapter 1. Introduction (nvidia.com)

80. CUDA AI
mac999 (Tae wook) (github.com)
mac999/computer_vision_deeplearning: computer vision based on deep learning
lecture materials (github.com)

81. CUDA AI –smart building management
RANSAC

82. CUDA AI
Linear(4, out = 6)
ReLU()
BatchNorm1d(in = 6)
Linear(in = 6, out = 16)
ReLU()
BatchNorm1d(in = 16)
Linear(in = 16, out = 9)
ReLU()
…
ReLU()
BatchNorm1d(in = 3)
Linear(in = 3, out = 1))
epochs = 5000
learning_rate= 0.001
batch_size= 64

83. CUDA AI
Acc = 0.95

84. 1.Shafiee,Mohammad Javad,BrendanChywl,FrancisLi,andAlexanderWong."FastYOLO:AFastYouOnlyLookOnce
SystemforReal-timeEmbeddedObjectDetectioninVideo."arXivpreprintarXiv:1709.05943(2017).
2.AlanSafe,2016.2.12,HowtheInternetofThingsisImpactingtheConstructionIndustry,ForConstructionPros.com
3.RachelBurger,2015.7.28,ThreeWaystheInternetofThingsCanBenefitYourConstructionProject,Construction
Management
4.JacqiLevy,2016.4.28,4BIGwaystheIoTisimpactingdesignandconstruction,InternetofThingsblog,IBM
5.whitelightgroup,2014.8.18,HowtheInternetofThingsistransformingtheconstructionindustry
6.RachelBurger,2016.8.5,How"TheInternetofThings"isAffectingtheConstructionIndustry,thebalance.com
7.AIG,HumanConditionSafety:UsingSensorstoImproveWorkerSafety
8.NiinaGromov,2015.11.23,OfferingValuethroughInternetofThingsCase:ConstructionCompaniesinFinland,Schoolof
Science,AaltoUniversity
9.WiproDigital,2016.4.1,CASESTUDY:INCREASINGCUSTOMERVALUETHROUGHIOTFORJCBINDIA
10.MonitorDeloitte,2015.7,Every.Thing.Connected.
11.LauraBlack,2015.8.12,AnInsideLookatAutodesk’sProjectAquila,ConstructionTech
12.JeffWalsh,2015.10.1,HumanConditionAimstoTransformConstruction-SiteSafetyWithWearables,Lineshapespace.com
13.Insights,IoTLogisticsAreTransformingtheTruckingIndustry
14.ChrisTopham,2015.9.10,CaseStudy:NorthumbriaSpecialistCareHospitalPushesKNXintotheIoT,AbtecBuilding
Technologies
15.MikeChino,2015.11.6,Intel’sSmartTinyHousepacksfuturistictechnologyinto264squarefeet,inhabitat
16.WandaLau,2016.5.9,KieranTimberlakeOffersaNewToolforArchitectsWantinganInonIoT
17.CADDIGEST,2016.7.7,IBMWatsonIoTPlatformtoAddIntelligencetoBuildingsWorldwide
18.http://daddynkidsmakers.blogspot.com/
19.OxfordRoboticsInstitute,2017,Vote3Deep:FastObjectDetectionin3DPointCloudsUsingEfficientConvolutionalNeural
Networks,ICRA
20.강태욱,임지순역,2015.2,스마트홈오토메이션 ,씨아이알
21.강태욱,현소영역,2014.12,스마트빌딩시스템,씨아이알
22.강태욱,2011,BIM의원리
Reference
KICT

85. Reference

87. Thanks
Daddy Makers
(daddynkidsmakers.blogspot.com)
Computer graphics digest on Apple Podcasts
ApplePodcasts －《BIM digest》
Software engineering digest on Apple
Podcasts
[email protected]