3.6.3. 控制Load-Store Units

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

下载 PDF

ID 683521

日期 9/26/2022

版本 22.3

Public

3.6.3. 控制Load-Store Units

Intel® FPGA SDK for OpenCL™ Offline Compiler允许您通过一组可以从全局存储器加载和对其储存的内置调用指令来控制针对全局存储器而生成的LSU的类型。

Load Built-ins

下表总结了load built-in的类型：

表 10. Load Built-ins
Built-in	实现的LSU类型
`__pipelined_load()`	流水线式（如果可能）
`__prefetching_load()`	预取式（如果可能）
`__burst_coalesced_load()`	突发合并
`__burst_coalesced_cached_load()`	突发合并缓存（如果可能）

所有变体都需要以下自变量：

表 11. Load Built-in自变量
Built-in	类型	描述
Argument #1	Pointer	要从中加载的存储器位置。
Argument #2	Integer	仅适用于 `__burst_coalesced_cached_load()`函数。以字节描述LSU cache的大小。。非负编译时间常量整数。
Return value（返回值）	Object	指针自变量所指向的数据。与指针自变量的基础类型相同的类型。

Store Built-ins

下表总结了store built-in的类型：

表 12. Store Built-ins
Built-in	实现的LSU类型
`__pipelined_store()`	流水线式（如果可能）
`__burst_coalesced_store()`	突发合并式

所有变体都需要以下自变量：

表 13. Store Built-in Arguments
Built-in	类型	描述
Argument #1	Pointer	要存储到的存储器位置。
Argument #2	与指针的基础类型相同	要存储的值。

注：所有的store built-in变体都是“无返回值”类型。

示例

以下是一个OpenCL实例，描述了load-和store built-ins的不同变体：


kernel void oclTest(global int * restrict in, 
                    global int * restrict out) {
    int i = get_global_id(0);
  
    int a1 = __pipelined_load(in + 3*i+0); // Uses a pipelined LSU
    // Uses a burst-coalesced LSU with a cache of size 1024 bytes
    int a2 = __burst_coalesced_cached_load(&in[3*i+1], 1024);  
    int a3 = __prefetching_load(&in[3*i+2]); // Uses a prefetching LSU

    __burst_coalesced_store(&out[3*i+0], a3); // Uses a burst-coalesced LSU
}

注：

编译器不允许您在需要LSU的情形下选择一个可能导致该情况下出现功能性错误结果的LSU。例如，如果您请求在易失性指针上的预取LSU，则编译器会出错。如果在cache（对于LSU来说是局部）因为其他LSU写入存储器而变得不连贯的情况下而请求进行缓存，那么编译器也会出错。
预取LSU不适用于英特尔 Stratix 10器件

选择您的语言

使用 Intel.com 搜索

快速链接

最近搜索

高级搜索

仅搜索

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

3.6.3. 控制Load-Store Units

Load Built-ins

Store Built-ins

示例