在内核 OpenCL 中实现 FIFO 的最佳方法
Best approach to FIFO implementation in a kernel OpenCL
目标：在 OpenCL 中实现下图。 OpenCl 内核需要做的主要事情是将系数数组和临时数组相乘，然后在最后将所有这些值累加为一个。 （这可能是最耗时的操作，并行性在这里真的很有帮助）。
我正在为内核使用一个辅助函数来执行乘法和加法（我希望这个函数也是并行的）。
图片说明：
一次一个，这些值被传递到与系数数组大小相同的数组（临时数组）中。现在 每次 将单个值传递到此数组时，临时数组与系数数组并行相乘，然后将每个索引的值连接成一个元素。这将一直持续到输入数组到达它的最后一个元素。
我的代码会怎样？
对于输入的 60 个元素，它需要超过 8000 毫秒！！我总共有 120 万个输入仍然需要传递。我知道有一个更好的解决方案来完成我正在尝试的事情。下面是我的代码。
这里有一些我知道他的代码肯定有问题的地方。当我尝试将系数值与临时数组相乘时，它崩溃了。这是因为global_id。我想让这条线做的只是将两个数组并行相乘。
我试图弄清楚为什么执行 FIFO 功能需要这么长时间，所以我开始注释掉这些行。我首先注释了除 FIFO 函数的第一个 for 循环之外的所有内容。结果这花了 50 毫秒。然后，当我取消注释下一个循环时，它跳到了 8000 毫秒。所以延迟与数据传输有关。
是否有我可以在 OpenCl 中使用的寄存器移位？也许对整数数组使用一些逻辑移位方法？ （我知道有一个“>>”运算符）。
float constant temp[58];
float constant tempArrayForShift[58];
float constant multipliedResult[58];
float fifo(float inputValue, float *coefficients, int sizeOfCoeff) {
//take array of 58 elements (or same size as number of coefficients)
//shift all elements to the right one
//bring next element into index 0 from input
//multiply the coefficient array with the array thats the same size of coefficients and accumilate
//store into one output value of the output array
//repeat till input array has reached the end
int globalId = get_global_id(0);
float output = 0.0f;
//Shift everything down from 1 to 57
//takes about 50ms here
for(int i=1; i<58; i++){
tempArrayForShift[i] = temp[i];
}
//Input the new value passed from main kernel. Rest of values were shifted over so element is written at index 0.
tempArrayForShift[0] = inputValue;
//Takes about 8000ms with this loop included
//Write values back into temp array
for(int i=0; i<58; i++){
temp[i] = tempArrayForShift[i];
}
//all 58 elements of the coefficient array and temp array are multiplied at the same time and stored in a new array
//I am 100% sure this line is crashing the program.
//multipliedResult[globalId] = coefficients[globalId] * temp[globalId];
//Sum the temp array with each other. Temp array consists of coefficients*fifo buffer
for (int i = 0; i < 58; i ++) {
// output = multipliedResult[i] + output;
}
//Returned summed value of temp array
return output;
}
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) {
//Initialize the temporary array values to 0
for (int i = 0; i < 58; i ++) {
temp[i] = 0;
tempArrayForShift[i] = 0;
multipliedResult[i] = 0;
}
//fifo adds one element in and calls the fifo function. ALL I NEED TO DO IS SEND ONE VALUE AT A TIME HERE.
for (int i = 0; i < 60; i ++) {
Output[i] = fifo(Array[i], coefficients, 58);
}
}
我在使用 OpenCl 时遇到这个问题很长时间了。我不确定如何同时执行并行指令和顺序指令。
我正在考虑的另一种选择
在主 cpp 文件中，我正在考虑在那里实现 fifo 缓冲区并让内核执行乘法和加法。但这意味着我必须在一个循环中调用内核 1000 多次。这是更好的解决方案吗？还是完全没有效率。
要获得 GPU 的良好性能，您需要将您的工作并行化到多个线程。在您的代码中，您只使用单个线程，GPU 每个线程的速度非常慢，但如果多个线程同时 运行，则速度可能非常快。在这种情况下，您可以为每个输出值使用一个线程。您实际上不需要通过数组移动值：对于每个输出值，考虑 58 个值的 window，您可以从内存中获取这些值，将它们与系数相乘并写回结果。
一个简单的实现是（启动与输出值一样多的线程）：
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output)
{
int globalId = get_global_id(0);
float sum=0.0f;
for (int i=0; i< 58; i++)
{
float tmp=0;
if (globalId+i > 56)
{
tmp=Array[i+globalId-57]*coefficient[57-i];
}
sum += tmp;
}
output[globalId]=sum;
}
这并不完美，因为它生成的内存访问模式不是 GPU 的最佳选择。缓存可能会有所帮助，但显然还有很大的优化空间，因为值会被多次重用。您尝试执行的操作称为卷积 (1D)。 NVidia 在他们的 GPU Computing SDK 中有一个名为 oclConvolutionSeparable 的二维示例，它显示了一个优化版本。您适应使用他们的 convolutionRows 内核进行一维卷积。
这是您可以试用的另一个内核。有很多同步点（障碍），但这应该表现得相当好。 65 项工作组不是很理想。
步骤：
- 将本地值初始化为 0
- 将系数复制到局部变量
遍历要计算的输出元素：
- 移动现有元素（仅限工作项 > 0）
- 复制新元素（仅限工作项 0）
- 计算点积
5a.乘法 - 每个工作项一个
5b.计算总和的缩减循环
- 将点积复制到输出（仅限 WI 0）
- 最后一道屏障
代码：
__kernel void lowpass(__global float *Array, __constant float *coefficients, __global float *Output, __local float *localArray, __local float *localSums){
int globalId = get_global_id(0);
int localId = get_local_id(0);
int localSize = get_local_size(0);
//1 init local values to 0
localArray[localId] = 0.0f
//2 copy coefficients to local
//don't bother with this id __constant is working for you
//requires another local to be passed in: localCoeff
//localCoeff[localId] = coefficients[localId];
//barrier for both steps 1 and 2
barrier(CLK_LOCAL_MEM_FENCE);
float tmp;
for(int i = 0; i< outputSize; i++)
{
//3 shift elements (+barrier)
if(localId > 0){
tmp = localArray[localId -1]
}
barrier(CLK_LOCAL_MEM_FENCE);
localArray[localId] = tmp
//4 copy new element (work item 0 only, + barrier)
if(localId == 0){
localArray[0] = Array[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//5 compute dot product
//5a multiply + barrier
localSums[localId] = localArray[localId] * coefficients[localId];
barrier(CLK_LOCAL_MEM_FENCE);
//5b reduction loop + barrier
for(int j = 1; j < localSize; j <<= 1) {
int mask = (j << 1) - 1;
if ((localId & mask) == 0) {
localSums[local_index] += localSums[localId +j]
}
barrier(CLK_LOCAL_MEM_FENCE);
}
//6 copy dot product (WI 0 only)
if(localId == 0){
Output[i] = localSums[0];
}
//7 barrier
//only needed if there is more code after the loop.
//the barrier in #3 covers this in the case where the loop continues
//barrier(CLK_LOCAL_MEM_FENCE);
}
}
更多的工作组呢？
这被稍微简化以允许单个 1x65 工作组计算机整个 1.2M 输出。要允许多个工作组，您可以使用 /get_num_groups(0) 来计算每个组应该做的工作量 (workAmount)，并调整 i for-loop:
for (i = workAmount * get_group_id(0); i< (workAmount * (get_group_id(0)+1) -1); i++)
步骤 #1 也必须更改为 localArray 初始化为正确的起始状态，而不是全为 0。
//1 init local values
if(groupId == 0){
localArray[localId] = 0.0f
}else{
localArray[localSize - localId] = Array[workAmount - localId];
}
这两项更改应该允许您使用更优化的工作组数量；我建议设备上计算单元数量的一些倍数。不过，请尝试将每个小组的工作量保持在数千个。尝试解决这个问题，有时在 运行.
的情况下，在高级别上看似最佳的做法会对内核造成损害
优势
几乎在这个内核的每个点上，工作项都有事情要做。唯一少于 100% 的项目在工作的时间是在步骤 5b 的减少循环期间。 Read more here about why that is a good thing.
缺点
由于障碍的本质，障碍会减慢内核的速度：暂停一个工作项，直到其他工作项到达该点。也许有一种方法可以以更少的障碍实现它，但我仍然认为这是最佳的，因为你正在尝试解决这个问题。
每个组没有容纳更多工作项的空间，65 个也不是最佳大小。理想情况下，您应该尝试使用 2 的幂或 64 的倍数。但这不会是一个大问题，因为内核中有很多障碍使得它们都相当规律地等待。
目标：在 OpenCL 中实现下图。 OpenCl 内核需要做的主要事情是将系数数组和临时数组相乘，然后在最后将所有这些值累加为一个。 （这可能是最耗时的操作，并行性在这里真的很有帮助）。
我正在为内核使用一个辅助函数来执行乘法和加法（我希望这个函数也是并行的）。
图片说明：
一次一个，这些值被传递到与系数数组大小相同的数组（临时数组）中。现在 每次 将单个值传递到此数组时，临时数组与系数数组并行相乘，然后将每个索引的值连接成一个元素。这将一直持续到输入数组到达它的最后一个元素。
我的代码会怎样？
对于输入的 60 个元素，它需要超过 8000 毫秒！！我总共有 120 万个输入仍然需要传递。我知道有一个更好的解决方案来完成我正在尝试的事情。下面是我的代码。
这里有一些我知道他的代码肯定有问题的地方。当我尝试将系数值与临时数组相乘时，它崩溃了。这是因为global_id。我想让这条线做的只是将两个数组并行相乘。
我试图弄清楚为什么执行 FIFO 功能需要这么长时间，所以我开始注释掉这些行。我首先注释了除 FIFO 函数的第一个 for 循环之外的所有内容。结果这花了 50 毫秒。然后，当我取消注释下一个循环时，它跳到了 8000 毫秒。所以延迟与数据传输有关。
是否有我可以在 OpenCl 中使用的寄存器移位？也许对整数数组使用一些逻辑移位方法？ （我知道有一个“>>”运算符）。
float constant temp[58];
float constant tempArrayForShift[58];
float constant multipliedResult[58];
float fifo(float inputValue, float *coefficients, int sizeOfCoeff) {
//take array of 58 elements (or same size as number of coefficients)
//shift all elements to the right one
//bring next element into index 0 from input
//multiply the coefficient array with the array thats the same size of coefficients and accumilate
//store into one output value of the output array
//repeat till input array has reached the end
int globalId = get_global_id(0);
float output = 0.0f;
//Shift everything down from 1 to 57
//takes about 50ms here
for(int i=1; i<58; i++){
tempArrayForShift[i] = temp[i];
}
//Input the new value passed from main kernel. Rest of values were shifted over so element is written at index 0.
tempArrayForShift[0] = inputValue;
//Takes about 8000ms with this loop included
//Write values back into temp array
for(int i=0; i<58; i++){
temp[i] = tempArrayForShift[i];
}
//all 58 elements of the coefficient array and temp array are multiplied at the same time and stored in a new array
//I am 100% sure this line is crashing the program.
//multipliedResult[globalId] = coefficients[globalId] * temp[globalId];
//Sum the temp array with each other. Temp array consists of coefficients*fifo buffer
for (int i = 0; i < 58; i ++) {
// output = multipliedResult[i] + output;
}
//Returned summed value of temp array
return output;
}
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output) {
//Initialize the temporary array values to 0
for (int i = 0; i < 58; i ++) {
temp[i] = 0;
tempArrayForShift[i] = 0;
multipliedResult[i] = 0;
}
//fifo adds one element in and calls the fifo function. ALL I NEED TO DO IS SEND ONE VALUE AT A TIME HERE.
for (int i = 0; i < 60; i ++) {
Output[i] = fifo(Array[i], coefficients, 58);
}
}
我在使用 OpenCl 时遇到这个问题很长时间了。我不确定如何同时执行并行指令和顺序指令。
我正在考虑的另一种选择
在主 cpp 文件中，我正在考虑在那里实现 fifo 缓冲区并让内核执行乘法和加法。但这意味着我必须在一个循环中调用内核 1000 多次。这是更好的解决方案吗？还是完全没有效率。
要获得 GPU 的良好性能，您需要将您的工作并行化到多个线程。在您的代码中，您只使用单个线程，GPU 每个线程的速度非常慢，但如果多个线程同时 运行，则速度可能非常快。在这种情况下，您可以为每个输出值使用一个线程。您实际上不需要通过数组移动值：对于每个输出值，考虑 58 个值的 window，您可以从内存中获取这些值，将它们与系数相乘并写回结果。
一个简单的实现是（启动与输出值一样多的线程）：
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output)
{
int globalId = get_global_id(0);
float sum=0.0f;
for (int i=0; i< 58; i++)
{
float tmp=0;
if (globalId+i > 56)
{
tmp=Array[i+globalId-57]*coefficient[57-i];
}
sum += tmp;
}
output[globalId]=sum;
}
这并不完美，因为它生成的内存访问模式不是 GPU 的最佳选择。缓存可能会有所帮助，但显然还有很大的优化空间，因为值会被多次重用。您尝试执行的操作称为卷积 (1D)。 NVidia 在他们的 GPU Computing SDK 中有一个名为 oclConvolutionSeparable 的二维示例，它显示了一个优化版本。您适应使用他们的 convolutionRows 内核进行一维卷积。
这是您可以试用的另一个内核。有很多同步点（障碍），但这应该表现得相当好。 65 项工作组不是很理想。
步骤：
- 将本地值初始化为 0
- 将系数复制到局部变量
遍历要计算的输出元素：
- 移动现有元素（仅限工作项 > 0）
- 复制新元素（仅限工作项 0）
- 计算点积
5a.乘法 - 每个工作项一个
5b.计算总和的缩减循环 - 将点积复制到输出（仅限 WI 0）
- 最后一道屏障
代码：
__kernel void lowpass(__global float *Array, __constant float *coefficients, __global float *Output, __local float *localArray, __local float *localSums){
int globalId = get_global_id(0);
int localId = get_local_id(0);
int localSize = get_local_size(0);
//1 init local values to 0
localArray[localId] = 0.0f
//2 copy coefficients to local
//don't bother with this id __constant is working for you
//requires another local to be passed in: localCoeff
//localCoeff[localId] = coefficients[localId];
//barrier for both steps 1 and 2
barrier(CLK_LOCAL_MEM_FENCE);
float tmp;
for(int i = 0; i< outputSize; i++)
{
//3 shift elements (+barrier)
if(localId > 0){
tmp = localArray[localId -1]
}
barrier(CLK_LOCAL_MEM_FENCE);
localArray[localId] = tmp
//4 copy new element (work item 0 only, + barrier)
if(localId == 0){
localArray[0] = Array[i];
}
barrier(CLK_LOCAL_MEM_FENCE);
//5 compute dot product
//5a multiply + barrier
localSums[localId] = localArray[localId] * coefficients[localId];
barrier(CLK_LOCAL_MEM_FENCE);
//5b reduction loop + barrier
for(int j = 1; j < localSize; j <<= 1) {
int mask = (j << 1) - 1;
if ((localId & mask) == 0) {
localSums[local_index] += localSums[localId +j]
}
barrier(CLK_LOCAL_MEM_FENCE);
}
//6 copy dot product (WI 0 only)
if(localId == 0){
Output[i] = localSums[0];
}
//7 barrier
//only needed if there is more code after the loop.
//the barrier in #3 covers this in the case where the loop continues
//barrier(CLK_LOCAL_MEM_FENCE);
}
}
更多的工作组呢？
这被稍微简化以允许单个 1x65 工作组计算机整个 1.2M 输出。要允许多个工作组，您可以使用 /get_num_groups(0) 来计算每个组应该做的工作量 (workAmount)，并调整 i for-loop:
for (i = workAmount * get_group_id(0); i< (workAmount * (get_group_id(0)+1) -1); i++)
步骤 #1 也必须更改为 localArray 初始化为正确的起始状态，而不是全为 0。
//1 init local values
if(groupId == 0){
localArray[localId] = 0.0f
}else{
localArray[localSize - localId] = Array[workAmount - localId];
}
这两项更改应该允许您使用更优化的工作组数量；我建议设备上计算单元数量的一些倍数。不过，请尝试将每个小组的工作量保持在数千个。尝试解决这个问题，有时在 运行.
的情况下，在高级别上看似最佳的做法会对内核造成损害优势
几乎在这个内核的每个点上，工作项都有事情要做。唯一少于 100% 的项目在工作的时间是在步骤 5b 的减少循环期间。 Read more here about why that is a good thing.
缺点
由于障碍的本质，障碍会减慢内核的速度：暂停一个工作项，直到其他工作项到达该点。也许有一种方法可以以更少的障碍实现它，但我仍然认为这是最佳的，因为你正在尝试解决这个问题。
每个组没有容纳更多工作项的空间，65 个也不是最佳大小。理想情况下，您应该尝试使用 2 的幂或 64 的倍数。但这不会是一个大问题，因为内核中有很多障碍使得它们都相当规律地等待。