跳转至

2022

  • 分类于 Programming
  • 需要 1 分钟阅读时间

Go Install and Command

Install

wget https://go.dev/dl/go1.18.3.linux-amd64.tar.gz
rm -rf /usr/local/go && tar -C /usr/local -xzf go1.18.3.linux-amd64.tar.gz
(maybe need sudo)
sudo rm -rf /usr/local/go && sudo tar -C /usr/local -xzf go1.18.3.linux-amd64.tar.gz
export PATH=$PATH:/usr/local/go/bin
go version

Command usage

$ cd $HOME/go/src/hello
$ go run main.go #直接运行
Hello, World!!
$ go build # 产生可执行文件
$ ./hello
Hello, World!!

包管理

Packages

Go packages are folders that contain one more go files.

Modules

A modules (starting with vgo and go 1.11) is a versioned collection of packages.

go get github.co­m/a­nda­nhm­/go­-pr­ett­ytimee
go mod init github.co­m/a­nda­nhm­/go­-pr­ett­ytime

go list -m -u all 来检查可以升级的package，

使用go get -u need-upgrade-package 升级后会将新的依赖版本更新到go.mod

也可以使用 go get -u 升级所有依赖

作者：若与 链接：https://www.jianshu.com/p/760c97ff644c 来源：简书 著作权归作者所有。商业转载请联系作者获得授权，非商业转载请注明出处。

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://devhints.io/go

  • 分类于 Programming
  • 需要 1 分钟阅读时间

Go mod

简介

go modules 是 golang 1.11 新加的特性。现在1.12 已经发布了，是时候用起来了。Modules官方定义为：

模块是相关Go包的集合。modules是源代码交换和版本控制的单元。 go命令直接支持使用modules，包括记录和解析对其他模块的依赖性。modules替换旧的基于GOPATH的方法来指定在给定构建中使用哪些源文件。

使用

初始化项目

mkdir Gone
cd Gone
go mod init Gone
对应go.mod文件 
module Gone
go 1.14
go.mod文件一旦创建后，它的内容将会被go toolchain全面掌控。

go toolchain会在各类命令执行时，比如go get、go build、go mod等修改和维护go.mod文件。

go.mod 提供了module, require、replace和exclude 四个命令

module 语句指定包的名字（路径） require 语句指定的依赖项模块 replace 语句可以替换依赖项模块 exclude 语句可以忽略依赖项模块

自动添加依赖

对于main.go里的import 

package main

import (
    "crypto/hmac"
    "crypto/sha1"
    "encoding/hex"
    "encoding/json"
    "fmt"
    "io/ioutil"
    "log"
    "net/http"
    "os"
    "os/exec"
    "strings"
)

……
执行 go run main.go 运行代码会发现 go mod 会自动查找依赖自动下载，并修改go.mod（安装 package 的原則是先拉最新的 release tag，若无tag则拉最新的commit）

自己发布module包

结合github很简单实现

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://www.jianshu.com/p/760c97ff644c

  • 分类于 Programming
  • 需要 1 分钟阅读时间

Golang Syntax

为什么要学习go语言

  1. 同步方式轻松实现高并发，充分利用多核
  2. 基于消息传递的通信方式
  3. 适合服务器和网络编程
  4. 有垃圾回收机制
  5. 静态语言，有编译过程，和独立的静态可执行文件，只依赖glibc
  6. 不像python要安装各种库，java也要JRE
  7. 兼顾python的易开发性和c的性能
  8. 内存占用极小，支持10W+的并行

一些缺点

  1. 实际运行时，由于GC的影响，延迟会比较严重
  2. 代码会有很多重复的地方

有趣的工具

  1. gofmt
  2. gofix
  3. govet

