reminder

SST和所有外部组件用相同编译器编译

术语

SM : streaming Multiprocessor
基础模块，有4block ( 16 fp32, 8 fp64, 16 int32, 128K L1 cache, 64k register）

SIMT : single instruction multiple threads
一组32 cores执行相同指令不同数据

gridDim.x
blockIdx.x
blockDim.x
threadIdx.x

nvcc --ptxas-options=-v : 显示寄存器和 shared mem 信息

概念

• Page-Locked Memory : Unified Virtual Addressing (UVA)，CPU 和 GPU 共享 VA，实现 zero copy，所有对 CPU 内存访问走 PCIe 在 CPU 使用 page-locked host memory 使用 cudaHostAlloc 和 cudaFreeHost 分配新内存 使用 cudaHostRegister 将 malloc 的内存变为 page-locked
• Unified Memory : GPU 自动管理内存，page fault 时从 CPU 拷贝到 GPU

usage

golbal : host调用device执行函数
__syncthreads() 同步一个block内的thread
shared int a[] : 每个block一个共享内存中的a
T __ldg(const T address) 从address读取一个数据T (从read-only data cache中读取)

thrust:

cuda 中类似 C++ STL

thrust::device_ptr<int> d_arr = thrust::device_pointer_cast<int>(d_bin_counter);
// host指针转device_ptr
thrust::exclusive_scan(d_arr, d_arr + SEGBIN_NUM, d_arr); // 每个位置变为自己之前不包括自己的前缀和

atomicAdd(int* address, int val) 将 val 原子加到 address

cudaDeviceSynchronize() block直到所有设备同步

GPU中所有 active thread 被分配了单独的寄存器，当切换线程时不需要交换寄存器

cuda 中所有 kernel launch 都为异步
在开始和结束 CPU timer 之前都需要调用 cudaDeviceSynchronize

cuda 中 default stream (stream 0) 可以保证所有之前任何 stream 的调用均在前执行，且所有之后任何 stream 的调用均在后执行 (serializing)

事件计时：

cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

只能在 stream 0 中用此方法计时，其他 stream 无法保证顺序

多GPU

int deviceCount;
cudaGetDeviceCount(&deviceCount); // device 数量
cudaGetDeviceProperties(&deviceProp, 1); // 获取 device 1 的 properties
cudaSetDevice(0); // 设置当前 device 为 0，之后的内存分配和 kernel launch 等操作均在 device 0 上进行

特性

在 CUDA 中用户可以创建多个 stream，但同时不会有超过 32 个 kernel 执行
在一些 device 上， L1 cache 和 shared memory 共享硬件资源

optimization

#pragma unroll [factor] : 强制循环展开，factor为展开参数，factor=1为不展开
restrict : 指针重名优化（Pointer Aliasing），显示指定没有重名指针

A100

tensor core 理论 315TFLOPS
A100 dram 1500GB/s
A100 L2 40MB 所有SM共享 V100 L2 6MB

compute capability :
A100 : 8.0
V100 : 7.0
P100 : 6.0

额，请在编译时指定正确的 compute capability

hip

hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream*/, a, b, c, n /*kernel函数参数*/);
// Replace MyKernel<<<dim3(gridDim), dim3(gridDim), 0, 0>>> (a,b,c,n);

我想要拥有力量
这样当面对现实的压迫时才可以选择不去妥协，才可以去反抗
否则只能看着所爱之物被践踏
这会成为我前进的理由之一
但是请记住伴随力量而来的责任，面对弱者时也不要忘记过去同样弱小的自己

又到了一年NOI的时间
到了现在，即使是OI相关的回忆，也变得温暖起来了呢

在京吹之前也看过一些京阿尼的作品：
冰菓，中二病，紫罗兰，clannad，轻音 可能还有一些吧

从这些作品中我可以感受到京阿尼在作品中的投入，人物塑造的饱满，不同风格的尝试，制作的精良
京阿尼的作品总是散发着一种文艺的忧伤感，但又使人有一种莫名的距离感，如同小心地欣赏一件艺术品。

此前京阿尼一直给我一种值得尊敬的感觉。京阿尼不会去制作烂俗的青春校园喜剧，作画一向保持着很高的水准，但我却一直都没有对京阿尼抱有太多特殊的感情。

但京吹让我改变了这个想法。京吹中的人物透露着很强的真实感，其中的事件也给人带来了一种很强的共鸣感。
北宇治不是一个完美的中学。主角们不是过着无忧无虑，每天在社团悠哉嬉笑的生活。她们是一群为梦想而奋斗着的普通中学生。她们也面临着激烈的比赛与竞争，她们需要为了练习而早起晚归。她们有着普通中学生的苦恼，来自家长同学的压力。但也正因如此，她们在舞台上的身影才如此地闪耀，她们为梦想而奋斗的样子才这样地打动人心。