Hello World

发表于 2026-01-24

Welcome to Hexo! This is your very first post. Check documentation for more info. If you get any problems when using Hexo, you can find the answer in troubleshooting or you can ask me on GitHub.

Quick Start

Create a new post

1	$ hexo new "My New Post"

More info: Writing

Run server

1	$ hexo server

More info: Server

Generate static files

1	$ hexo generate

More info: Generating

Deploy to remote sites

1	$ hexo deploy

More info: Deployment

Client 实现分析

time：2026_1_21

app 依赖包

ClientRender
可视化

GameConfig
client 基础信息参数后续可以扩展到server 公用一个确保准确

InputPredicition
预测 x,y 的动作

主体部分

ClientCtx

lab::net::UdpSocket sock; 发收
lab::net::UdpAddr server{}; server 接受

localHist 本地
remoteHist 预测
stateHist 权威

BuildCmdVec （localPid localCmd remoteCmds）
通过pid识别是否是本段，保存一个vectorcmds size = 2的状态，本段保存输入，对端保存预测
return （输入，预测）// vectorcmds size 2

GetRemoteCmdForTick ctx pid t

return (预测)

ApplyAuthoritativeState ctx st

整理

主要就就两条思路接受处理，发收处理维护不同的状态

OnUdp
Recv：只负责收包 -> 更新 ack/state -> 更新对手预测/触发回滚

网络开发总结

模块

NetCodec

Input
ID,count, client(本端)最新tick, clientAckServerTick（server确认最新tick） cmd包
Ack
权威tick，最后tick 以及hash验证
Statae
权威状态包
Start
游戏开始通知

NetStub

废弃 tick 模拟状态发送和打印监控，状态发送已改为键盘输入

Packets

PacketType 状态，input ack state start
InputPacket 包冗余发松多个InputCmd 包以防丢包
start 包
对局状态信息

UdpSocket

udp 连接比tcp简单，持续监听端口

服务器分析

time：2026_1_20

ClientConn
历史缓存
tick
最后输入tick
最后权威tick
最后输入包
玩家数量

ServerCtx
基础信息
对局设置
帧同步设置
世界设置
sock 网络
clients 存储客户端状态
函数
AssignSlot 分配id 1，2
GetPlayer 从ServerCtx->clients 取出ClientConn
OnlineCount 在线玩家统计
GetCmdForTick 从ClientConn中获取动作信息，正常接受，如果出现延迟但小于延迟设置，着延续上一次包，但是出现过大延迟直接采用默认信息

OnUdp 收包加入缓存
MaybeStartMatch for pid 轮训发收start包，其中通过利用getplayer 吧serverctx中的clientconn取出，同时检测ctx中的状态已start则跳过

OnTick
    整体事件监听
    推进世界 + 给两边发 State/Ack

模拟

发表于 2026-01-19 分类于学习，实时竞技游戏开发

渲染

发表于 2026-01-19 分类于学习，实时竞技游戏开发

bufferevent

发表于 2026-01-07 分类于学习， Linux高性能服务器编程

bufferevent

1. 基本概念

bufferevent：在套接字之上封装了输入/输出缓冲与回调机制的 I/O 对象，避免手写 read/write 和边沿触发细节。
事件循环：event_base 驱动，所有 I/O、超时、信号都在此 loop 内调度。

2. 关键 API

buferevent_socket_new()

struct bufferevent* bufferevent_socket_new(
    struct event_base *base,
    evutil_socket_t fd,           // 通常传 -1，让 libevent 自己创建 socket
    int options                   // 常用：BEV_OPT_CLOSE_ON_FREE | BEV_OPT_DEFER_CALLBACKS
);

说明

fd = -1：由 libevent 负责创建 socket，更简洁。
BEV_OPT_CLOSE_ON_FREE：在 bufferevent_free() 时自动关闭 socket。
BEV_OPT_DEFER_CALLBACKS：把回调延迟到事件循环安全点执行，减少重入问题。

buferevent_socket_connect()

int bufferevent_socket_connect(
    struct bufferevent *bev,
    struct sockaddr *addr,
    int addrlen
);

说明

非阻塞发起连接；成功返回 0，失败返回 -1（同时触发 event 回调中的 BEV_EVENT_ERROR）。
连接成功后会触发 event 回调中的 BEV_EVENT_CONNECTED。

如果要用主机名而非 IP，推荐 bufferevent_socket_connect_hostname()（需配合 evdns_base 进行异步解析）。

3. 最小可用示例（IPv4）

#include <event2/event.h>
#include <event2/bufferevent.h>
#include <event2/buffer.h>
#include <event2/util.h>
#include <arpa/inet.h>
#include <string.h>
#include <stdio.h>

static void read_cb(struct bufferevent* bev, void* ctx) {
    char buf[4096];
    size_t n;
    while ((n = bufferevent_read(bev, buf, sizeof(buf))) > 0) {
        fwrite(buf, 1, n, stdout);
    }
}