数据类型

  • int8类型 表示 -128～127
  • Channel 类型
  • 切片类型 (可变长数组

变量声明

第一种，指定变量类型，如果没有初始化，则变量默认为零值。

//var v_name v_type
var b, c int = 1, 2

//特殊
var a *int
var a []int
var a map[string] int
var a chan int
var a func(string) int
var a error // error 是接口

第二种，根据值自行判定变量类型。

//var v_name = value
var d = true

第三种，使用声明符号:=

但是如果变量已经使用 var 声明过了，再使用 := 声明变量，就产生编译错误，格式：

v_name := value

循环语句

for key, value := range oldMap {
    newMap[key] = value
}

并发和通道通讯

go函数

Go 语言支持并发，我们只需要通过 go 关键字来开启 goroutine 即可。

goroutine 是轻量级线程，goroutine 的调度是由 Golang 运行时进行管理的。

goroutine 语法格式：go 函数名( 参数列表 )

Go 允许使用 go 语句开启一个新的运行期线程， 即 goroutine，以一个不同的、新创建的 goroutine 来执行一个函数。 同一个程序中的所有 goroutine 共享同一个地址空间。

通道（channel）

通道可用于两个 goroutine 之间通过传递一个指定类型的值来同步运行和通讯。操作符 <- 用于指定通道的方向，发送或接收。如果未指定方向，则为双向通道。

ch <- v    // 把 v 发送到通道 ch
v := <-ch  // 从 ch 接收数据
           // 并把值赋给 v

声明一个通道很简单，我们使用chan关键字即可，通道在使用前必须先创建：

ch := make(chan int)

example

1

func countGoodRectangles(rectangles [][]int) int {
    cnt, maxLen := 0, 0
    for _, rectangle := range rectangles {
        k := int(math.Min(float64(rectangle[0]), float64(rectangle[1])))
        if k == maxLen {
            cnt++
        }
        if k > maxLen {
            maxLen, cnt = k, 1
        }
    }
    return cnt
}

webhook

https://github.com/swangeese/acsa-web/tree/webhook

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://www.runoob.com/go/go-concurrent.html

  • 分类于 Tutorials
  • 需要 3 分钟阅读时间

Game Streaming & Video Streaming

视频传输编码

在初步接触了视频传输编码之后，我开始好奇Streaming采用的哪种视频编码呢？

详见 Streaming Protocol一文

Moonlight for IPAD

Nvidia Geforece界面

  1. 常规中开启分享
  3. SHEILD 开启
  5. 添加，按照C:\Windows\System32\mstsc.exe这个地址，将mstsc.exe添加进去，mstsc.exe就是你的桌面，等会串流，可以用手机直接操控你的电脑桌面。

串流画面问题

moonlight找不到电脑

https://www.bilibili.com/read/cv10239020

netsh advfirewall firewall add rule name="GameStream UDP" dir=in protocol=udp localport=5353,47995,47998-48010 action=allow
netsh advfirewall firewall add rule name="GameStream TCP" dir=in protocol=tcp localport=47984,47989,47995,48010 action=allow

还是不行，猜测是

但是这个是wifi6 11ax

尝试安装Internet-Hosting-Tool，运行有提示失败，建议重装也不行。

sjf的解决办法

  1. 卸载当前版本，然后安装3.19
  2. 打开服务
  3. 找到属性 单击打开。找到登录 复选框 把里面的登录身份选择成 本地系统账户 运行服务与桌面交换。勾选然后应用
  4. moonlight可以搜索到电脑
  5. https://pan.baidu.com/s/1x83Uk4kkYQritiNAqg_vLg [/url]提取码：1111 获得NvContainerNetworkService服务注册表文件
  6. 官网下载更新GF到最新

实际解决

  1. 官网下载更新GF到最新
  2. 通过上面的注册表添加NvContainerNetworkService服务，启动
  3. 在任务计划程序里设置, 设置开机启动moonlightNVNetStart任务
  4. 程序"C:\Program Files\NVIDIA Corporation\NvContainer\nvcontainer.exe"
  5. 参数-s NvContainerNetworkService -f "C:\ProgramData\NVIDIA\NvContainerNetworkService.log" -l 3 -d "C:\Program Files\NVIDIA Corporation\NvContainer\plugins\NetworkService" -r -p 30000 -st "C:\Program Files\NVIDIA Corporation\NvContainer\NvContainerTelemetryApi.dll"
  7. 成功
  8. 修改分辨率为ipad分辨率，全屏应用
  9. 修改英伟达控制面板的分辨率为IPAD 2388*1688 macbook 2560*1600

问题：Nvidia控制面板没有显示一项

  • 如果显卡驱动装好，且显卡都开了，但就是没有显示选项。
  • 打开服务，找到NVDisplay.ContainerLocalSystem，点登录项，将“允许服务与桌面交互（W）”前的勾打上，重启NVDisplay.ContainerLocalSystem服务，
  • 返回桌面，右键-显示设置，将分辩率任意改一个可用的-应用，
  • 桌面右键-N..控制面板，就有了显示选项，可以改2K分辩率啦。

IPAD moonlight 串流控制

type ESC and mouse scroll

对于实体键盘可以修改映射, 但是滚轮就不好用了。

与其这样不如换个思路，添加手柄，看其能不能支持滚轮和ESC。初步尝试，滚轮可以只是灵敏度有点低。AntiMicroX完美解决了这个问题，配置文件路径 I:\BT\GAME\x18Game\moonlightAntiMicroX.gamecontroller.amgp

体验十分丝滑，任意程序也可以添加。支持PS4手柄(长按PS和share键配对)

share时一定要登录steam

任意应用全屏

应用串流至少将某应用窗口转发，所以只需要停止流式传输，然后调整分辨率就行了。

晚上关闭屏幕，不休眠

  1. 方法一：管理员运行代码 @powercfg -h off

云原神测试

高画质60帧 1.6MB/s

最低画质30帧 500KB/s

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

无

  • 分类于 Tutorials
  • 需要 3 分钟阅读时间

Crawler

如何获取请求链接

这个api是怎么来的呢？ lesson_info_url = "https://www.eeo.cn/saasajax/webcast.ajax.php?action=getLessonLiveInfo"

感谢大佬回答

输入

返回数据

PHP源文件

PHP是后端语言，前端是无法查看的，前端看到的是最终运算之后的结果，PHP源代码是无法查看的。

使用

将 header改一下就能用了，注意不要开代理

from requests import Session

session = Session()
lesson_info_url = "https://www.eeo.cn/saasajax/webcast.ajax.php?action=getLessonLiveInfo"

headers = {
    'Accept': 'text/html,application/xhtml+xml,application/xml;q=0.9,*/*;q=0.8',
    'Accept-Language': 'zh-CN,zh;q=0.9,zh-TW;q=0.8,en;q=0.7',
    'User-Agent': 'Mozilla/5.0 (Windows NT 10.0; Win64; x64) AppleWebKit/537.36 (KHTML, like Gecko) Chrome/80.0.3987.149 Safari/537.36'
}

data = {
    'lessonKey': lessonKey
}

resp = session.post(url=lesson_info_url, headers=headers, data=data)
text = resp.json()
CourseName = text['data']['courseName']

urllib.request下载视频

from urllib import request

base_url = 'https://f.us.sinaimg.cn/001KhC86lx07laEy0PtC01040200y8vC0k010.mp4?label=mp4_hd&template=640x360.28&Expires=1528689591&ssig=qhWun5Mago&KID=unistore,video'
#下载进度函数
def report(a,b,c):
    '''
    a:已经下载的数据块
    b:数据块的大小
    c:远程文件的大小
    '''
    per = 100.0 * a * b / c
    if per > 100:
        per = 100
    if per % 1 == 1:
        print ('%.2f%%' % per)
#使用下载函数下载视频并调用进度函数输出下载进度
request.urlretrieve(url=base_url,filename='weibo/1.mp4',reporthook=report,data=None)

例子一

小白尝试 扒学校的资源网址(http://wlkt.ustc.edu.cn/)

爬取List读取

正则匹配video/detail出视频网址后缀

网页视频位置

正则匹配mp4.php得到视频位置http://wlkt.ustc.edu.cn/mp4.php?file=HXMEV11IQNB2ZXPM6BVWY77AJ2HZTM4U 但是不打开网站没有php返回，网页只能得到。 可通过下面API返回需要的, 可以见github代码 

opener = urllib.request.FancyURLopener({})
f = opener.open(taskUrl)
content = f.read()

#1.得到beautifulsoup对象
soup = BeautifulSoup(content,'html.parser')

#通过指定的 属性获取对象
ic(soup.find(id=glv._get(taskType)["data1id"]).attrs['value'])#单个对象
data输入 返回数据

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

科大BB clashIn 想爬录像。但是网上的两个都用不了了，想自学，改一下

https://github.com/aoxy/ClassIn-Video-Download

https://github.com/JiangGua/classin-downloader

参考文献

https://blog.csdn.net/qq_37275405/article/details/80780925

  • 分类于 Programming
  • 需要 3 分钟阅读时间

Cuda Optimize : Vectorized Memory Access

baseline

__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) { 
  int idx = blockIdx.x * blockDim.x + threadIdx.x; 
  for (int i = idx; i < N; i += blockDim.x * gridDim.x) { 
    d_out[i] = d_in[i]; 
  } 
} 

void device_copy_scalar(int* d_in, int* d_out, int N) 
{ 
  int threads = 128; 
  int blocks = min((N + threads-1) / threads, MAX_BLOCKS);  
  device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N); 
}

