基础知识的介绍在这里就不赘述了,下面就来记录一些实操
向量数据的声明
【来自 ziyang 的瞎总结】
向量类型的格式【可能】:基本类型(如half、float)+v+此向量所容纳的基本类型数据个数(仅16或者32)
目前看到的向量类型有:halfv16
:包含16个半精度浮点数(half precision floating point)的向量floatv16
:包含16个单精度浮点数(floating point)的向量intv16
:包含16个整数(int)的向量uintv16
:16个无符号整数(unsigned int)的向量shortv32
:包含32个短整数(short int)的向量ushortv32
:包含32个无符号短整数(unsigned short int)的向量
向量数据的装载
simd_load
【来自ziyang的瞎总结】
向量类型~数组“类型”
向量类型声明的变量类似于数组,可以用数组的方式进行访问,并且向量变量名指向该向量开辟的空间的首地址【like数组“变量”】
函数原型:void simd_load(vector_type1 &place, const simple_type *place2)
例如:
int a = 0;
int b = 0;
intv16 vec_tem = 0;
simd_load(vec tem, &a);
simd_load(vec tem, &b);
simd_loadu
不对齐向量装入,将64字节(halfv16类型的装入是32字节)长度的数据从连续内存区域中装入到一个向量变量中。
函数原型:void simd_store(vector_type vec, const simple_type *place2)
例如:
__global__ void func(int *addr_0, half *addr_1)
{
// 64Byte整型数据装入 intv16
intv16 va_0;
simd_loadu(va_0, addr_0);
// 32Byte 16位浮点型数据装入 halfv16
halfv16 va_1;
simd_loadu(va_1, addr_1);
}
向量地址的装载
【来自ziyang的瞎总结】
这个对其应该是表示数据传输的时候数据地址是否连续simd_store
对齐向量存储,将一个向量变量中的数据存储到64字节(halfv16类型的存储是32字节)连续内存区域中。
函数原型:void simd_store(vector_type vec, const simple_type *place2)
例如:
// addr_0 64B 对齐,addr_1 32B 对齐
__global__ void func(int *addr_0, half *addr_1)
{
// 将 intv16 的数据写入到地址addr_0
intv16 va_0 = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
simd_store(va_0, addr_0);
// 将 halfv16 的数据写入到地址addr_1
halfv16 va_1 = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9, 1.0, 1.1, 1.2, 1.3, 1.4, 1.5, 1.6};
simd_store(va_1, addr_1);
}
simd_storeu
不对齐向量存储,将一个向量变量中的数据存储到64字节(halfv16类型的存储是32字节)连续内存区域中。
函数原型:void simd_store(vector_type vec, const simple_type *place2)
例如:
__global__ void func(int *addr_0, half *addr_1)
{
// 将 intv16 的数据写入到地址addr_0
intv16 va_0 = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
simd_storeu(va_0, addr_0);
// 将 halfv16 的数据写入到地址addr_1
halfv16 va_1 = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9, 1.0, 1.1, 1.2, 1.3, 1.4, 1.5, 1.6};
simd_storeu(va_1, addr_1);
}
# simd_seleq、simd_selle、simd_sellt
【来自 ziyang 的瞎总结】
符合前面提及的向量间返回的判断结果【布尔值】
可以考虑先判断然后使用这个函数输出