OpenCl 重点总结

half 数据类型

• 符号 IEEE 754-2008 标准
• vstore_half 和 vload_half 存在 float->half / half->float 的转换操作
• half 类型仅用于声明一个指针（指向 half value 的内存）

合法使用案例

void
bar (__global half *p)
{
    ...
}

__kernel void
foo (__global half *pg, __local half *pl)
{
    __global half *ptr;
    int offset;

    ptr = pg + offset;
    bar(ptr);
}

非法使用案例

half a;
half b[100];
half *p;
a = *p; //  not allowed. must use *vload_half* function

数据类型转换及算术运算规则

• 如果一个矢量和标量进行算术运算，opencl 会将标量扩充为和矢量元素相当的矢量进行算术运算

• 后缀和前缀递增和递减运算符（-- 和 ++）适用于内置的标量和矢量类型，但不适用于内置的标量和矢量浮点类型

• 算术运算符加（+）、减（-）、乘（*）和除（/）适用于内置的整数和浮点标量以及矢量数据类型。取余（%）操作适用于内置的整数标量和整数矢量数据类型。

• 按位与（&）、按位或（|）、按位异或（^）和按位取反（~）运算符适用于除了内置标量和矢量浮点类型之外的所有标量和矢量内置类型。

• 右移位运算符（>>）、左移位运算符（<<）适用于除了内置标量和矢量浮点类型之外的所有标量和矢量内置类型。对于右移位（>>）和左移位（<<）运算符，如果第一个操作数是标量，则最右边的操作数必须是标量；如果第一个操作数是矢量，则最右边的操作数可以是矢量或标量。

地址空间

• 通用地址空间：OpenCL 2.0引入了未命名的通用地址空间。如果指针声明时没有指定具名地址空间，那么它将默认指向通用地址空间。

• 函数指针参数和返回值：如果函数中的指针参数或返回值没有声明具名地址空间，那么它们被默认定义为指向通用地址空间。

• 内核函数参数：内核函数中声明为指针或数组类型的参数必须指向具名地址空间，可以是__global、__local或__constant中的一个。

• 具名地址空间：__global、__local和__constant是具名地址空间的示例。这些地址空间是通用地址空间的子集，但常量地址空间除外。

• 指针赋值：指针之间的赋值有一些限制。指针A只能被赋值给与其相同地址空间A的指针，或者是通用地址空间的指针。如果A和B是具名地址空间，并且A不同于B，则将地址空间A的指针转换为地址空间B的指针是非法的。

• 图像对象：图像对象总是从全局地址空间分配的，因此不应为图像类型指定__global或global限定符。图像对象的元素不能直接访问，而是提供了用于读取和写入图像对象的内置函数。

• 初始化：在全局地址空间中声明的变量可以进行初始化，但只能使用常量表达式进行初始化。这意味着初始化表达式必须在编译时就能够确定其值，而不能依赖于运行时计算。

• 在程序范围（全局）和函数内部的静态变量也可以在全局地址空间中声明。这些变量可以是任何OpenCL C支持的数据类型，但不能是“其他内置数据类型”中的类型（例如 event_t，不能声明成 global）。这意味着它们可以是用户自定义的类型，或者是指向用户自定义类型的指针。

global int foo;         // OK.
int foo;                // OK. Declared in the global address space
global uchar buf[512];  // OK.
global int baz = 12;    // OK. Initialization is allowed
static global int bat;  // OK. Internal linkage

static int foo;         // OK. Declared in the global address space
static global int foo;  // OK.

int *foo;               // OK. foo is allocated in global address space.
                        // pointer to foo in generic address space

void func(...)
{
    int *foo;           // OK. foo is allocated in private address space.
                        // foo points to a location in generic address space.
    ...
}

global int * global ptr;          // OK.
int * global ptr;                 // OK.
constant int *global ptr=&baz;    // error since baz is in global address
                                  // space.
global int * constant ptr = &baz; // OK

// Pointers work. Also, initialization to a constant known at
// program load time
global int *global baz_ptr = &baz;

global image2d_t im; // Error. Invalid type for program scope
                     // variables

global event_t ev;   // Error. Invalid type for program scope variables

global int *bad_ptr; // Error. No implicit address space

• 局部内存只能声明在 kernel function 内部，并且不能被初始化

kernel void my_func(...)
{
    local float a;     // A single float allocated
                       // in local address space

    local float b[10]; // An array of 10 floats
                       // allocated in local address space.

    if (...)
    {
        // example of variable in __local address space but not
        // declared at __kernel function scope.
        local float c; // not allowed.
    }
}

kernel void my_func(...)
{
    local float a = 1; // not allowed

    local float b;
    b = 1;             // allowed
}

• __constant 地址空间位于 global memory 中，属于 read-only 类型，字符串类型的数据应该存储在常量地址空间中，必须使用编译时的常数进行初始化
• 未使用地址空间限定符声明的 kernel function 内的变量、非 kernel function 内的所有变量以及所有函数参数都位于 __private 或私有地址空间中。
• 指向 global, local 或者 private 地址空间的指针可以隐式转换到通用地址空间，反之不行；
• 指针（没有指向空间）可以从 global, local 或者 private 转换到通用地址空间，反之也行；
• constant 不能进行这种转换

int *ptr;
global int g;
ptr = &g; // legal

local int l;
ptr = &l; // legal

private int p;
ptr = &p; // legal

constant int c;
ptr = &c; // illegal

global int *gp;
local int *lp;
private int *pp;

int *p;
p = gp; // legal
p = lp; // legal
p = pp; // legal