简单的分块拷贝。

通过cuobjdump -sass executable.得到对应的标量copy对应的SASS代码

/*0058*/ IMAD R6.CC, R0, R9, c[0x0][0x140]                
/*0060*/ IMAD.HI.X R7, R0, R9, c[0x0][0x144]              
/*0068*/ IMAD R4.CC, R0, R9, c[0x0][0x148]               
/*0070*/ LD.E R2, [R6]                                   
/*0078*/ IMAD.HI.X R5, R0, R9, c[0x0][0x14c]              
/*0090*/ ST.E [R4], R2

（SASS不熟悉，请看SASS一文）

其中4条IMAD指令计算出读取和存储的指令地址R6:R7和R4:R5。第4和6条指令执行32位的访存命令。

Vector way1: CUDA C/C++ standard headers

通过使用int2, int4, or float2

比如将int的指针d_in类型转换然后赋值。

reinterpret_cast<int2*>(d_in)
// simple in C99
(int2*(d_in))

但是需要注意对齐问题，比如

reinterpret_cast<int2*>(d_in+1)

这样是非法的。

Vector way2: structures

通过使用对齐的结构体来实现同样的目的。

struct Foo {int a, int b, double c}; // 16 bytes in size
Foo *x, *y;
…
x[i]=y[i];

