2013-09-01 26 views
0

我目前正在试验使用GPU的OpenCL代码的性能,并在CPU上使用C++。我编写了计算总和z = x + y的程序,其中z,x和y是GPU和CPU的二维数组(矩阵)。在测试这些程序之后,我发现由于GPU和CPU之间的PCI总线中的数据传输缓慢,所以CPU在计算这一总和方面比GPU更高效。现在我想确定为了使GPU比CPU更有效,还需要多少钱。我打算通过增加z = x + y到z = x + y + y + y + y + ...等来做到这一点。在GPU和CPU上添加2D阵列的性能

仅通过增加此特定问题的总和数量,是否有可能使GPU的使用效率高于CPU?

正如仅供参考:我使用的是nVIDIA GeForce GT 640图形卡和i5 Intel核心CPU。

任何帮助将不胜感激。

编辑:

下面我附着在CPU上我的代码:

int main(int argc, const char * argv[]) 
{ 

    //This value determines the size of the nxn (square array)    
    int n = 1000; 

    //Allocating the memory for the nxn arrays of floats. 
    float **x = (float**)malloc(sizeof(float*)*n); 
    float **y = (float**)malloc(sizeof(float*)*n); 
    float **z = (float**)malloc(sizeof(float*)*n); 


    //Initializing the arrays. 
    for(int i = 0; i<n; i++){ 
     x[i] = (float*)malloc(sizeof(float)*n); 
     y[i] = (float*)malloc(sizeof(float)*n); 
     z[i] = (float*)malloc(sizeof(float)*n); 

     for(int j = 0; j<n; j++){ 
      x[i][j] = i+j; 
      y[i][j] = i+j; 

     } 
    } 

    for(int i = 0; i<n; i++){ 
     for(int j = 0; j<n; j++){ 

      z[i][j] = x[i][j] + y[i][j]; 
      for(int k = 0; k < 100; k++){ 
       z[i][j] += y[i][j]; 
      } 
     } 
    } 

    return 0; 

} 

这里是C++使用OpenCL的:(用于复制数据,并执行对GPU内核)

int n = 1000; 

for(int i = 0; i<n; i++) 
    { 
     //Writing the data from the host to the device 
     err = clEnqueueWriteBuffer(queue, d_xx, CL_TRUE, 0, sizeof(float)*n, h_xx[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not write to buffer d_xx" << std::endl; 
      exit(1); 
     } 

     err = clEnqueueWriteBuffer(queue, d_yy, CL_TRUE, 0, sizeof(float)*n, h_yy[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not write to buffer d_yy" << std::endl; 
      exit(1); 
     } 

     //Setting the Kernel Arguments 
     err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_xx); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_xx." << std::endl; 
      exit(1); 
     } 

     err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_yy); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_yy." << std::endl; 
      exit(1); 
     } 

     err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_zz); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not set kernel argument h_zz." << std::endl; 
     } 

     work_units_per_kernel = n; 

     //Executing the Kernel 
     err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_units_per_kernel, NULL, 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not execute kernel." << std::endl; 
      exit(1); 
     } 

     //Reading the Data from the Kernel 
     err = clEnqueueReadBuffer(queue, d_zz, CL_TRUE, 0, n*(sizeof(float)), h_zz[i], 0, NULL, NULL); 
     if(err != CL_SUCCESS){ 
      std::cout << "Error: Could not read data from kernel." << std::endl; 
      exit(1); 
     } 

    } 

而且最后在GPU上执行的内核代码:

__kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc) 
{ 

    int i = get_global_id(0); 

    d_cc[i] = d_aa[i] + d_bb[i]; 


    for(int j = 0; j < 100; j++){ 
     d_cc[i] += d_bb[i]; 
    } 


} 
+0

你为什么不在这里发布你的代码? GPU如果与你并列,就会获胜。 GPU也可能具有更快的内存。你的程序是否并行?因为一个cpu线程会比一个GPU“线程”执行得更好。 – SigTerm

+0

@SigTerm感谢您的回复。我附加了一些我的代码片段。我希望他们帮助澄清我是否在我的程序中并行化。 – user2736519

+1