// it is illegal to convert from a generic pointer
// to an explicit address space pointer without a cast:
gp = p; // compile-time error
lp = p; // compile-time error
pp = p; // compile-time error

• 当两个操作数是指针时，相等运算符的行为受到一些约束
  • 如果其中一个操作数是指针，而另一个是空指针常量，则空指针常量将被转换为指针的类型
  • 如果两个操作数都是指针，并且它们指向不同的地址空间，那么这两个地址空间必须有重叠部分，否则行为是未定义的
  • 如果两个操作数是指向不同地址空间的指针，那么首先将其中一个指针转换为指向另一个指针的地址空间，使得它们指向相同的地址空间

bool callee(int *p1, int *p2)
{
    if (p1 == p2)
        return true;
    return false;
}

void caller()
{
    global int *gptr = 0xdeadbeef;
    private int *pptr = 0xdeadbeef;

    // behavior of callee is undefined
    bool b = callee(gptr, pptr);
}

• 条目运算符
  • 如果第二个和第三个操作数是指向不同地址空间的指针，则这两个地址空间必须重叠。

kernel void test1()
{
    global int arr[5] = { 0, 1, 2, 3, 4 };
    int *p = &arr[1];
    global int *q = &arr[3];
    local int *r = NULL;
    int *val = NULL;

    // legal. 2nd and 3rd operands are in address spaces
    // that overlap
    val = (q >= p) ? q : p;

    // compiler error. 2nd and 3rd operands are in disjoint
    // address spaces
    val = (q >= p) ? q : r;
}

• 结构体或联合体的成员不能具有地址空间限定符。
• 一个类型不能同时具有不同地址空间的限定符。

访问限定符

image object 可以使用 read_only、write_only、read_write 修饰，如果在 kernel function 中没有访问限定符，默认是 read_only

• 无采样器的 image object 可以使用 read_write 修饰
• 有采样器的 image object 不能使用 read_write 修饰，否则会报编译时错误

kernel attribute

__kernel 限定符可以与关键字 attribute 一起使用，用于声明关于内核函数的附加信息，如 attribute((vec_type_hint())) 是对编译器的提示，旨在表示 __kernel 的计算宽度，并且应作为编译器在自动向量化代码时计算处理器带宽利用率的依据。

// autovectorize assuming float4 as the
// basic computation width
__kernel __attribute__((vec_type_hint(float4)))
void foo( __global float4 *p ) { ... }

// autovectorize assuming double as the
// basic computation width
__kernel __attribute__((vec_type_hint(double)))
void foo( __global float4 *p ) { ... }

// autovectorize assuming int (default)
// as the basic computation width
__kernel
void foo( __global float4 *p ) { ... }

• attribute((work_group_size_hint(X, Y, Z))) 是对编译器的提示，旨在指定可能使用的工作组大小， 即最可能由 local_work_size 参数传递给 clEnqueueNDRangeKernel

• 可选的 attribute((reqd_work_group_size(X, Y, Z))) 是必须作为 local_work_size 参数传递给 clEnqueueNDRangeKernel 的工作组大小。这允许编译器为此内核适当地优化生成的代码。

• 可选的 attribute((nosvm)) 限定符可与指针变量一起使用，通知编译器该指针不引用共享虚拟内存区域。

存储说明符

支持 typedef、extern 和 static 存储类说明符。不支持 auto 和 register 存储类说明符。

• extern 存储类说明符只能用于程序作用域中声明的函数（内核函数和非内核函数）和全局变量，或者用于函数内部（内核函数和非内核函数）声明的变量。
• static 存储类说明符只能用于非内核函数、程序作用域中声明的全局变量以及在全局或常量地址空间中声明的函数内的变量。

extern constant float4 noise_table[256];
static constant float4 color_table[256];

extern kernel void my_foo(image2d_t img);
extern void my_bar(global float *a);

kernel void my_func(image2d_t img, global float *a)
{
    extern constant float4 a;
    static constant float4 b = (float4)(1.0f); // OK.
    static float c;  // Error: No implicit address space
    global int hurl; // Error: Must be static
    ...
    my_foo(img);
    ...
    my_bar(a);
    ...
    while (1)
    {
        static global int inside; // OK.
        ...
    }
    ...
}

规则限制

• 指针
  • 在程序中声明的内核函数的参数如果是指针，必须用 __global、__constant 或 __local 限定符声明。
  • 用 __constant 限定符声明的指针只能分配给相应用 __constant 限定符声明的指针。
  • 不允许使用指向函数的指针。
• image object
  • 图像类型（image2d_t、image3d_t、image2d_array_t、image1d_t、image1d_buffer_t 或 image1d_array_t）只能用作函数参数的类型。图像函数参数不可修改。只能使用内置的图像读取和写入函数访问图像的元素。
  • 图像类型不能用于声明变量、结构体或联合体字段、图像数组、指向图像的指针，也不能用作函数的返回类型
  • 图像类型不能与 __global、__private、__local 和 __constant 地址空间限定符一起使用。
• simpler_t
  • 采样器类型（sampler_t）只能用作函数参数的类型或在程序范围或内核函数的最外层范围中声明的变量类型。在内核函数的非最外层范围中声明的采样器变量的行为由实现定义。采样器参数或变量不可修改。
  • 采样器类型不能用于声明结构体或联合体字段、采样器数组、指向采样器的指针，也不能用作函数的返回类型。
  • 采样器类型不能与 __local 和 __global 地址空间限定符一起使用。