实际修改LD.E.64

执行for循环次数减半，注意边界处理。

__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
  }

  // in only one thread, process final element (if there is one)
  if (idx==N/2 && N%2==1)
    d_out[N-1] = d_in[N-1];
}

void device_copy_vector2(int* d_in, int* d_out, int n) {
  threads = 128; 
  blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS); 

  device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

对应汇编可以看出

/*0088*/                IMAD R10.CC, R3, R5, c[0x0][0x140]              
/*0090*/                IMAD.HI.X R11, R3, R5, c[0x0][0x144]            
/*0098*/                IMAD R8.CC, R3, R5, c[0x0][0x148]             
/*00a0*/                LD.E.64 R6, [R10]                                      
/*00a8*/                IMAD.HI.X R9, R3, R5, c[0x0][0x14c]           
/*00c8*/                ST.E.64 [R8], R6

变成了LD.E.64

实际修改LD.E.128

执行for循环次数减半，注意边界处理。

__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
  }

  // in only one thread, process final elements (if there are any)
  int remainder = N%4;
  if (idx==N/4 && remainder!=0) {
    while(remainder) {
      int idx = N - remainder--;
      d_out[idx] = d_in[idx];
    }
  }
}

void device_copy_vector4(int* d_in, int* d_out, int N) {
  int threads = 128;
  int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);

  device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

对应汇编可以看出

/*0090*/                IMAD R10.CC, R3, R13, c[0x0][0x140]              
/*0098*/                IMAD.HI.X R11, R3, R13, c[0x0][0x144]            
/*00a0*/                IMAD R8.CC, R3, R13, c[0x0][0x148]               
/*00a8*/                LD.E.128 R4, [R10]                               
/*00b0*/                IMAD.HI.X R9, R3, R13, c[0x0][0x14c]             
/*00d0*/                ST.E.128 [R8], R4

变成了LD.E.128

