go中len和Count的区别

最近使用go lang, 用着特别爽。但最近开发一个模块,一直不符合预期。后面通过看go源码,发现趟着大坑了。

len或者Count都能获取字符串、字节数组长度。如

data := “hello”

data_len := len(data)

data_len := strings.Count(data, “”) – 1

content := []byte(“hello”)

content_len := len(content)

content_len := bytes.Count(content, nil) – 1

但是len和Count还是有区别。以bytes为例,使用bytes.Count(data, nil) – 1获取字节数组data长度,长度值为utf8.RuneCount(data) + 1,返回的utf-8编码的长度;而len返回实际字节长度。字符串类似。

c++ virtual table

class Base {

public:
	Base(int a): a_(a) {}

	virtual void f() { cout << "Base::f" << endl; }

	virtual void g() { cout << "Base::g" << endl; }

	virtual void h() { cout << "Base::h" << endl; }
private:
	int a_;
};

int main() {
	// your code goes here
	typedef void(*Fun)(void);

	Base b(5);

	Fun pFun = NULL;
	
	cout << "a_:" << *((int*)&b + 2) << endl;

	cout << "虚函数表地址:" << (int*)(&b) << endl;

	int* vtable = (int*)*(int*)(&b);
	cout << "虚函数表 — 第一个函数地址:" << vtable << endl;

	// Invoke the first virtual function

	//pFun = (Fun)*((int*)*(int*)(&b));
	pFun = (Fun)*(vtable);
	pFun();
	pFun = (Fun)*((int*)*(int*)(&b) + 2);
	pFun();
	pFun = (Fun)*(vtable + 4);
	pFun();
        cout << "vtable end: " << (Fun)*(vtable + 6) << std::endl;

       // vtable and end of vtable
       int** pVtab = (int**)&b;
       pFun = (Fun)pVtab[0][0];
       pFun();  // base::f
       pFun = (Fun)pVtab[0][1];
       cout << pFun << endl;
       pFun = (Fun)pVtab[0][2];
       pFun();  // base::g
       pFun = (Fun)pVtab[0][3];
       cout << pFun << endl;
       pFun = (Fun)pVtab[0][4];
       pFun();  // base::h
       pFun = (Fun)pVtab[0][5];
       cout << pFun << endl;
       pFun = (Fun)pVtab[0][6];
       cout << pFun << endl;
       pFun = (Fun)pVtab[0][7];
       cout << pFun << endl;


	return 0;
}
输出结果:

a_:5
虚函数表地址:0x7ffe3f0a5970
虚函数表 ? 第一个函数地址:0x400e70
Base::f
Base::g
Base::h
vtable end: 1
Base::f
0
Base::g
0
Base::h
0
1
1

类static_cast、dynamic_cast与RTTI

1、static_cast

class Base {
    public:
        virtual void f() {
            cout << "Base::f() " << endl;
        }

};

class Derive: public Base {
    public:
        virtual void f() {
            cout << "Derive::f() " << endl;
        }

        virtual void f2() {
            cout << "Derive::f2() " << endl;
        }
};

int main(){
    // static_cast
    Base *pb1 = new Derive();
    Derive *pd1 = static_cast<Derive*>(pb1);
    pd1 -> f();

    Base *pb2 = new Base();
    Derive *pd2 = static_cast<Derive*>(pb2);
    pd2 -> f();
    //pd2 -> f2();  // core, base no f2()

    delete pb1;
    delete pb2;
    return 0;
}

static_cast可以在基类和派生类之间转换(偏移指针),编译时确定,不保证类型转换安全。

2、dynamic_cast

class Base {
    public:
        virtual void f() {
            cout << "Base::f() " << endl;
        }

};

class Derive: public Base {
    public:
        virtual void f() {
            cout << "Derive::f() " << endl;
        }

        virtual void f2() {
            cout << "Derive::f2() " << endl;
        }
};

int main(){
// dynamic_cast
    Base *pb3 = new Derive();
    Derive *pd3 = dynamic_cast<Derive*>(pb3);   // down cast
    pd3 -> f();

    Base *pb4 = new Base();
    Derive *pd4 = dynamic_cast<Derive*>(pb4);   // up cast 
    if(pd4){                    // pd4 is NULL here
        pd4 -> f();
        pd4 -> f2(); 
    }

    delete pb3;
    delete pb4;

    return 0;
}

专门用于用于有继承关系类之间转换,尤其是向下转换,运行时确定,是类型安全的。

3、RTTI

CentOS7 防火墙配置

最近在CentOS 7上使用rsync,莫名奇妙发现无法连通,最后发现是防火墙打开了。。。
而且CentOS7上防火墙进行了升级,不是使用的iptables
开启防火墙 (关闭stop)

systemctl start firewalld.service

添加端口

firewall-cmd –permanent –zone=public –add-port=873/tcp

(添加后,记得firewall-cmd reload)