static void event_cb(struct bufferevent* bev, short what, void* ctx) {
    if (what & BEV_EVENT_CONNECTED) {
        printf("[OK] connected\n");
        const char* msg = "GET / HTTP/1.1\r\nHost: 127.0.0.1\r\n\r\n";
        bufferevent_write(bev, msg, strlen(msg));
    }
    if (what & BEV_EVENT_EOF) {
        printf("[*] server closed\n");
        bufferevent_free(bev);
    }
    if (what & BEV_EVENT_ERROR) {
        int err = EVUTIL_SOCKET_ERROR();
        fprintf(stderr, "[ERR] %s\n", evutil_socket_error_to_string(err));
        bufferevent_free(bev);
    }
    if (what & BEV_EVENT_TIMEOUT) {
        printf("[!] timeout\n");
        bufferevent_free(bev);
    }
}

int main() {
    struct event_base* base = event_base_new();

    //创建套字 socket
    struct bufferevent* bev = bufferevent_socket_new(
        base, -1, BEV_OPT_CLOSE_ON_FREE | BEV_OPT_DEFER_CALLBACKS);
    // 将read_cb event_cb 回调函数先绑定防止漏听
    bufferevent_setcb(bev, read_cb, NULL, event_cb, NULL);
    bufferevent_enable(bev, EV_READ | EV_WRITE);

    // 可选：设置读/写超时（连接阶段通常走写超时）
    struct timeval rto = {10, 0}, wto = {10, 0};
    bufferevent_set_timeouts(bev, &rto, &wto);

    // 目标地址
    struct sockaddr_in sin;
    memset(&sin, 0, sizeof(sin));
    sin.sin_family = AF_INET;
    sin.sin_port   = htons(8080);
    evutil_inet_pton(AF_INET, "127.0.0.1", &sin.sin_addr);

    if (bufferevent_socket_connect(bev, (struct sockaddr*)&sin, sizeof(sin)) < 0) {
        perror("connect");
        bufferevent_free(bev);
        event_base_free(base);
        return 1;
    }

    event_base_dispatch(base);

    event_base_free(base);
    return 0;
}

4. 使用要点与常见坑

回调必须先设再 connect
先 bufferevent_setcb() + bufferevent_enable()，再调用 bufferevent_socket_connect()，避免早到事件丢失。
超时设置
- bufferevent_set_timeouts() 同时设置读/写超时。
- 连接阶段通常依赖写超时触发（因为 connect 过程中会等待可写）。
水位（可选）
- 读水位：bufferevent_setwatermark(bev, EV_READ, low, high)
  当输入缓冲超过 high 时会暂时禁用读事件，防止内存暴涨。
- 写水位：同理控制输出缓冲积压。
线程安全
- libevent 可多线程，但一个 event_base 通常只在一个线程内跑；跨线程提交任务请用 event_base_once() 或管道/通知机制。
DNS 支持（推荐）
- 异步解析：evdns_base_new(base, 1) + bufferevent_socket_connect_hostname(bev, dns, AF_UNSPEC, "example.com", 80)
- 避免阻塞式 getaddrinfo()。
清理顺序
- 一般在 event_cb 收到 EOF/ERROR/TIMEOUT 后 bufferevent_free()；loop 退出后 event_base_free()。
错误定位
- BEV_EVENT_ERROR 时使用 EVUTIL_SOCKET_ERROR() + evutil_socket_error_to_string() 打印 errno。
- 常见：ECONNREFUSED（服务未起）、ETIMEDOUT（网络/防火墙/超时）、EHOSTUNREACH。
选项建议
- BEV_OPT_CLOSE_ON_FREE：免漏句柄
- BEV_OPT_DEFER_CALLBACKS：减少重入与回调嵌套
- 如需逐线程回调可考虑 BEV_OPT_THREADSAFE（需要启用 evthread_use_pthreads() 等）。

5. 主机名连接（带 DNS）的示例（简版）

#include <event2/dns.h>

struct evdns_base* dns = evdns_base_new(base, 1); // 1=初始化内置解析器
int rc = bufferevent_socket_connect_hostname(
    bev, dns, AF_UNSPEC, "example.com", 80);
if (rc < 0) { /* handle error */ }

6. 小结（流程速记）

event_base_new()
bufferevent_socket_new(base, -1, BEV_OPT_CLOSE_ON_FREE | BEV_OPT_DEFER_CALLBACKS)
bufferevent_setcb() → bufferevent_enable(EV_READ|EV_WRITE) → bufferevent_set_timeouts()
bufferevent_socket_connect()（或 _connect_hostname()）
在 event_cb 处理 BEV_EVENT_CONNECTED / EOF / ERROR / TIMEOUT
退出时 bufferevent_free()、event_base_free()

7. 进阶建议

TLS：用 bufferevent_openssl_socket_new() 包装 SSL 会话，接口与普通 bev 基本一致。
背压控制：配合水位、分块写入，避免写缓冲过大。
协议层：在 read_cb 内从 evbuffer 取数据，做半包/粘包处理（如定长头 + 可变体）。

libevent 笔记

发表于 2026-01-07 分类于学习， Linux高性能服务器编程

libevent 笔记

TCP evconnlistener

evconnlistener机制替代了之前socket bind listen 监听套字的创立

创建evconnlistener 接口

struct evconnlistener *evconnlistener_new(struct event_base *base,
    evconnlistener_cb cb, void *ptr, unsigned flags, int backlog,
    evutil_socket_t fd);
struct evconnlistener *evconnlistener_new_bind(struct event_base *base,
    evconnlistener_cb cb, void *ptr, unsigned flags, int backlog,
    const struct sockaddr *sa, int socklen);
void evconnlistener_free(struct evconnlistener *lev);