• 其他
  • 位字段结构成员目前不受支持。
  • 不支持变长数组和具有灵活（或无大小）数组的结构体。
  • 除了 printf 和 enqueue_kernel 外，不支持可变参数宏和函数。
  • 如果函数声明中的参数列表为空，则该函数不接受任何参数。这是由于上述可变原型的限制。
  • 除非在 OpenCL 规范中定义，否则 C99 标准头文件 assert.h、ctype.h、complex.h、errno.h、fenv.h、float.h、inttypes.h、limits.h、locale.h、setjmp.h、signal.h、stdarg.h、stdio.h、stdlib.h、string.h、tgmath.h、time.h、wchar.h 和 wctype.h 中定义的库函数、宏、类型和常量不可用，也不能被程序包含。
  • 不支持 auto 和 register 存储类说明符。
  • 不支持递归。
  • 内核函数的返回类型必须是 void。
  • 程序中的内核函数的参数不能声明为内置标量类型 bool、size_t、ptrdiff_t、intptr_t 和 uintptr_t，也不能声明为包含被声明为这些内置标量类型之一的字段的结构体和/或联合体。（这些类型的字节大小是实现定义的，而且在 OpenCL 设备和主机处理器上也可能不同，这使得难以分配缓冲区对象作为传递给指针声明为这些类型之一的内核的参数。）
  • 不支持 half 标量，因为 half 只能用作存储格式（可以传递 half 指针），不能用作可执行浮点运算的数据类型。
  • 支持 C99 规范定义的类型限定符 const、restrict 和 volatile。这些限定符不能与 image2d_t、image3d_t、image2d_array_t、image2d_depth_t、image2d_array_depth_t、image1d_t、image1d_buffer_t 和 image1d_array_t 类型一起使用。除指针类型外的其他类型不得使用 restrict 限定符。
  • 事件类型 (event_t) 不能用作内核函数参数的类型。事件类型不能用于声明程序范围变量。事件类型不能用于声明结构体或联合体字段。事件类型不能与 __local、__constant 和 __global 地址空间限定符一起使用。
  • 不能将 clk_event_t、ndrange_t 和 reserve_id_t 类型用作从主机入队的内核函数的参数。clk_event_t 和 reserve_id_t 类型不能在程序范围内声明。
  • 对 queue_t、clk_event_t、ndrange_t 和 reserve_id_t 类型应用 sizeof 运算符返回的值是实现定义的。
  • OpenCL 程序中的函数不能被称为 main。
  • 不支持隐式函数声明。
  • 由主机入队的内核必须继续具有指向命名地址空间的类型声明的指针参数。

隐式函数声明指的是在调用函数之前没有提前声明函数的情况下直接调用函数。在 C 语言中，如果在调用函数之前没有进行函数声明或者函数原型的提前声明，编译器会假定函数返回类型为 int，并且可以接受任意数量和类型的参数。

这种情况下编译器不会检查函数的返回类型和参数类型是否正确，容易导致编译错误或者运行时错误。因此，隐式函数声明是一种不推荐的编程做法，应该在调用函数之前提供函数的声明或者原型。

预处理指令 及 宏

#pragma 指令的描述如下，记得增加 OPENCL 标记，否则会报错

// on-off-switch is one of ON, OFF, or DEFAULT
#pragma OPENCL FP_CONTRACT on-off-switch

#pragma OPENCL EXTENSION extensionname : behavior

#pragma OPENCL EXTENSION all : behavior

其他的宏：
FILE
当前源文件的假定名称（一个字符字符串字面值）。

LINE
当前源行（在当前源文件中）的假定行号（一个整数常量）。

OPENCL_VERSION
替换为反映由OpenCL设备支持的OpenCL版本号的整数。本文档中描述的OpenCL版本将使__OPENCL_VERSION__替换为整数200。

CL_VERSION_1_0
替换为整数100，反映OpenCL 1.0版本。

CL_VERSION_1_1
替换为整数110，反映OpenCL 1.1版本。

CL_VERSION_1_2
替换为整数120，反映OpenCL 1.2版本。

CL_VERSION_2_0
替换为整数200，反映OpenCL 2.0版本。

OPENCL_C_VERSION
如果未指定-cl-std构建选项，则在为每个设备编译程序时，将使用每个设备支持的最高的OpenCL C 1.x语言版本作为OpenCL C的版本。如果指定了-cl-std=CL2.0，则本文档中描述的OpenCL C版本将使__OPENCL_C_VERSION__替换为整数200。

ENDIAN_LITTLE
用于确定OpenCL设备是小端架构还是大端架构（如果设备是小端，则为1的整数常量；否则未定义）。还参考CL_DEVICE_ENDIAN_LITTLE设备查询的值。

kernel_exec(X, type__n) (and kernel_exec(X, typen))

在这里插入代码片__kernel __attribute__((work_group_size_hint(X, 1, 1))) \
    __attribute__((vec_type_hint(type__n__)))

IMAGE_SUPPORT
用于确定OpenCL设备是否支持图像。如果支持图像，则这是一个整数常量为1，否则未定义。还参考CL_DEVICE_IMAGE_SUPPORT设备查询的值。

FAST_RELAXED_MATH
用于确定是否在提供给clBuildProgram或clCompileProgram的构建选项中指定了-cl-fast-relaxed-math优化选项。如果指定了-cl-fast-relaxed-math构建选项，则这是一个整数常量为1，否则未定义。

NULL宏扩展为一个空指针常量。 值为0的整数常量表达式，或将此类表达式转换为void *类型，称为空指针常量。