一段简单c++代码

最近遇到一个有意思的问题,代码很简单,涉及到的东西会有一些,包括编译、虚函数等。

class A {
    public:
        A() {}
        A(int i) {}
        virtual void print(int j = 1){
            cout << "base: " << j << endl;
        }
};
class B : public A {
    public:
        virtual void print(int j = 4){
            cout << "derived: " << j << endl;
        }
    B() {}
    A a;
};
int main(){
    B b;
    A *a = &b;
    a -> print();
    return 0;
}

上面输出是多少

md5算法原理

MD5还是好久之前学过,早已经忘了。。。最近项目需要,又拾起来看了看,记录下来。
MD5,中文名为消息摘要算法,常用于数据完整校验。通过特定hash散列算法将文本信息转换为简短的消息摘要。
md5以512位分组处理信息,每一分组又被划分为16个32位子分组,变换处理后,输出由4个32位分组组成,级联生成一个128位散列值。
大致过程如下:

  • 填充

假设输入信息长度为len,则以len%512=448公式对数据进行填充

  • 填充消息长度

将消息长度上一步结果的后64位,如果消息长度大于2^64,则取低64位

  • 变换处理

常数 A=0X67452301L,B=0XEFCDAB89L,C=0X98BADCFEL,D=0X10325476L
4个变换函数

F(X,Y,Z)=(X&Y)|((~X)&Z)
G(X,Y,Z)=(X&Z)|(Y&(~Z))
H(X,Y,Z)=X^Y^Z
I(X,Y,Z)=Y^(X|(~Z))

4轮变换的操作

FF(a,b,c,d,Mj,s,ti)表示a=b+((a+F(b,c,d)+Mj+ti)<<<s)
GG(a,b,c,d,Mj,s,ti)表示a=b+((a+G(b,c,d)+Mj+ti)<<<s)
HH(a,b,c,d,Mj,s,ti)表示a=b+((a+H(b,c,d)+Mj+ti)<<<s)
II(a,b,c,d,Mj,s,ti)表示a=b+((a+I(b,c,d)+Mj+ti)<<<s)
Mj表示消息的第j个子分组

每轮变换后A、B、C、D分别加上a、b、c、d,然后进行一下轮
 

无锁编程

无锁编程有两个关键点:

  • 原子操作或者cas原语
  • aba问题
inline bool CAS(pointer_t *addr, pointer_t &old_value, pointer_t &new_value)
{
    bool ret;
    __asm__ __volatile__(
    "lock cmpxchg16b %1;\n"
    "sete %0;\n"
    :"=m"(ret),"+m" (*(volatile pointer_t *) (addr))
    :"a" (old_value.ptr), "d" (old_value.tag), "b" (new_value.ptr), "c" (new_value.tag));
    return ret;
}

lock锁地址总线
aba问题:在读取v值和进行cas操作前,如果有一个线程将v值由a改成b,另外一个线程在将b改成了a,就会cas操作混乱。解决方法是用标记或版本编号与进行CAS操作的每个值相关联,并原子的更新值和标记。
 
入队列

EnQueue (x) //进队列改良版
{
    q = new record ();
    q->value = x;
    q->next = NULL;
    p = tail;
    oldp = p
    do {
        while (p->next != NULL)
            p = p->next;
    } while( CAS (p.next, NULL, q) != TRUE); //如果没有把结点链上,再试
    CAS (tail, oldp, q); //置尾结点
}

出队列