evconnlistener_new和evconnlistener_new_bind都用于分配和返回一个新的用于监听连接的对象。每个用于监听的对象都使用event_base来维护一个建立的连接。当有新的TCP 连接发生时，回调函数就会调用.

base： event_base对象。参考event_base。
cb：回调函数，当有新的TCP连接发生时，会唤醒回调函数。
ptr：传递给回调函数的参数。
flags：一些标志，后面会进一步介绍。
backlog：监听队列允许容纳的最大连接数。
fd：函数evconnlistener_new假设我们已经绑定了套接字到要监听的端口上，fd参数就是我们已经绑定的socket套接字。
sa： evconnlistener_new_bind帮助我们绑定监听地址。sa就是传入的监听地址。
socklen： sa的长度。

一般通常情况下flags设置LEV_OPT_CLOSE_ON_FREE，LEV_OPT_REUSEABLE

回调函数

1 2	typedef void (evconnlistener_cb)(struct evconnlistener listener, evutil_socket_t sock, struct sockaddr addr, int len, void ptr);

监听开启和关闭

1 2	int evconnlistener_disable(struct evconnlistener lev); int evconnlistener_enable(struct evconnlistener lev);

从enconnlistener 中获取信息

// 获得一个evconnlistener对象的套接字
evutil_socket_t evconnlistener_get_fd(struct evconnlistener *lev);
// 获得一个evconnlistener对象event_base对象
evutil_socket_t evconnlistener_get_base(struct evconnlistener *lev);

完整示范

#include <event2/listener.h>
#include <event2/bufferevent.h>
#include <event2/buffer.h>

#include <arpa/inet.h>

#include <string.h>
#include <stdlib.h>
#include <stdio.h>
#include <errno.h>

static void
echo_read_cb(struct bufferevent *bev, void *ctx)
{
        /*获取读入读出缓冲区*/
        struct evbuffer *input = bufferevent_get_input(bev);
        struct evbuffer *output = bufferevent_get_output(bev);

        /* 将输入缓冲区的内容添加到输出缓冲区，实现回声功能 */
        evbuffer_add_buffer(output, input);
}

static void
echo_event_cb(struct bufferevent *bev, short events, void *ctx)
{
        if (events & BEV_EVENT_ERROR)
                perror("Error from bufferevent");
        if (events & (BEV_EVENT_EOF | BEV_EVENT_ERROR)) {
                bufferevent_free(bev);
        }
}

static void
accept_conn_cb(struct evconnlistener *listener,
    evutil_socket_t fd, struct sockaddr *address, int socklen,
    void *ctx)
{
        struct event_base *base = evconnlistener_get_base(listener);//获取事件
        struct bufferevent *bev = bufferevent_socket_new(
                base, fd, BEV_OPT_CLOSE_ON_FREE);//建立连接

        bufferevent_setcb(bev, echo_read_cb, NULL, echo_event_cb, NULL);//设置read和错误回调函数，传入参数为0

        bufferevent_enable(bev, EV_READ|EV_WRITE);//设置出发条件，EV_READ 读入时触发，EV_WRITE 错误是触发
}

static void
accept_error_cb(struct evconnlistener *listener, void *ctx)
{
        struct event_base *base = evconnlistener_get_base(listener);
        int err = EVUTIL_SOCKET_ERROR();
        fprintf(stderr, "Got an error %d (%s) on the listener. "
                "Shutting down.\n", err, evutil_socket_error_to_string(err));

        event_base_loopexit(base, NULL);
}

int
main(int argc, char **argv)
{
        struct event_base *base;
        struct evconnlistener *listener;
        struct sockaddr_in sin;

        int port = 9876;

        if (argc > 1) {
                port = atoi(argv[1]);
        }
        if (port<=0 || port>65535) {
                puts("Invalid port");
                return 1;
        }

        base = event_base_new();
        if (!base) {
                puts("Couldn't open event base");
                return 1;
        }

        memset(&sin, 0, sizeof(sin));

        sin.sin_family = AF_INET;
        sin.sin_addr.s_addr = htonl(0);
        sin.sin_port = htons(port);

        listener = evconnlistener_new_bind(base, accept_conn_cb, NULL,
            LEV_OPT_CLOSE_ON_FREE|LEV_OPT_REUSEABLE, -1,
            (struct sockaddr*)&sin, sizeof(sin));
        if (!listener) {
                perror("Couldn't create listener");
                return 1;
        }
        evconnlistener_set_error_cb(listener, accept_error_cb);

        event_base_dispatch(base);
        return 0;
}

bufferevents

bufferevent 由一个底层传输系统（比如socket），一个读缓冲区和一个写缓冲区组成

bufferevent和evbuffers

每一个bufferevent 都有一个输入缓冲区和输出缓冲区，这些缓冲区（buffer）都是struct evbuffer 类型

基于socket 的bufferevent

bufferevent_socket_new()
struct bufferevent *bufferevent_socket_new(
struct event_base *base,
evutil_socket_t fd,
enum bufferevent_options options);

正常通过ip host进行访问后续可以使用socket进行构建client客户端连接简化实现
int bufferevent_socket_connect(struct bufferevent *bev,
struct sockaddr *address, int addrlen);
通过域名进行解析访问
int bufferevent_socket_connect_hostname(struct bufferevent *bev,
struct evdns_base *dns_base, int family, const char *hostname,
int port);
int bufferevent_socket_get_dns_error(struct bufferevent *bev);