在我看来,这是一个很少计算的经典案例,整个操作都是内存绑定的,所以除非您可以对GPU上的数据做更多的事情,否则总线速度将比GPU本身获得的更多。 –

回答

2

对于n = 1000 * 1000,您正在接近复制,操作和复制回值这一点。正如DarkZero指出的那样,全局内​​存并不是最优的,所以如果你可以将全局内存缓存到本地内存或线程内存并使用本地工作组,这对CPU和GPU都将有很大的帮助。

让我们从内核开始吧。来自全局内存的d_cc被引用100次。这种情况下的简单更改是将全局内存缓存到线程内存中,然后在最后将本地内存复制到全局内存中。

__kernel void arraysum(__global const float *d_aa, __global const float *d_bb, __global float *d_cc) 
{ 

    int i = get_global_id(0); 

    float t_d_cc = d_aa[i] + d_bb[i]; //make a thread only version of d_cc 

    for(int j = 0; j < 100; j++){ 
     t_d_cc += d_bb[i]; 
    } 

    d_cc[i] = t_d_cc; //copy the thread only back to global 
} 

另一个变化,依赖于硬件,是缓存d_aa和d_bb到本地内存。这让OpenCL可以利用全局内存中的批量拷贝。这可能会更具挑战性,因为每个OpenCL设备具有不同的大小和可以使用的本地工作组大小的倍数。

例如,我的i5的最大工作组大小为1024,工作组倍数为1,因此我的本地工作组可以是1到1024之间的任何值。我的ATI-7970的值分别为256和64,所以我的本地Worksgroups需要是64,128等,这是更严格的。

__kernel void arraysum(__global const float *d_aa, 
         __local float *l_d_aa, 
         __global const float *d_bb, 
         __local float *l_d_bb, 
         __global float *d_cc, 
         __local float *l_d_cc) 
{ 

//In this example, the global_id(1) is the number of rows and global_id(0) is the columns 
//So when the kernel is called, the local work group size needs to be the size of the 
//number of columns 

int i = get_global_id(1)*get_global_size(0) + get_global_id(0); //Index of the row 
int j = get_local_id(0); 

l_d_aa[get_local_id(0)] = d_aa[i]; 
l_d_bb[get_local_id(0)] = d_bb[i]; 

read_mem_fence(CLK_LOCAL_MEM_FENCE); 

float l_d_cc[get_local_id(0)] = l_d_aa[get_local_id(0)] + l_d_bb[get_local_id(0)]; 

for(int j = 0; j < get_global_size(0); j++){ 
    l_d_cc[get_local_id(0)] += l_d_bb[j]; 
} 

d_cc[i] = l_d_cc[get_local_id(0)]; //copy the thread only back to global 

}

我道歉,如果我得到的算法不对,但希望它传达如何缓存全局内存本地内存。同样,在i5上,本地工作组大小可以是1到1024,但ATI7970仅限于列大小64,128等。

这在概念上要困难得多,但OpenCL的性能很差,使用这种方法时要好得多。

社区,请随时清理内核。

+1

因为累加器是每个工作项目,所以使用本地内存没什么意义。但是,私有内存内核应该工作得很好。内核+非阻塞呼叫+非常高的N值=非常好的速度。 – DarkZeros

+0

我认为这是值得的,如果他可以缓存整个行到本地,他可以使用CPU。使用高速缓存将编译器(gcc或MSVC)代码与CPU上的OpenCL编译代码进行比较会很有趣。 – Austin

1

很多事情让你放慢脚步:

1-滥用使用全局内存。每个全局内存访问速度都是400倍,并且只能使用全局内存(如200个读/写)。全局内存只能用于在开始时读取,并在最后写入,而不能作为中间值。

2-您的N长度很短。 CPU将以1000条指令完成,而GPU中的所有延迟都比这慢得多。由于100MB副本比1字节副本更有效,因此复制操作会有开销。可能CPU编码正在被编译器优化成乘法运算,而GPU编码不能,因为它正在访问像全局变量这样的易失性变量。内存对设备的读取/写入操作非常昂贵,如果将其包含在calc中,CPU将很容易获胜。 OpenCL缓冲区和内核创建也非常昂贵。请注意,您也正在使用阻止写入调用,这比非阻止调用慢得多。