预定义标识符__func__可用。

attribute

• types
  • 对于枚举、结构或联合类型，您可以在枚举、结构或联合标签和类型名称之间或在定义的右大括号后指定属性。前一种语法更受推荐。

align:
struct S { short f[3]; } __attribute__ ((aligned (8)));
struct __attribute__ ((aligned (8))) S { short f[3]; };
typedef int more_aligned_int __attribute__ ((aligned (8)));

pack:
struct my_unpacked_struct
{
    char c;
    int i;
};

struct __attribute__ ((packed)) my_packed_struct
{
    char c;
    int i;
    struct my_unpacked_struct s;
};

• variables

align:
int x __attribute__ ((aligned (16))) = 0;
// double-word aligned
struct foo { int x[2] __attribute__ ((aligned (8))); };

pack:
struct foo
{
    char a;
    int x[2] __attribute__ ((packed));
};

放置在用户定义类型开始处的属性列表适用于该类型的变量，而不是类型本身，而跟在类型主体后面的属性适用于类型本身。

/* a has alignment of 128 */
__attribute__((aligned(128))) struct A {int i;} a;

/* b has alignment of 16 */
__attribute__((aligned(16))) struct B {double d;}
__attribute__((aligned(32))) b ;

struct A a1; /* a1 has alignment of 4 */

struct B b1; /* b1 has alignment of 32 */

endian 属性确定变量的字节顺序。endiantype 可设置为 host，表示变量使用主机处理器的字节顺序，或设置为 device，表示变量使用将要执行内核的设备的字节顺序。默认值为 device。

• endian 属性只能应用于位于 global 或 constant 地址空间中的指针类型。
• endian 属性不能用于非指针类型的变量。
• 当一个指针被赋给另一个指针时，两个指针的 endian 属性值必须相同。

global float4 *p __attribute__ ((endian(host)));

• Unrolling Loops

attribute((opencl_unroll_hint)) 和 attribute((opencl_unroll_hint(n))) 属性限定符可用于指定循环（for、while 和 do 循环）可以展开。此属性限定符可用于指定完全展开或按指定数量部分展开。这是一个编译器提示，编译器可能会忽略此指令。

n 是循环展开因子，必须是正整数的编译时常量表达式。展开因子为 1 禁用展开。如果未指定 n，则编译器确定循环的展开因子。

__attribute__((opencl_unroll_hint(2)))
while (*s != 0)
    *p++ = *s++;

• Funciton
  之前已经介绍过了，可以制定 type 和 work_size

Blocks

你可以使用 ^ 运算符来声明一个 Block 变量，并表示一个 Block 字面量的开始。Block 本身的主体包含在 {} 中，如下例所示（和 C 语言一样，; 表示语句的结束）(理解成匿名函数和函数指针吧)：

Block 变量持有对 Block 的引用。你可以使用与声明指向函数的指针类似的语法来声明它们，只是使用 ^ 替代 *。Block 类型与 C 类型系统完全兼容。以下是有效的 Block 变量声明：

void (^blockReturningVoidWithVoidArgument)(void);
int (^blockReturningIntWithIntAndCharArguments)(int, char);

• Block参数与void：如果一个Block不接受任何参数，那么在声明时必须在参数列表中指定void，表示参数为空。
• 不能进行指针解引用：不能使用指针解引用操作符 * 来解引用Block引用，因此无法在编译时计算Block的大小。
• 编译器可以检查 Block 的参数类型是否与其声明一致，以及返回值的类型是否与 Block 的返回类型匹配。这样可以避免在运行时出现类型不匹配的错误，提高代码的安全性和可维护性。
  -可以为Blocks创建自定义的类型，尤其是当同一签名的Block在多个地方重复使用时，这是一种良好的实践方式。如下面的代码：

typedef float (^MyBlockType)(float, float);

MyBlockType myFirstBlock = // ...;
MyBlockType mySecondBlock = // ...;

一个 Block 字面表达式产生一个指向 Block 的引用。它通过使用 ^ 符号作为一元运算符来引入。

^ block_decl compound_statement_body

其中，block_decl 可以是空的，也可以是参数列表或类型表达式。对于类型表达式，可以扩展允许使用 ^ 来表示一个指向 Block 的引用，类似于函数引用中使用 * 的方式。

^ void (void) { printf("hello world**\n**"); }

上面的表达式产生一个不带参数且没有返回值的 Block 的引用。等价于下面两行代码

^ ( void ) { printf("hello world**\n**"); }
^ { printf("hello world**\n**"); }

• 复合语句体在父级范围内建立了新的词法作用域，其中的变量按照正常的方式与 Block 绑定。但对于自动（栈）存储中的变量，会被捕获到 Block 中作为 const 副本。
• 编译器会在评估 Block 字面表达式时执行变量捕获，但如果编译器可以证明对变量的引用实际上不会被评估，它就不需要捕获变量。
• 在 Block 中声明的变量的生命周期与函数的生命周期相同。
• Block 字面表达式可以嵌套，所有被任何嵌套 Block 捕获的变量也隐式地被捕获到它们的封闭 Block 的作用域中。
• Block 字面表达式可以用作全局或本地静态范围中 Block 变量的初始化值，并且也可以在程序范围内声明为全局字面量。

int GlobalInt = 0;

int (^getGlobalInt)(void) = ^{ return GlobalInt; };