注：
bufferevent_socket_connect和evconnlistener_new_bind的区别
bufferevent_socket_connect：用于主动建立连接，即作为客户端去连接服务器。
evconnlistener_new_bind：用于被动接受连接，即作为服务器监听端口并接受客户端的连接。

回调函数 & 水位线操作

水位线设置
void bufferevent_setwatermark(struct bufferevent *bufev, short events,
size_t lowmark, size_t highmark);
events参数传入决定设置不同的水危险
EV_WRITE、 EV_READ

读、写、报错回调函数调用
void bufferevent_setcb(struct bufferevent *bufev,
bufferevent_data_cb readcb, bufferevent_data_cb writecb,
bufferevent_event_cb eventcb, void *cbarg);

#include <event2/event.h>
#include <event2/bufferevent.h>
#include <event2/buffer.h>
#include <event2/util.h>

#include <stdlib.h>
#include <errno.h>
#include <string.h>

struct info {
    const char *name;
    size_t total_drained;
};

void read_callback(struct bufferevent *bev, void *ctx)
{
    struct info *inf = ctx;
    struct evbuffer *input = bufferevent_get_input(bev);
    size_t len = evbuffer_get_length(input);
    if (len) {
        inf->total_drained += len;
        evbuffer_drain(input, len);// 和evbuffer_remove()操作类似，但不执行copy工作仅删除长度len的缓冲
        printf("Drained %lu bytes from %s\n",
             (unsigned long) len, inf->name);
    }
}

void event_callback(struct bufferevent *bev, short events, void *ctx)
{
    struct info *inf = ctx;
    struct evbuffer *input = bufferevent_get_input(bev);
    int finished = 0;

    if (events & BEV_EVENT_EOF) {
        size_t len = evbuffer_get_length(input);
        printf("Got a close from %s.  We drained %lu bytes from it, "
            "and have %lu left.\n", inf->name,
            (unsigned long)inf->total_drained, (unsigned long)len);
        finished = 1;
    }
    if (events & BEV_EVENT_ERROR) {
        printf("Got an error from %s: %s\n",
            inf->name, evutil_socket_error_to_string(EVUTIL_SOCKET_ERROR()));
        finished = 1;
    }
    if (finished) {
        free(ctx);
        bufferevent_free(bev);
    }
}

struct bufferevent *setup_bufferevent(void)
{
    struct bufferevent *b1 = NULL;
    struct info *info1;

    info1 = malloc(sizeof(struct info));
    info1->name = "buffer 1";
    info1->total_drained = 0;

    /* ... Here we should set up the bufferevent and make sure it gets
       connected... */

    /* Trigger the read callback only whenever there is at least 128 bytes
       of data in the buffer. */
    bufferevent_setwatermark(b1, EV_READ, 128, 0);

    bufferevent_setcb(b1, read_callback, NULL, event_callback, info1);

    bufferevent_enable(b1, EV_READ); /* Start reading. */
    return b1;
}

bufferevent中的数据操作

获取不同缓冲区的evbuffer描述符
struct evbuffer *bufferevent_get_input(struct bufferevent *bufev);
struct evbuffer *bufferevent_get_output(struct bufferevent *bufev);

写入读入操作
int bufferevent_write(struct bufferevent bufev, const voiddata, size_t size);
int bufferevent_write_buffer(struct bufferevent bufev, struct evbufferbuf);

size_t bufferevent_read(struct bufferevent *bufev, void *data, size_t size);
int bufferevent_read_buffer(struct bufferevent *bufev,
struct evbuffer *buf);

程序自我通信

int bufferevent_pair_new(struct event_base *base, int options,
struct bufferevent *pair[2]);

pair[0]和pair[1] 设置为一对相互连接的bufferevent

evbuffer缓冲IO的实用功能

创建和释放
struct evbuffer *evbuffer_new(void);
void evbuffer_free(struct evbuffer *buf);

CUDA 开启的 GPU 编程

发表于 2026-01-04 分类于学习，高性能C++并行编程

CUDA 开启的 GPU 编程

time：2025_12_26

1. 工程与编译（CMake / nvcc）

1.1 最小 CMake 工程

cmake_minimum_required(VERSION 3.18)
project(hellocuda LANGUAGES CXX CUDA)

add_executable(main main.cu)

set_target_properties(main PROPERTIES
  CUDA_STANDARD 17
  CUDA_STANDARD_REQUIRED ON
)

# 需要跨 .cu 调用 device 函数/做 device link 时开启
set_target_properties(main PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

# 常用 CUDA 编译选项（按需启用）
target_compile_options(main PUBLIC
  $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
  $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
)

1.2 推荐的基础错误检查宏

避免依赖 sample 的 helper_cuda.h，直接自带一个最小版：

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

#define CUDA_CHECK(call) do {                                  \
  cudaError_t err = (call);                                    \
  if (err != cudaSuccess) {                                    \
    std::fprintf(stderr, "CUDA error %s:%d: %s\n",             \
      __FILE__, __LINE__, cudaGetErrorString(err));            \
    std::exit(1);                                              \
  }                                                            \
} while(0)

2. CUDA 基础：函数修饰符与执行位置

2.1 `global / device / host`