summary

(个人感觉，提升也不大吗？也没有两倍和四倍的效果)

绝大部分情况，向量比标量好， increase bandwidth, reduce instruction count, and reduce latency. 。

但是会增加额外的寄存器(SASS里也没有看到？？)和降低并行性(什么意思？？？)

参考文献

https://developer.nvidia.com/blog/cuda-pro-tip-increase-performance-with-vectorized-memory-access/#entry-content-comments

  • 分类于 Architecture
  • 需要 3 分钟阅读时间

cuda Assembly:PTX & SASS

两种汇编

  1. parallel thread execution (PTX) 内联汇编有没有关系
  2. PTX是编程人员可以操作的最底层汇编，原因是SASS代码的实现会经常根据GPU架构而经常变换
  3. https://docs.nvidia.com/cuda//pdf/Inline_PTX_Assembly.pdf
  4. ISA指令手册 https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#instruction-set
  6. SASS
  7. Streaming ASSembly(Shader Assembly?) 没有官方的证明
  8. 没有官方详细的手册，有基本介绍：https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html#ampere
  9. https://zhuanlan.zhihu.com/p/161624982
  10. 从可执行程序反汇编SASS
    1. https://www.findhao.net/easycoding/2339.html

SASS 指令基本信息

对于Ampere架构

指令方向 

(instruction) (destination) (source1), (source2) ...

各种寄存器说明 * RX for registers * URX for uniform registers * SRX for special system-controlled registers * PX for predicate registers * c[X][Y] for constant memory

SASS 举例说明1

SASS的难点在于指令的后缀。由于手册确实，需要结合PTX的后缀查看

/*0028*/         IMAD R6.CC, R3, R5, c[0x0][0x20]; 
/*0030*/         IMAD.HI.X R7, R3, R5, c[0x0][0x24]; 
/*0040*/         LD.E R2, [R6]; //load

line1

/*0028*/ IMAD R6.CC, R3, R5, c[0x0][0x20];
Extended-precision integer multiply-add: multiply R3 with R5, sum with constant in bank 0, offset 0x20, store in R6 with carry-out.

c[BANK][ADDR] is a constant memory。

.CC means “set the flags”

line2

/*0030*/ IMAD.HI.X R7, R3, R5, c[0x0][0x24];
Integer multiply-add with extract: multiply R3 with R5, extract upper half, sum that upper half with constant in bank 0, offset 0x24, store in R7 with carry-in.

line3

/*0040*/         LD.E R2, [R6]; //load
LD.E is a load from global memory using 64-bit address in R6,R7(表面上是R6，其实是R6 与 R7 组成的地址对)

summary

R6 = R3*R5 + c[0x0][0x20], saving carry to CC
R7 = (R3*R5 + c[0x0][0x24])>>32 + CC
R2 = *(R7<<32 + R6)
寄存器是32位的原因是 SMEM的bank是4字节的。c数组将32位的基地址分开存了。

first two commands multiply two 32-bit values (R3 and R5) and add 64-bit value c[0x0][0x24]<<32+c[0x0][0x20],

leaving 64-bit address result in the R6,R7 pair

对应的代码是 

kernel f (uint32* x) // 64-bit pointer
{
   R2 = x[R3*R5]
}

SASS Opt Code分析2

  • LDG - Load form Global Memory
  • ULDC - Load from Constant Memory into Uniform register
  • USHF - Uniform Funnel Shift （猜测是特殊的加速shift）
  • STS - Store within Local or Shared Window

流水STS

观察 偏移 * 4 * 2060(delta=2056) * 4116(delta=2056) * 8228(delta=2 * 2056) * 6172(delta=-1 * 2056) * 10284(delta=2 * 2056) * 12340(delta=2056)

可见汇编就是中间写反了，导致不连续，不然能隐藏更多延迟

STS缓存寄存器来源