• 不支持具有可变参数的 Blocks
• 不支持 Blocks 的数组
• 不支持将 Blocks 用作结构体和联合体的成员
• 不支持 __block 存储类型： 在一些编程语言中（比如 Objective-C、Swift 等），__block 存储类型用于在 Block 内部修改外部变量的值。但在 OpenCL C 中，不支持这种存储类型，意味着不能在 Block 内部修改外部变量的值。
• 不支持 Block_copy() 和 Block_release() 函数： 在一些其他编程语言中，特别是在 Objective-C 中，Block 可以被复制和释放。Block_copy() 函数用于复制 Block 对象，Block_release() 函数用于释放 Block 对象。然而，在 OpenCL C 中，由于不支持这些函数，因此无法进行 Block 的复制和释放操作。
• 一个 Block 不能是函数的返回值。
• Blocks 不能作为三元选择运算符（?:）的表达式。
• 不能使用一元操作符（* 和 &）与 Block 一起。
• 不允许指向 Blocks 的指针。
• 一个 Block 不能捕获在外部作用域中声明的另一个 Block 变量。
• 在一个作用域中分配的 Block 变量只能与相同作用域或任何嵌套作用域一起使用。
• 在 OpenCL C 中，程序作用域的 Block 变量不能使用 extern 关键字进行声明。
• 在 OpenCL C 中，所有的 Block 变量的声明都隐式地带有 const 修饰符，这意味着它们必须在声明时进行初始化，并且不能在后续的代码中重新赋值。

block 作为参数

void foo(int *x, int (^bar)(int, int))
{
    *x = bar(*x, *x);
}

kernel
void k(global int *x, global int *z)
{
    if (some expression)
        foo(x, ^int(int x, int y){return x+y+*z;}); // legal
    else
        foo(x, ^int(int x, int y){return (x*y)-*z;}); // legal
}

block 在声明时就必须赋值

kernel
void k(global int *x, global int *z)
{
    int ^(tmp)(int, int);
    if (some expression)
    {
        tmp = ^int(int x, int y){return x+y+*z;}); // illegal
    }
    *x = foo(x, tmp);
}

block 不能被重新赋值，或者声明为 extern

int GlobalInt = 0;
int (^getGlobalInt)(void) = ^{ return GlobalInt; }; // legal
int (^getAnotherGlobalInt)(void);                   // illegal
extern int (^getExternGlobalInt)(void);             // illegal

void foo()

{
    ...
    getGlobalInt = ^{ return 0; }; // illegal - cannot assign to
                                   // a global block variable
    ...
}

一个 Block 不能捕获在外部作用域中声明的另一个 Block 变量，可以捕获全局的 block 变量

void (^bl0)(void) = ^{
    ...
};

kernel void k()
{
    void(^bl1)(void) = ^{
        ...
    };

    void(^bl2)(void) = ^{
        bl0(); // legal because bl0 is a global
               // variable available in this scope
        bl1(); // illegal because bl1 would have to be captured
    };
}

opencl 数值一致性

包括 inf、nan、误差等，参考官方文档：Opencl Numerical Compliance

• -cl-unsafe-math-optimizations：这是一个编译器选项，当在编译OpenCL程序时指定这个选项，编译器会进行一些优化，这些优化可能会牺牲一些数学运算的精度以换取性能的提升。
• 在某些实现中，powr() 或 pown() 可能比 pow() 执行得更快。如果已知 x 大于等于0，考虑使用 powr() 替代 pow()；如果已知 y 是一个整数，考虑使用 pown() 替代 pow()。
• 对于那些C99规范的F.9节或其他浮点数精度的类似节规定了结果的值，舍入误差或溢出行为的通常允许不适用。这些值必须产生完全规定的结果，没有其他可能。例如，ceil(-1 < x < 0) 返回 -0。如果使用了±符号，符号应该被保留。例如，sin(±0) = ±0 应该被解释为 sin(+0) 是 +0，sin(-0) 是 -0。
• half_函数的行为：带有half_前缀的函数（例如half_sqrt表示半精度的平方根函数）应该与没有前缀的同名函数行为相同。它们必须符合相同的边界情况要求，如C99规范的F.9节和G.6节所述。对于其他情况，除非另有说明，这些单精度函数允许有高达8192 ULPs的误差（以单精度结果为测量标准），尽管鼓励更好的精度。
• NaN的返回：如果一个函数有多个NaN操作数，它应该返回其中一个NaN操作数。如果函数返回一个NaN操作数，它可能会忽略这个NaN（如果它是一个信号NaN）。非信号NaN应该被转换为非信号NaN，信号NaN应该被转换为普通的NaN，并且最好是转换为非信号NaN。NaN的有效载荷位的其他部分或NaN的符号如何转换是未定义的。

如果非规格化数被强制舍入为零，则函数可能返回以下四种结果之一：

• 对于非强制舍入为零模式的任何符合要求的结果。
• 如果1中的结果在舍入前是一个次规格化数，则它可能被强制舍入为零。
• 如果函数的一个或多个次规格化操作数被强制舍入为零，则函数可以返回任何非强制舍入的符合要求的结果。
• 如果3中的结果在舍入前是一个次规格化数，则结果可能被强制舍入为零。

在上述每种情况下，如果一个操作数或结果被强制舍入为零，则零的符号是未定义的。

如果非规格化数被强制舍入为零，设备可以选择符合以下情况而不是列在额外要求部分中的情况：

nextafter(+最小正规数, y < +最小正规数) = +0。
nextafter(-最小正规数, y > -最小正规数) = -0。
nextafter(-0, y > 0) 返回最小的正规数值。
nextafter(+0, y < 0) 返回最小的负规数值。