__global__：核函数（kernel）
- 在 GPU 上并行执行
- 由主机端（CPU）发起 <<<...>>>
- 返回类型必须为 void（通常通过指针写回结果）
__device__：设备函数
- 在 GPU 上执行
- 只能从 device/global 调用
__host__：主机函数
- 在 CPU 上执行
- 未标注的普通函数默认就是 host

组合：

__host__ __device__：同一函数在 CPU/GPU 两侧都可用（注意 device 侧不支持完整的 C++ 标准库能力）

2.2 `__CUDA_ARCH__`（区分 device/host 编译路径）

__CUDA_ARCH__ 只在 device 编译路径中定义，值为计算能力架构号（例如 750/800 等）
常用于同一个函数在 host/device 的条件编译

__host__ __device__ inline int where_am_i() {
#ifdef __CUDA_ARCH__
    return __CUDA_ARCH__;   // device：架构号
#else
    return -1;              // host：标记
#endif
}

2.3 `constexpr` 与 device 代码

如需更宽松的 constexpr 在 device 侧工作，常用 --expt-relaxed-constexpr
device 侧 lambda 扩展常用 --expt-extended-lambda

3. Kernel 启动、线程块模型与索引

3.1 Kernel 启动语法

1	kernel<<<grid, block, shared_bytes, stream>>>(args...);

常用：<<<grid, block>>>
shared_bytes：动态共享内存字节数（默认 0）
stream：CUDA 流（默认 0）

3.2 线程/块索引

threadIdx.{x,y,z}：线程在块内索引
blockIdx.{x,y,z}：块在网格内索引
blockDim.{x,y,z}：每块线程数维度
gridDim.{x,y,z}：网格块数维度

3.3 典型打印示例（便于理解执行模型）

#include <cstdio>
#include <cuda_runtime.h>

__global__ void kernel() {
    printf("block %d/%d, thread %d/%d\n",
           blockIdx.x, gridDim.x,
           threadIdx.x, blockDim.x);
}

int main() {
    kernel<<<2, 3>>>();
    cudaDeviceSynchronize();
    return 0;
}

3.4 Grid-Stride Loop（通用遍历范式）

适用于任意大小数据与任意 grid/block 配置：

__global__ void work(int* a, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n;
         i += blockDim.x * gridDim.x) {
        a[i] = i;
    }
}

4. 同步与错误处理

4.1 CPU/GPU 默认异步

kernel launch 对 host 来说通常是异步的
常用同步：
- cudaDeviceSynchronize()：等待当前设备上所有已提交工作完成
- cudaStreamSynchronize(stream)：等待某个流完成

4.2 推荐的 launch 后检查模板

1
2
3

kernel<<<grid, block>>>(...);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

5. 内存管理（Host/Device/Unified）

5.1 经典模式：`cudaMalloc + cudaMemcpy`

#include <cstdio>
#include <cuda_runtime.h>

#define CUDA_CHECK(call) do { \
  cudaError_t err = (call); \
  if (err != cudaSuccess) { \
    std::fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err)); \
    std::exit(1); \
  } \
} while(0)

__global__ void kernel(int* out) { *out = 42; }

int main() {
    int* d_out = nullptr;
    CUDA_CHECK(cudaMalloc(&d_out, sizeof(int)));

    kernel<<<1,1>>>(d_out);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    int h_out = 0;
    CUDA_CHECK(cudaMemcpy(&h_out, d_out, sizeof(int), cudaMemcpyDeviceToHost));

    std::printf("ret=%d\n", h_out);

    CUDA_CHECK(cudaFree(d_out));
    return 0;
}

5.2 统一内存（Unified Memory）：`cudaMallocManaged`

一份指针同时可被 CPU/GPU 访问
常配合同步；性能敏感时可用预取提升稳定性

#include <cstdio>
#include <cuda_runtime.h>

__global__ void fill(int* a, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        a[i] = i;
    }
}

int main() {
    int n = 32;
    int* a = nullptr;
    cudaMallocManaged(&a, sizeof(int) * n);

    fill<<<1, 128>>>(a, n);
    cudaDeviceSynchronize();

    for (int i = 0; i < n; ++i) std::printf("a[%d]=%d\n", i, a[i]);
    cudaFree(a);
    return 0;
}

5.3 预取（Prefetch）与驻留优化（进阶但常用）

int dev = 0;
cudaGetDevice(&dev);
cudaMemPrefetchAsync(a, sizeof(int)*n, dev, 0);     // 预取到 GPU
// kernel ...
cudaMemPrefetchAsync(a, sizeof(int)*n, cudaCpuDeviceId, 0); // 预取回 CPU
cudaDeviceSynchronize();

6. C++ 封装：RAII 与可复用接口

6.1 RAII 管理 Unified Memory 指针（简单、可靠、适合入门）

#include <cuda_runtime.h>
#include <cstddef>

template<class T>
struct ManagedArray {
    T* p{nullptr};
    size_t n{0};

    explicit ManagedArray(size_t n_) : n(n_) {
        cudaMallocManaged(&p, sizeof(T) * n);
    }
    ~ManagedArray() {
        if (p) cudaFree(p);
    }
    T* data() { return p; }
    const T* data() const { return p; }
    T& operator[](size_t i) { return p[i]; }
    const T& operator[](size_t i) const { return p[i]; }
};

6.2 结合 kernel 使用

