Xeon-Phi从主机openMP并行区异步卸载

Question

我在主机openMP代码中使用了英特尔的卸载编译指示。代码如下所示

int s1 = f(a,b,c);

#prama offload singnal(s1) in (...) out(x:len)
{
    for (int i = 0; i < len; ++i)
    {
        x[i] = ...
    }   
}

#pragma omp parallel default(shared)
{
    #pragma omp for schedule(dynamic) nowait
    for (int i = 0; i < count; ++i)
    {
        /* code */
    }

    #pragma omp for schedule(dynamic) 
    for (int j = 0; j < count2; ++j)
    {
        /* code */
    }
}

#pragma offload wait(s1)
{
    /* code */
}

$x$到MIC的代码卸载计算。代码通过将一些openMP分配给CPU核心来使自己保持忙碌。上面的代码按照预期工作。然而，第一个卸载杂注需要很长时间，并且已经成为瓶颈。然而，总的来说，将$x$的计算卸载到MIC是值得的。我正在尝试的一种可能克服此延迟问题的方法如下

int s1 = f(a,b,c);

#pragma omp parallel default(shared)
{
    #pragma omp single nowait
    {
        #prama offload singnal(s1) in (...) out(x:len)
        {
            for (int i = 0; i < len; ++i)
            {
                x[i] = ...
            }   
        }

    }

    #pragma omp for schedule(dynamic) nowait
    for (int i = 0; i < count; ++i)
    {
        /* code */
    }

    #pragma omp for schedule(dynamic) 
    for (int j = 0; j < count2; ++j)
    {
        /* code */
    }
}

#pragma offload wait(s1)
{
    /* code */
}

因此，这段新代码分配了一个线程来执行卸载，而其他openmp线程可以用于其他工作共享构造。但是，此代码不起作用。我收到以下错误消息

device 1 does not have a pending signal for wait(0x1)

卸载报告指出，上述代码是罪魁祸首。一种临时的解决办法是使用一个常量作为信号，即signal(0)，这是有效的。然而，我需要一个更持久的解决方案。有没有人知道我的代码出了什么问题？

谢谢

Answer 1

让我来补充一下Taylor的回答。

第一次卸载确实比后续卸载需要更多的时间，因为初始化工作正在进行。泰勒描绘了那里正在发生的一些事情。您可以使用环境变量OFFLOAD_INIT=on_start来避免虚拟卸载。这应该会让运行时系统提前完成所有的初始化。这种开销不会消失，但它会从您的第一次卸载转移到应用程序初始化。

您的第二个代码片段的问题似乎是您的卸载针对不同的设备。仅当信号和等待发生在同一目标设备上时，信令和等待才有效。由于您没有对卸载显式使用target(mic:0)子句，因此运行时系统很可能选择不同的目标设备。

我想提出的一个建议是，不要使用纯整数作为信令。通常，该信号指示某个缓冲区已准备就绪。在这些情况下，最好使用缓冲区指针作为信号句柄，因为它对于使用不同缓冲区的并发卸载是唯一的。

干杯，-michael

Answer 2

我不能评论第二个代码块。我对第一个问题有一些观察。

第一次卸载总是需要更长的时间，因为它还设置了卸载基础架构。这种结构包括传递环境变量、复制libomp5的mic实现、设置线程池等。

避免这种情况的方法是首先设置一个虚拟卸载，这意味着它实际上不会做任何事情，也不是计算块的一部分。

在software.intel.com/mic developer的training选项卡下有一组关于xeon phi协处理器优化的优秀参考资料。

还要看一下software.intel.com/en-us/articles/programming-and-compiling-for-intel-many-integrated-core-architecture，software.intel.com/en-us/articles/optimization-and-performance-tuning-for-intel-xeon-phi-coprocessors-part-1-optimization，和software.intel.com/en-us/articles/optimization-and-performance-tuning-for-intel-xeon-phi-coprocessors-part-1-optimization.

很抱歉我的URL很长，但是stackoverflow不允许我包含两个以上的链接，因为我是新的。