为了明确起见，非规格化数或次规格化数被定义为范围在 0 < x < TYPE_MIN 和 -TYPE_MIN < x < -0 的可表示数集。它们不包括 ±0。如果一个非零数在舍入前被归一化，其基数-2的指数小于(TYPE_MIN_EXP - 1)。

在这里，TYPE_MIN 和 TYPE_MIN_EXP 应该由适用于考虑的浮点类型的常量替换，例如 float 的 FLT_MIN 和 FLT_MIN_EXP。

image object 回顾

采样器 simpler_t

在OpenCL中，采样器（Sampler）是用于在图像上执行采样操作的对象。它定义了图像采样的方式和参数，以及采样过程中使用的过滤器类型。采样器通常与read_image函数一起使用，用于从图像中获取特定坐标处的像素值。

具体来说，采样器的作用包括：

• 定义采样方式： 采样器确定了在图像上执行采样操作的方式，例如采样点的位置和采样范围。
• 指定过滤器类型： 采样器可以指定在进行采样时所使用的过滤器类型，例如最近邻插值、线性插值或者其他更复杂的插值算法。
• 控制边界处理： 采样器可以定义在图像边界处的采样行为，例如是否对越界的像素进行镜像、重复或者截断处理。
• 处理归一化坐标： 采样器可以将输入的坐标映射到图像的实际像素坐标上，并确保采样操作在图像范围内进行。

采样器的地址映射和滤波模式

如果寻址模式不是 CLK_ADDRESS_REPEAT no 或 CLK_ADDRESS_MIRRORED_REPEAT.

• Filter Mode CLK_FILTER_NEAREST
  当滤波模式为 CLK_FILTER_NEAREST 时，获取图像中距离 (u,v,w) 最近（曼哈顿距离）的图像元素。这意味着位于位置 (i,j,k) 的图像元素的值成为图像元素值，其中

i = address_mode((int)floor(u))
j = address_mode((int)floor(v))
k = address_mode((int)floor(w))

Addressing Mode	Result of address_mode(coord)
CLK_ADDRESS_CLAMP_TO_EDGE	clamp (coord, 0, size - 1)
CLK_ADDRESS_CLAMP	clamp (coord, -1, size)
CLK_ADDRESS_NONE	coord

clamp 的定义如下：

clamp(a, b, c) = return (a < b) ? b : ((a > c) ? c : a)

• Filter Mode CLK_FILTER_LINEAR

当滤波模式设置为 CLK_FILTER_LINEAR 时，对于2D图像，会选择一个2×2的图像元素正方形；对于3D图像，会选择一个2×2×2的图像元素立方体。这个2×2的正方形或2×2×2的立方体的获取方式如下。

i0 = address_mode((int)floor(u - 0.5))
j0 = address_mode((int)floor(v - 0.5))
k0 = address_mode((int)floor(w - 0.5))
i1 = address_mode((int)floor(u - 0.5) + 1)
j1 = address_mode((int)floor(v - 0.5) + 1)
k1 = address_mode((int)floor(w - 0.5) + 1)
a = frac(u - 0.5)
b = frac(v - 0.5)
c = frac(w - 0.5)

在这里，frac(x) 表示 x 的小数部分，并且计算方式为 x - floor(x)。

对于3D图像，图像元素的值是这样找到的：

T = (1 - a) * (1 - b) * (1 - c) * T_i0j0k0
    + a * (1 - b) * (1 - c) * T_i1j0k0
    + (1 - a) * b * (1 - c) * T_i0j1k0
    + a * b * (1 - c) * T_i1j1k0
    + (1 - a) * (1 - b) * c * T_i0j0k1
    + a * (1 - b) * c * T_i1j0k1
    + (1 - a) * b * c * T_i0j1k1
    + a * b * c * T_i1j1k1

在这里，T_ijk 表示3D图像中位置为 (i,j,k) 的图像元素。

对于 2D 图像：

T = (1 - a) * (1 - b) * T_i0j0
    + a * (1 - b) * T_i1j0
    + (1 - a) * b * T_i0j1
    + a * b * T_i1j1

在这里，T_ij 表示2D图像中位置为 (i,j) 的图像元素。

如果在上述方程中的任何选定的 T_ijk 或 T_ij 指向图像外的位置，那么边界颜色将被用作 T_ijk 或 T_ij 的颜色值。

如果图像通道类型是 CL_FLOAT 或 CL_HALF_FLOAT，并且任何图像元素 T_ijk 或 T_ij 是 INF（无穷大）或 NaN（不是一个数字），那么内置图像读取函数的行为是未定义的。

寻址模式是 CLK_ADDRESS_REPEAT.

如果 (s,t,r) 中的值是 INF 或 NaN，那么内置图像读取函数的行为是未定义的。

• Filter Mode CLK_FILTER_NEAREST
  图像位于（i,j,k）的元素计算如下：

u = (s - floor(s)) * w_t
i = (int)floor(u)
if (i > w_t - 1)
    i = i - w_t

v = (t - floor(t)) * h_t
j = (int)floor(v)
if (j > h_t - 1)
    j = j - h_t

w = (r - floor(r)) * d_t
k = (int)floor(w)
if (k > d_t - 1)
    k = k - d_t

其中 (s,t,r)表示图像当前的索引，（$w_t$, $h_t$, $d_t$），表示图像宽，高，深度

• Filter Mode CLK_FILTER_LINEAR
  当滤波模式设置为 CLK_FILTER_LINEAR 时，对于2D图像，会选择一个2×2的图像元素正方形；对于3D图像，会选择一个2×2×2的图像元素立方体。这个2×2的正方形或2×2×2的立方体的获取方式如下。