__global__ void init(int* a, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        a[i] = i;
    }
}

int main() {
    int n = 1000;
    ManagedArray<int> a(n);
    init<<<32, 128>>>(a.data(), n);
    cudaDeviceSynchronize();
    // CPU 侧直接读
    return 0;
}

allocator 方式也可把 unified memory 接入 std::vector，但 allocator 细节较多；建议先把 RAII 指针与 .data() 传参掌握牢。

7. Thrust 库：容器与算法（高层 CUDA）

7.1 常用容器

thrust::host_vector<T>：主机端 vector
thrust::device_vector<T>：设备端 vector
通过赋值可触发 H2D / D2H 拷贝（更准确地说：构造/赋值会在 host/device 容器之间进行数据迁移）

7.2 AXPY 示例（device_vector + 自写 kernel）

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <cuda_runtime.h>
#include <cstdio>

__global__ void axpy(float* x, const float* y, float a, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        x[i] = a * x[i] + y[i];
    }
}

int main() {
    int n = 1 << 20;
    float a = 3.14f;

    thrust::host_vector<float> hx(n), hy(n);
    for (int i = 0; i < n; ++i) { hx[i] = i * 0.001f; hy[i] = 1.0f; }

    thrust::device_vector<float> dx = hx;
    thrust::device_vector<float> dy = hy;

    axpy<<<256, 256>>>(thrust::raw_pointer_cast(dx.data()),
                       thrust::raw_pointer_cast(dy.data()),
                       a, n);
    cudaDeviceSynchronize();

    hx = dx;
    std::printf("hx[0]=%f, hx[n-1]=%f\n", hx[0], hx[n-1]);
    return 0;
}

8. 原子操作（Atomic）

8.1 常用原子

atomicAdd / atomicSub
atomicAnd / atomicOr / atomicXor
atomicMin / atomicMax
atomicCAS：Compare-And-Swap，可用于构造自定义原子操作

8.2 用 CAS 实现自定义原子加

__device__ __forceinline__ int my_atomic_add(int* dst, int val) {
    int old = *dst;
    int assumed;
    do {
        assumed = old;
        old = atomicCAS(dst, assumed, assumed + val);
    } while (assumed != old);
    return old;
}

8.3 朴素并行求和（全局原子累加）

__global__ void parallel_sum(int* sum, const int* arr, int n) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        atomicAdd(sum, arr[i]);
    }
}

9. 线程块与共享内存（Shared Memory）

9.1 核心概念

__shared__：块内共享内存（一个 block 内所有线程可见）
__syncthreads()：块内同步屏障（必须保证同一 block 的线程都能到达）

共享内存常用于：

块内复用数据（减少 global memory 访问）
块内归约（reduce）
tile-based 计算（矩阵乘、卷积、图像算子）

9.2 块内归约：每块只做一次全局原子

#include <cuda_runtime.h>

__global__ void reduce_sum(const int* arr, int n, int* out) {
    extern __shared__ int sdata[]; // 动态共享内存
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + tid;

    sdata[tid] = (i < n) ? arr[i] : 0;
    __syncthreads();

    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) sdata[tid] += sdata[tid + s];
        __syncthreads();
    }

    if (tid == 0) atomicAdd(out, sdata[0]);
}

启动方式（动态共享内存大小）：

1
2
3

int threads = 256;
int blocks = (n + threads - 1) / threads;
reduce_sum<<<blocks, threads, threads * sizeof(int)>>>(arr, n, out);

9.3 Tile 示例：2D 图像 3x3 均值滤波（共享内存加速范式）

适用于图像/矩阵类任务（tile + halo）：

#include <cuda_runtime.h>

__global__ void mean3x3(const float* in, float* out, int H, int W) {
    // blockDim = (Bx, By)
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;

    // tile 尺寸：块大小 + halo(上下左右各1)
    const int Bx = blockDim.x;
    const int By = blockDim.y;
    extern __shared__ float tile[];

    // tile 索引函数
    auto t = [&](int ty, int tx) -> float& {
        return tile[ty * (Bx + 2) + tx];
    };

    // 对应 tile 坐标（+1 是为了留 halo）
    const int tx = threadIdx.x + 1;
    const int ty = threadIdx.y + 1;

    // 读主区域
    float v = 0.f;
    if (x < W && y < H) v = in[y * W + x];
    t(ty, tx) = v;

    // 读 halo（边界处做 clamp 或置零，这里用置零策略）
    if (threadIdx.x == 0) {
        float lv = (x > 0 && y < H) ? in[y * W + (x - 1)] : 0.f;
        t(ty, 0) = lv;
    }
    if (threadIdx.x == Bx - 1) {
        float rv = (x + 1 < W && y < H) ? in[y * W + (x + 1)] : 0.f;
        t(ty, Bx + 1) = rv;
    }
    if (threadIdx.y == 0) {
        float uv = (y > 0 && x < W) ? in[(y - 1) * W + x] : 0.f;
        t(0, tx) = uv;
    }
    if (threadIdx.y == By - 1) {
        float dv = (y + 1 < H && x < W) ? in[(y + 1) * W + x] : 0.f;
        t(By + 1, tx) = dv;
    }

    // 角落 halo（四个角）
    if (threadIdx.x == 0 && threadIdx.y == 0) {
        t(0,0) = (x>0 && y>0) ? in[(y-1)*W + (x-1)] : 0.f;
    }
    if (threadIdx.x == Bx-1 && threadIdx.y == 0) {
        t(0,Bx+1) = (x+1<W && y>0) ? in[(y-1)*W + (x+1)] : 0.f;
    }
    if (threadIdx.x == 0 && threadIdx.y == By-1) {
        t(By+1,0) = (x>0 && y+1<H) ? in[(y+1)*W + (x-1)] : 0.f;
    }
    if (threadIdx.x == Bx-1 && threadIdx.y == By-1) {
        t(By+1,Bx+1) = (x+1<W && y+1<H) ? in[(y+1)*W + (x+1)] : 0.f;
    }

    __syncthreads();

    if (x < W && y < H) {
        float sum = 0.f;
        sum += t(ty-1, tx-1); sum += t(ty-1, tx); sum += t(ty-1, tx+1);
        sum += t(ty,   tx-1); sum += t(ty,   tx); sum += t(ty,   tx+1);
        sum += t(ty+1, tx-1); sum += t(ty+1, tx); sum += t(ty+1, tx+1);
        out[y * W + x] = sum / 9.f;
    }
}