那么这些寄存器是怎么来的呢？感觉就是写反了 

      IMAD.WIDE.U32 R16, R16, R19, c[0x0][0x168] 
      LDG.E R27, [R16.64] 
      IMAD.WIDE R30, R19, c[0x0][0x164], R16 
      LDG.E R31, [R30.64] 
      IMAD.WIDE R32, R19, c[0x0][0x164], R30 
      LDG.E R39, [R32.64] 
      # important R41 R37
      IMAD.WIDE R34, R19, c[0x0][0x164], R32 
      IMAD.WIDE R40, R19, c[0x0][0x164], R34 
      LDG.E R41, [R40.64] 
      LDG.E R37, [R34.64]

Fix

原因是前面是手动展开的，假如等待编译器自动展开for循环就不会有这个问题

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

参考文献

https://forums.developer.nvidia.com/t/solved-sass-code-analysis/41167/2

https://stackoverflow.com/questions/35055014/how-to-understand-the-result-of-sass-analysis-in-cuda-gpu

  • 分类于 Tutorials
  • 需要 2 分钟阅读时间

Cuda Optimize : Stencil

课程报告PPT

有对应的PPT，代码。

最终将1000ms程序优化到1~2ms

乔良师兄有根据知乎介绍如何利用寄存器文件缓存

SMEM难点: 跨线程访存

  1. 不仅每个线程需要访问自己划分对应区域之外的元素
  2. 而且访问的总个数也不是线程数对应的倍数

导致Embarrassingly Parallel Problems

1D 梯度计算 Stencil实例

计算某点的梯度，需要前后的function值。

Halo/Ghost Cells 光晕

问题: 对于边界上的cells，需要访问相邻区域的元素。

解决办法: 将他们也加入进当前block的SMEM

Indexing with Halo Cells

  1. Stencil问题的半径 radius (RAD) 是边缘元素需要的某方向的额外元素
  2. 在梯度的例子里是1
  3. SMEM声明的大小，需要在每个维度上都增加 2*RAD的个数
  4. 这导致SMEM的index的每个维度需要增加RAD. s_idx = threadIdx.x + RAD;

code

int main() {
    const float PI = 3.1415927;
    const int N = 150;
    const float h = 2 * PI / N;
    float x[N] = { 0.0 };
    float u[N] = { 0.0 };
    float result_parallel[N] = { 0.0 };
    for (int i = 0; i < N; ++i) {
        x[i] = 2 * PI*i / N;
        u[i] = sinf(x[i]);
    }
    ddParallel(result_parallel, u, N, h);
}

Kernel Launching

#define TPB 64
#define RAD 1 // radius of the stencil
…
void ddParallel(float *out, const float *in, int n, float h) {
    float *d_in = 0, *d_out = 0;
    cudaMalloc(&d_in, n * sizeof(float));
    cudaMalloc(&d_out, n * sizeof(float));
    cudaMemcpy(d_in, in, n * sizeof(float), cudaMemcpyHostToDevice);

    // Set shared memory size in bytes
    const size_t smemSize = (TPB + 2 * RAD) * sizeof(float);
    ddKernel<<<(n + TPB - 1)/TPB, TPB, smemSize>>>(d_out, d_in, n, h);
    cudaMemcpy(out, d_out, n * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_in);
    cudaFree(d_out);
}

Kernel Definition

__global__ void ddKernel(float *d_out, const float *d_in, int size, float h) {
    const int i = threadIdx.x + blockDim.x * blockIdx.x;
    if (i >= size) return;

    const int s_idx = threadIdx.x + RAD;
    extern __shared__ float s_in[];

    // Regular cells
    s_in[s_idx] = d_in[i];
    // Halo cells
    if (threadIdx.x < RAD) {
        s_in[s_idx - RAD] = d_in[i - RAD];
        s_in[s_idx + blockDim.x] = d_in[i + blockDim.x];
    }
    __syncthreads();
    d_out[i] = (s_in[s_idx-1] - 2.f*s_in[s_idx] + s_in[s_idx+1])/(h*h);
}

需要进一步的研究学习

暂无

遇到的问题

暂无

开题缘由、总结、反思、吐槽~~

研一下USTC并行计算自己的选题

参考文献

https://dumas.ccsd.cnrs.fr/dumas-00636254/document

https://indico.fysik.su.se/event/6743/contributions/10338/attachments/4175/4801/4.CUDA-StencilsSharedMemory-Markidis.pdf