u = (s - floor(s)) * w_t
i0 = (int)floor(u - 0.5)
i1 = i0 + 1
if (i0 < 0)
    i0 = w_t + i0
if (i1 > w_t - 1)
    i1 = i1 - w_t

v = (t - floor(t)) * h_t
j0 = (int)floor(v - 0.5)
j1 = j0 + 1
if (j0 < 0)
    j0 = h_t + j0
if (j1 > h_t - 1)
    j1 = j1 - h_t

w = (r - floor(r)) * d_t
k0 = (int)floor(w - 0.5)
k1 = k0 + 1
if (k0 < 0)
    k0 = d_t + k0
if (k1 > d_t - 1)
    k1 = k1 - d_t

a = frac(u - 0.5)
b = frac(v - 0.5)
c = frac(w - 0.5)

在这里，frac(x) 表示 x 的小数部分，并且计算方式为 x - floor(x)。

对于3D图像，图像元素的值是这样找到的：

T = (1 - a) * (1 - b) * (1 - c) * T_i0j0k0
    + a * (1 - b) * (1 - c) * T_i1j0k0
    + (1 - a) * b * (1 - c) * T_i0j1k0
    + a * b * (1 - c) * T_i1j1k0
    + (1 - a) * (1 - b) * c * T_i0j0k1
    + a * (1 - b) * c * T_i1j0k1
    + (1 - a) * b * c * T_i0j1k1
    + a * b * c * T_i1j1k1

在这里，T_ijk 表示3D图像中位置为 (i,j,k) 的图像元素。

对于 2D 图像：

T = (1 - a) * (1 - b) * T_i0j0
    + a * (1 - b) * T_i1j0
    + (1 - a) * b * T_i0j1
    + a * b * T_i1j1

在这里，T_ij 表示2D图像中位置为 (i,j) 的图像元素。

如果图像通道类型是 CL_FLOAT 或 CL_HALF_FLOAT，并且任何图像元素 T_ijk 或 T_ij 是 INF（无穷大）或 NaN（不是一个数字），那么内置图像读取函数的行为是未定义的。

寻址模式为 CLK_ADDRESS_MIRRORED_REPEAT

当寻址模式设置为 CLK_ADDRESS_MIRRORED_REPEAT 时，图像的读取方式会模拟一种镜像重复的效果。这意味着图像会在每个整数坐标的边界上进行镜像，形成一个连续的、无缝的图案。在这种模式下，当你尝试访问图像的某个坐标时，如果这个坐标超出了图像的实际边界，它会在图像的另一侧找到对应的镜像坐标。

如果 (s,t,r) 中的值是 INF 或 NaN，那么内置图像读取函数的行为是未定义的。

• Filter Mode CLK_FILTER_NEAREST
  图像中（i,j,k）的计算如下：

s' = 2.0f * rint(0.5f * s)
s' = fabs(s - s')
u = s' * w_t
i = (int)floor(u)
i = min(i, w_t - 1)

t' = 2.0f * rint(0.5f * t)
t' = fabs(t - t')
v = t' * h_t
j = (int)floor(v)
j = min(j, h_t - 1)

r' = 2.0f * rint(0.5f * r)
r' = fabs(r - r')
w = r' * d_t
k = (int)floor(w)
k = min(k, d_t - 1)

• Filter Mode CLK_FILTER_LINEAR
  当滤波模式设置为 CLK_FILTER_LINEAR 时，对于2D图像，会选择一个2×2的图像元素正方形；对于3D图像，会选择一个2×2×2的图像元素立方体。这个2×2的正方形或2×2×2的立方体的获取方式如下。

s' = 2.0f * rint(0.5f * s)
s' = fabs(s - s')
u = s' * w_t
i0 = (int)floor(u - 0.5f)
i1 = i0 + 1
i0 = max(i0, 0)
i1 = min(i1, w_t - 1)

t' = 2.0f * rint(0.5f * t)
t' = fabs(t - t')
v = t' * h_t
j0 = (int)floor(v - 0.5f)
j1 = j0 + 1
j0 = max(j0, 0)
j1 = min(j1, h_t - 1)

r' = 2.0f * rint(0.5f * r)
r' = fabs(r - r')
w = r' * d_t
k0 = (int)floor(w - 0.5f)
k1 = k0 + 1
k0 = max(k0, 0)
k1 = min(k1, d_t - 1)

a = frac(u - 0.5)
b = frac(v - 0.5)
c = frac(w - 0.5)

在这里，frac(x) 表示 x 的小数部分，并且计算方式为 x - floor(x)。

对于3D图像，图像元素的值是这样找到的：

T = (1 - a) * (1 - b) * (1 - c) * T_i0j0k0
    + a * (1 - b) * (1 - c) * T_i1j0k0
    + (1 - a) * b * (1 - c) * T_i0j1k0
    + a * b * (1 - c) * T_i1j1k0
    + (1 - a) * (1 - b) * c * T_i0j0k1
    + a * (1 - b) * c * T_i1j0k1
    + (1 - a) * b * c * T_i0j1k1
    + a * b * c * T_i1j1k1

在这里，T_ijk 表示3D图像中位置为 (i,j,k) 的图像元素。

对于 2D 图像：

T = (1 - a) * (1 - b) * T_i0j0
    + a * (1 - b) * T_i1j0
    + (1 - a) * b * T_i0j1
    + a * b * T_i1j1