共享内存大小（字节）：

dim3 block(16, 16);
dim3 grid((W + block.x - 1) / block.x, (H + block.y - 1) / block.y);
size_t shared_bytes = (block.x + 2) * (block.y + 2) * sizeof(float);
mean3x3<<<grid, block, shared_bytes>>>(in, out, H, W);

10. CUDA Streams 与异步拷贝

10.1 为什么需要 streams

默认 stream（stream 0）会形成较强的串行依赖
多 stream 可以实现：
- H2D 拷贝与 kernel 重叠
- 多批次流水线（pipeline）
- 与 cudaMemcpyAsync 配合提升吞吐

10.2 pinned（页锁定）主机内存：提升异步拷贝效率

cudaMallocHost / cudaFreeHost
pinned 内存更利于 DMA，cudaMemcpyAsync 才更有意义

10.3 基本模板：两条 stream 流水搬运

#include <cstdio>
#include <cuda_runtime.h>

#define CUDA_CHECK(call) do {                                  \
  cudaError_t err = (call);                                    \
  if (err != cudaSuccess) {                                    \
    std::fprintf(stderr, "CUDA error %s:%d: %s\n",             \
      __FILE__, __LINE__, cudaGetErrorString(err));            \
    std::exit(1);                                              \
  }                                                            \
} while(0)

__global__ void scale(float* x, int n, float a) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < n; i += blockDim.x * gridDim.x) {
        x[i] *= a;
    }
}

int main() {
    const int n = 1 << 20;
    const size_t bytes = n * sizeof(float);

    // pinned host memory
    float* h = nullptr;
    CUDA_CHECK(cudaMallocHost(&h, bytes));

    // device memory
    float* d = nullptr;
    CUDA_CHECK(cudaMalloc(&d, bytes));

    // init host
    for (int i = 0; i < n; ++i) h[i] = 1.0f;

    cudaStream_t s;
    CUDA_CHECK(cudaStreamCreate(&s));

    // async H2D
    CUDA_CHECK(cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice, s));

    // kernel in same stream (will wait for H2D in this stream)
    scale<<<256, 256, 0, s>>>(d, n, 2.0f);
    CUDA_CHECK(cudaGetLastError());

    // async D2H
    CUDA_CHECK(cudaMemcpyAsync(h, d, bytes, cudaMemcpyDeviceToHost, s));

    // wait stream done
    CUDA_CHECK(cudaStreamSynchronize(s));

    std::printf("h[0]=%f, h[n-1]=%f\n", h[0], h[n-1]);

    CUDA_CHECK(cudaStreamDestroy(s));
    CUDA_CHECK(cudaFree(d));
    CUDA_CHECK(cudaFreeHost(h));
    return 0;
}

10.4 事件计时（event timing）

用于测量 GPU 端耗时：

cudaEvent_t st, ed;
cudaEventCreate(&st); cudaEventCreate(&ed);

cudaEventRecord(st, stream);
// kernel / memcpyAsync ...
cudaEventRecord(ed, stream);
cudaEventSynchronize(ed);

float ms = 0.f;
cudaEventElapsedTime(&ms, st, ed);

cmake环境构建注意事项

发表于 2025-12-25 更新于 2026-01-24 分类于杂七杂八

cmake环境构建注意事项

后续要装更多包：通用“镜像下载 + 安装模板”

你以后新增包基本分三类：CMake / Autotools / 纯 Make。建议统一规则：

下载：都落到 /opt/tp-downloads

解压构建：都在 /opt/src_clean/

安装：统一 –prefix=/opt/tp 或 -DCMAKE_INSTALL_PREFIX=/opt/tp

使用：工程端统一 CMAKE_PREFIX_PATH=/opt/tp

7.1 通用下载函数（带镜像）

建议你写个小函数（放到 ~/.bashrc 或单独脚本）：

fetch_gh () {

用法：fetch_gh owner repo ref output.tar.gz

例：fetch_gh sewenew redis-plus-plus refs/tags/1.3.15 redis-plus-plus-1.3.15.tar.gz

local owner=”$1” repo=”$2” ref=”$3” out=”$4”
mkdir -p /opt/tp-downloads
curl -L –fail
“https://gh-proxy.org/https://github.com/${owner}/${repo}/archive/${ref}.tar.gz“
-o “/opt/tp-downloads/${out}”
}

