到目前为止，在内核执行时间方面，我的GPU慢于我的CPU。我以为也许是因为我只测试了一个小样本，所以由于启动开销较小，CPU最终完成得更快。但是，当我用几乎是样本大小10倍的数据测试内核时，CPU仍然完成得更快，而GPU却落后了将近400毫秒。
使用2.39MB文件运行时CPU：43.511ms GPU：65.219ms
使用32.9MB文件的运行时CPU：289.541ms GPU：605.400ms
我尝试使用本地内存，尽管我100％确信使用了错误的内存，但遇到了两个问题。内核完成1000-3000ms之间的任何时间（取决于我为localWorkSize设置的大小），或者遇到状态代码-5，即CL_OUT_OF_RESOURCES。
这是一个SO成员同伴帮助我的内核。
__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< 65; i++)
{
float tmp=0;
if (globalId+i > 63)
{
tmp=Array[i+globalId-64]*coefficients[64-i];
}
sum += tmp;
}
Output[globalId]=sum;
}
这是我尝试使用本地内存的尝试。第一部分将是主机代码的一部分，随后的部分是内核。
//Set the size of localMem
status |= clSetKernelArg(
kernel,
2,
1024, //I had num_items*(float) but it gave me a -5. Num items is the amount of elements in my array (around 1.2 million elements)
null);
printf("Kernel Arg output status: %i \n", status);
//set a localWorkSize
localWorkSize[0] = 64;
//execute the kernel with localWorkSize included
status = clEnqueueNDRangeKernel(
cmdQueue,
kernel,
1,
NULL,
globalWorkSize,
localWorkSize,
0,
NULL,
&someEvent);
//Here is what I did to the kernel***************************************
__kernel void lowpass(__global float *Array, __global float *coefficients, __global float *Output, __local float *localMem) {
int globalId = get_global_id(0);
int localId = get_local_id(0);
localMem[localId] = globalId[globalId];
float sum=0.0f;
for (int i=0; i< 65; i++)
{
float tmp=0;
if (globalId+i > 63)
{
tmp=localMem[i+localId-64]*coefficients[64-i];
}
sum += tmp;
}
Output[globalId]=sum;
}
尝试设置局部变量时使用的参考链接：如何在OpenCL中使用局部内存？
用于查找kernelWorkGroupSize的链接（这就是为什么在kernelArg中设置1024的原因）：CL_OUT_OF_RESOURCES包含200万个具有1GB VRAM的浮点数？
我已经看到其他人也遇到类似的问题，其中GPU的速度比CPU慢，但是对于许多人来说，他们使用的是clEnqueueKernel而不是clEnqueueNDRangeKernel。
如果您需要有关此内核的更多信息，这是我之前的问题：在内核OpenCL中实现FIFO的最佳方法
还找到了一些针对GPU的优化技巧。https://developer.amd.com/wordpress/media/2012/10/Optimizations-ImageConvolution1.pdf
编辑的代码；错误仍然存在
__kernel void lowpass2(__global float *Array, __global float *coefficients, __global float *Output) {
int globalId = get_global_id(0);
float sum=0.0f;
float tmp=0.0f;
for (int i=64-globalId; i< 65; i++)
{
tmp = 0.0f;
tmp=Array[i]*coefficients[i];
sum += tmp;
}
Output[globalId]=sum;
}
为2400万个元素数组运行以下内核
__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< 65; i++)
{
float tmp=0;
if (globalId+i > 63)
{
tmp=Array[i+globalId-64]*coefficients[64-i];
}
sum += tmp;
}
Output[globalId]=sum;
}
25个计算单元设备池在200毫秒内完成，而8核CPU则在500毫秒内完成。
您有高端cpu和低端gpu，或者gpu驱动程序被卡住了，或者gpu的pci-e接口卡在pci-e 1.1 @ 4x带宽上，因此主机和设备之间的阵列副本受到限制。
另一方面，此优化版本：
__kernel void lowpass(__global __read_only float *Array,__constant float *coefficients, __global __write_only float *Output) {
int globalId = get_global_id(0);
float sum=0.0f;
int min_i= max(64,globalId)-64;
int max_i= min_i+65;
for (int i=min_i; i< max_i; i++)
{
sum +=Array[i]*coefficients[globalId-i];
}
Output[globalId]=sum;
}
对于cpu（8个计算单元）的时间少于150毫秒，对于gpu（25个计算单元）的时间少于80毫秒。每个项目的工作量只有65次。使用__constant和__read_only和__write_only参数说明符可以很容易地加速这种低数量的操作，并减少一些整数工作。
对于数组和输出，使用float4而不是float类型可以使cpu和gpu的速度提高％80，因为它们是SIMD类型和矢量计算单元。
该内核的瓶颈是：
一般来说：
还：
它还具有83 GB / s的L1和L2高速缓存带宽，因此除非计划升级计算机，否则就让它在__global驱动程序优化而不是LDS上起作用。 83 + 83 = 166 GB / s带宽。你可以试试。就银行冲突而言，也许二比二比交替更好。
使用系数作为__constant（166 GB / s）和数组作为__global应该给您166 + 83 = 249 GB / s的组合带宽。
每个系数元素每个线程仅使用一次，因此我不建议使用专用寄存器（499 GB / s）
本文收集自互联网，转载请注明来源。
如有侵权，请联系 [email protected] 删除。
我来说两句