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

ID 683521
日期 9/26/2022
Public
文档目录

10.4. 重复使用数据的片上存储

对于英特尔 Stratix 10设计,被缓存的LSU会阻止某些优化。 Intel® 建议您避免推断缓存LSU。

以下未优化的代码创建了一个已缓存的突发合并(burst-coalesced)LSU,它会消耗更多资源并禁用其它优化:

kernel void cached (global int * restrict in,
                    global int * restrict out) {
   int i = get_global_id(0);
   int idx = out[i];
   int cached_value = in[idx]; // Burst-coalesced cached LSU
   out[i] = cached_value;
}

为了放置缓存,请将该存储器指针标记为volatile,如以下优化的代码中所示:

kernel void not_cached (global volatile int * restrict in,
                        global int * restrict out) {
   int i = get_global_id(0);
   int idx = out[i];
   int not_cached_value = in[idx];
   out[i] = not_cached_value;
}

有关优化load-store units(加载-存储单元,LSU)的更多信息,请参阅Load-Store Units(加载存储单元)小节。

Intel® 还建议您使用片上存储来实现最佳效果。请注意,与非英特尔 Stratix 10器件相比,英特尔 Stratix 10具有更大的M20K与ALM比率,使得您可以创建更大的局部存储器系统。

以下未优化的代码有一个函数,它接收从全局存储器中的数组来的指针。该情况下,离线编译器会修改数组,随后将其回存到存储器中。然后,在外部循环的后续迭代中,重新使用该数组。

void cached_array(int N, int M, int BUF_SIZE, 
                  global float2 * global_buf)
{
   for (uint i = 0; i< N; i++) {
      float2 data_array[BUF_SIZE];
      for (uint j= 0; j < M; j++) {
	          
         data_array[i] = global_buf [j]; //Load value
		 
         ... //do work to modify data_array 
		 
         global_buf[j] = data_array[i]; 
      } 
   } 
}

为了防止不必要的全局存储器访问,请在内核中定义一个专用数组(private array)来声明片上的数组。函数访问专用数组而非访问片上声明的数组。由此,该数组接收片上局部存储器存储。访问该片上局部存储器不需要对全局存储器的访问。

优化的代码:

void local_array(int N, int M, int BUF_SIZE, 
                 global float2 * global_buf)
{
   float2 local_buf[BUF_SIZE];
   populate_buf(local_buf, global_buf); 

   for (uint i = 0; i< N; i++) {
      float2	data_array[BUF_SIZE];			
      for (uint j= 0; j < M; j++) {

         data_array[i] = local_buf[j];//Load value 

		 ... //do work to modify data_array

         local_buf[j] = data_array[i]; 

      }
   } 
}