7.2 CMake 类包安装模板
source /opt/tp/env.sh
cd /opt/src_clean/
rm -rf build && mkdir build && cd build
/opt/tp/bin/cmake .. -G Ninja
-DCMAKE_BUILD_TYPE=Release
-DCMAKE_INSTALL_PREFIX=/opt/tp
-DCMAKE_PREFIX_PATH=/opt/tp
/opt/tp/bin/cmake –build . -j”$(nproc)”
/opt/tp/bin/cmake –install .

7.3 Autotools 类包安装模板
./configure –prefix=/opt/tp
make -j”$(nproc)”
make install

下载 redis-plus-plus（走 GitHub 镜像）

redis-plus-plus 最新 release/tag 是 1.3.15。
GitHub

mkdir -p /opt/tp-downloads
cd /opt/tp-downloads

curl -L –fail
“https://gh-proxy.org/https://github.com/sewenew/redis-plus-plus/archive/refs/tags/1.3.15.tar.gz“
-o redis-plus-plus-1.3.15.tar.gz

解压、编译、安装到 /opt/tp

关键 CMake 选项：

-DREDIS_PLUS_PLUS_BUILD_TEST=OFF：不编测试，加速且避免额外依赖
GitHub
+1

-DCMAKE_PREFIX_PATH=/opt/tp：让它找到你已装的 hiredis（以及未来其他依赖）

source /opt/tp/env.sh

cd /opt/src_clean
rm -rf redis-plus-plus-1.3.15
tar -xf /opt/tp-downloads/redis-plus-plus-1.3.15.tar.gz
cd redis-plus-plus-1.3.15

rm -rf build
mkdir -p build
cd build

/opt/tp/bin/cmake .. -G Ninja
-DCMAKE_BUILD_TYPE=Release
-DCMAKE_INSTALL_PREFIX=/opt/tp
-DCMAKE_PREFIX_PATH=/opt/tp
-DREDIS_PLUS_PLUS_BUILD_TEST=OFF

/opt/tp/bin/cmake –build . -j”$(nproc)”
/opt/tp/bin/cmake –install .

Quick Start

Create a new post

Run server

Generate static files

Deploy to remote sites

Client 实现分析

app 依赖包

主体部分

整理

网络开发总结

模块

NetCodec

NetStub

Packets

UdpSocket

服务器分析

bufferevent

1. 基本概念

2. 关键 API

buferevent_socket_new()

buferevent_socket_connect()

3. 最小可用示例（IPv4）

4. 使用要点与常见坑

5. 主机名连接（带 DNS）的示例（简版）

6. 小结（流程速记）

7. 进阶建议

libevent 笔记

TCP evconnlistener

创建evconnlistener 接口

回调函数

监听开启和关闭

从enconnlistener 中获取信息

完整示范

bufferevents

bufferevent和evbuffers

基于socket 的bufferevent

回调函数 & 水位线操作

bufferevent中的数据操作

程序自我通信

evbuffer缓冲IO的实用功能

CUDA 开启的 GPU 编程

1. 工程与编译（CMake / nvcc）

1.1 最小 CMake 工程

1.2 推荐的基础错误检查宏

2. CUDA 基础：函数修饰符与执行位置

2.1 __global__ / __device__ / __host__

2.2 __CUDA_ARCH__（区分 device/host 编译路径）

2.3 constexpr 与 device 代码

3. Kernel 启动、线程块模型与索引

3.1 Kernel 启动语法

3.2 线程/块索引

3.3 典型打印示例（便于理解执行模型）

3.4 Grid-Stride Loop（通用遍历范式）

4. 同步与错误处理

4.1 CPU/GPU 默认异步

4.2 推荐的 launch 后检查模板

5. 内存管理（Host/Device/Unified）

5.1 经典模式：cudaMalloc + cudaMemcpy

5.2 统一内存（Unified Memory）：cudaMallocManaged

5.3 预取（Prefetch）与驻留优化（进阶但常用）

6. C++ 封装：RAII 与可复用接口

6.1 RAII 管理 Unified Memory 指针（简单、可靠、适合入门）

6.2 结合 kernel 使用

7. Thrust 库：容器与算法（高层 CUDA）

7.1 常用容器

7.2 AXPY 示例（device_vector + 自写 kernel）

8. 原子操作（Atomic）

8.1 常用原子

8.2 用 CAS 实现自定义原子加

8.3 朴素并行求和（全局原子累加）

9. 线程块与共享内存（Shared Memory）

9.1 核心概念

9.2 块内归约：每块只做一次全局原子

9.3 Tile 示例：2D 图像 3x3 均值滤波（共享内存加速范式）

10. CUDA Streams 与异步拷贝

10.1 为什么需要 streams

10.2 pinned（页锁定）主机内存：提升异步拷贝效率

10.3 基本模板：两条 stream 流水搬运

10.4 事件计时（event timing）

cmake环境构建注意事项

2.1 `global / device / host`

2.2 `__CUDA_ARCH__`（区分 device/host 编译路径）

2.3 `constexpr` 与 device 代码

5.1 经典模式：`cudaMalloc + cudaMemcpy`

5.2 统一内存（Unified Memory）：`cudaMallocManaged`