在这里，T_ij 表示2D图像中位置为 (i,j) 的图像元素。

对于 1D 图像

T = (1 - a) * T_i0
    + a * T_i1

在这里，T_i 表示2D图像中位置为 (i) 的图像元素。

如果图像通道类型是 CL_FLOAT 或 CL_HALF_FLOAT，并且任何图像元素 T_ijk 或 T_ij 是 INF（无穷大）或 NaN（不是一个数字），那么内置图像读取函数的行为是未定义的。

read_imagef 的行为

• 读取的图像类型 CL_UNORM_INT8 and CL_UNORM_INT16 -> 规范化浮点数，值域 【0.f，1.f】
• 读取的图像类型 CL_SNORM_INT8 and CL_SNORM_INT16 -> 规范化浮点数，值域 【-1.f，1.f】

转换过程：

CL_UNORM_INT8 (8-bit unsigned integer) → float
normalized float value = (float)c / 255.0f

CL_UNORM_INT_101010 (10-bit unsigned integer) → float
normalized float value = (float)c / 1023.0f

CL_UNORM_INT16 (16-bit unsigned integer) → float
normalized float value = (float)c / 65535.0f

CL_SNORM_INT8 (8-bit signed integer) → float
normalized float value = max(-1.0f, (float)c / 127.0f)

CL_SNORM_INT16 (16-bit signed integer) → float
normalized float value = max(-1.0f, (float)c / 32767.0f)

For CL_UNORM_INT8

• 0 must convert to 0.0f and
• 255 must convert to 1.0f

For CL_UNORM_INT_101010

• 0 must convert to 0.0f and
• 1023 must convert to 1.0f

For CL_UNORM_INT16

• 0 must convert to 0.0f and
• 65535 must convert to 1.0f

For CL_SNORM_INT8

• -128 and -127 must convert to -1.0f,
• 0 must convert to 0.0f and
• 127 must convert to 1.0f

For CL_SNORM_INT16

• -32768 and -32767 must convert to -1.0f,
• 0 must convert to 0.0f and
• 32767 must convert to 1.0f

剩下的看官方文档吧，懒的写了：opencl 文档

sRGBA to RGB 的转换

标准的RGB数据大致以线性的亮度级别显示颜色，opencl 处理 sRGB 图像会将其转换至 RGB，并在写出时，将 RGB 转换至 sRGB :

• sRGB颜色空间：sRGB是一种广泛使用的颜色空间，它模拟了人类视觉系统对亮度的感知。在sRGB空间中，颜色值通常是以8位无符号整型（即0到255的整数）来表示的。这些整数值在显示设备上映射到感知上均匀分布的亮度级别。例如，0对应于最暗（0.0f），255对应于最亮（1.0f），而中间的值则对应于介于两者之间的亮度。

• 从sRGB到线性RGB的转换：当使用OpenCL的read_imagef内置函数读取sRGB图像时，会自动将sRGB颜色值转换为线性RGB颜色值。这种转换是为了在后续的图像处理中使用线性颜色空间，这在数学运算和颜色混合中更为方便。

• alpha通道的处理：如果图像格式包含alpha通道（用于表示透明度），那么alpha数据会以线性颜色空间存储。这意味着alpha值也会从sRGB空间转换到线性空间。

• 写入sRGB图像：当使用write_imagef内置函数写入图像时，如果图像通道顺序是上述描述的sRGB值之一，并且设备支持写入sRGB图像，那么会自动将线性RGB颜色值转换回sRGB空间。

转换规则：以一个无符号 8 位整数转换至浮点数为例（read_imagef）：

if (c <= 0.04045),
    result = c / 12.92;
else
    result = powr((c + 0.055) / 1.055, 2.4);

从 RGB 转换到 sRGB 最终值必须加 0.5 (write_imagef):

if (c is NaN)
    c = 0.0;
if (c > 1.0)
    c = 1.0;
else if (c < 0.0)
    c = 0.0;
else if (c < 0.0031308)
    c = 12.92 * c;
else
    c = 1.055 * powr(c, 1.0/2.4) - 0.055;

scaled_reference_result = c * 255
channel_component = floor(scaled_reference_result + 0.5);

如何根据图像坐标值来选择2D图像数组中的特定层（layer）进行读取或写入操作。

• 2D图像数组中的坐标选择：

  • 当使用采样器（sampler）读取2D图像时，选择的图像层（layer）是通过以下方式计算的：
    • layer = clamp(rint(w), 0, dt - 1)
    • 这里，w 是未归一化的图像坐标值，rint 是四舍五入到最接近的整数的函数，clamp 函数用于确保结果在有效范围内（从0到图像深度 dt 减去1）。
  • 如果不使用采样器，选择的层（layer）直接由 w 的值决定：
    • layer = w
    • 由于 w 已经是整数，所以结果未定义（即行为不确定）如果 w 不是0, 1, …, dt - 1 中的一个整数。

• 1D图像数组中的坐标选择：

  • 当使用采样器读取1D图像时，选择的1D图像层（layer）是通过以下方式计算的：
    • layer = clamp(rint(v), 0, ht - 1)
    • 类似地，v 是未归一化的图像坐标值，ht 是图像的高度。
  • 如果不使用采样器，选择的层（layer）直接由 v 的值决定：
    • layer = v
    • 同样，如果 v 不是0, 1, …, ht - 1 中的一个整数，结果也是未定义的。

rint 是一个数学函数，用于将浮点数四舍五入到最接近的整数，关于 opencl built-in function 参考：opencl built-in-function