DeQueue () //出队列
{
    do{
        p = head;
        if (p->next == NULL){
            return ERR_EMPTY_QUEUE;
        }
    while( CAS (head, p, p->next);
    return p->next->value;
}

 
https://www.codeproject.com/Articles/153898/Yet-another-implementation-of-a-lock-free-circular
http://www.drdobbs.com/parallel/writing-lock-free-code-a-corrected-queue/210604448?pgno=1
https://blog.csdn.net/chenjiayi_yun/article/details/16333595
https://blog.csdn.net/mergerly/article/details/39009473 scsp无锁kfifo队列
 
 
 
 

struct宏定义一些用法

typedef struct {
    int name;
    int service;
    int server;
    char buf[0];
} item;
#define OFFSET(A, B) (int)&(((A*)0)->B)
#define CONTAINER(PTR, TYPE, FIELD) (TYPE*)((char*)PTR-OFFSE    T(TYPE, FIELD))

第一个,获取变量偏移地址;
第二个,获取变量对应结构体的地址。
 

CUDA Dynamic Parallelism

 
CUDA 5.0中引入动态并行化,使得在device端执行的kernel的线程也能跟在host上一样launch kernels,只有支持CC3.5或者以上的设备中才能支持。动态并行化使用CUDA Device Runtime library(cudadevrt),它是一个能在device code中调用的CUDA runtime子集。

编译链接

为了支持动态并行化,必须使用两步分离编译和链接的过程:首先,设定-c和-rdc=true(–relocatable-device-code=true)来生成relocatable device code来进行后续链接,可以使用-dc(–device -c)来合并这两个选项;然后将上一步目标文件和cudadevrt库进行连接生成可执行文件,-lcudadevrt。过程如下图Figure 4: The Separate Compilation and Linking Process for Dynamic Parallelism.

nvcc -arch=sm_35 -dc myprog.cu -o myprog.o
nvcc -arch=sm_35 myprog.o -lcudadevrt -o myprog

或者简化成一步

nvcc -arch=sm_35 -rdc=true myprog.cu -lcudadevrt -o myprog.o

执行、同步

在CUDA编程模型中,一组执行的kernel的线程块叫做一个grid。在CUDA动态并行化,parent grid能够调用child grids。child grid继承parant grid的特定属性和限制,如L1 cache、shared_memory、栈大小。如果一个parent grid有M个block和N个thread,如果对child kernel launch没有控制的话,那个将产生M*N个child kernel launch。如果想一个block产生一个child kernel,那么只需要其中一个线程launch a kernel就行。如下

if(threadIdx.x == 0) {
  child_k <<< (n + bs - 1) / bs, bs >>> ();
}

grid lanuch是完全嵌套的,child grids总是在发起它们的parent grids结束前完成,这可以看作是一个一种隐式的同步。

如果parent kernel需要使用child kernel的计算结果,也可以使用CudaDeviceSynchronize(void)进行显示的同步,这个函数会等待一个线程块发起的所有子kernel结束。往往不知道一个线程块中哪些子kernel已经执行,可以通过下述方式进行一个线程块级别的同步

void threadBlockDeviceSynchronize(void) {
  __syncthreads();
  if(threadIdx.x == 0)
    cudaDeviceSynchronize();
  __syncthreads();
}

CudaDeviceSynchronize(void)调用开销较大,不是必须的时候,尽量减少使用,同时不要在父kernel退出时调用,因为结束时存在上述介绍的隐式同步。

内存一致

当子 grids开始与结束之间,父grids和子grids有完全一致的global memory view。

当子kernel launch的时候,global memory视图不一致。

__device__ int v = 0;
__global__ void child_k(void) {
  printf("v = %d\n", v);
}
__global__ void parent_k(void) {
  v = 1;
  child_k <<< 1, 1 >>>> ();
  v = 2; // RACE CONDITION
  cudaDeviceSynchronize();
}

在子kernel launch之后,显示同步之前,parent grid不能对 child grid读取的内存做写入操作,否则会造成race condition。

向Child grids传递指针

指针的传递存在限制:

  • 可以传递的指针:global memory(包括__device__变量和malloc分配的内存),zero-copy host端内存,常量内存。
  • 不可以传递的指针:shared_memory(__shared__变量), local memory(包括stack变量)

Device Streams和Events

所有在device上创建的streams都是non-blocking的,不支持默认NULL stream的隐式同步。创建流的方式如下

cudaStream_t s;
cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);

一旦一个device stream被创建,它能被一个线程块中其他线程使用。只有当这个线程块完成执行的时候,这个stream才能被其他线程块或者host使用。反之亦然。
Event也是支持的,不过有限制,只支持在不同stream之间使用cudaStreamWaitEvent()指定执行顺序,而不能使用event来计时或者同步。

Recursion Depth和Device Limits

递归深度包括两个概念:

  • nesting depth:递归grids的最大嵌套层次,host端的为0;
  • synchronization depth:cudaDeviceSynchronize()能调用的最大嵌套层次,host端为1,cudaLimitDevRuntimeSyncDepth应该设定为maximum 所以你吃肉你咋体on depth加1,设定方式如 cudaDeviceLimit(cudaLimitDevRuntimeSyncDepth, 4).

maximum nesting depth有硬件限制,在CC3.5中, 对depth 的限制为24. synchronization depth也一样。
从外到内,直到最大同步深度,每一次层会保留一部分内存来保存父block的上下文数据,即使这些内存没有被使用。所以递归深度的设定需要考虑到每一层所预留的内存。
另外还有一个限制是待处理的子grid数量。pending launch buffer用来维持launch queue和追踪当前执行kernel的状态。通过

cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 32768);

来设定合适的限制。否则通过cudaGetLastError()调用可以返回CudaErrorLaunchPendingCountExceeded的错误。
动态并行化执行有点类似树的结构,但与CPU上树处理也有些不同。类似深度小,分支多,比较茂密的树的执行结构,比较适合动态并行化的处理。深度大,每层节点少的树的执行结构,则不适合动态并行化。

characteristic tree processing dynamic parallelism
node thin (1 thread) thick (many threads)
branch degree small (usually < 10) large (usually > 100)
depth large small

参考:http://devblogs.nvidia.com/parallelforall/cuda-dynamic-parallelism-api-principles/