OpenCL矩阵乘法速度

我写了一个小的OpenCL应用程序来计算两个矩阵的乘积。现在我注意到，如果矩阵的大小超过8192 x 8192，就会出现显著的性能下降(16384 x 16384的计算速度要慢80倍)，甚至串行实现也要快5倍以上。下面是主机代码:

/*Make some includes and definitions here*/
#include "stdafx.h"
#include <CL/cl.hpp>
#include <vector>
#include <iostream>
#include "util.hpp" // utility library
#define __CL_ENABLE_EXCEPTIONS
#define ROWS (16384)    // ROWS of vectors a, b, and c
#define COLUMNS (16384)
/*~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~*/
#include "metrics.h"
/*Start main()*/
int main(void)
{
    int A;
    // Fill vectors X and Y with random float values
    float* h_x = new float[ROWS*COLUMNS];
    for (int i = 0; i < ROWS; ++i){
        for (int j = 0; j < COLUMNS; ++j){
            h_x[j + i*COLUMNS] = rand() / (float)RAND_MAX;;
        }
    }
    float* h_y = new float[ROWS*COLUMNS];
    for (int i = 0; i < ROWS; ++i){
        for (int j = 0; j < COLUMNS; ++j){
            h_y[j + i*COLUMNS] = rand() / (float)RAND_MAX;;
        }
    }
    float* h_s = new float[ROWS*COLUMNS];
    for (int i = 0; i < ROWS; ++i){
        for (int j = 0; j < COLUMNS; ++j){
            h_s[j + i*COLUMNS] = 0.0;
        }
    }
    /*~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~*/
    // Get all platforms (drivers)
    std::vector<cl::Platform> all_platforms;
    cl::Platform::get(&all_platforms);

    if (all_platforms.size() == 0){ // Check for issues
        std::cout << " No platforms found. Check OpenCL installation!n";
        exit(1);
    }
    cl::Platform default_platform = all_platforms[0];
    std::cout << "Using platform: " << default_platform.getInfo<CL_PLATFORM_NAME>() << "n";
    // Get default device of the default platform
    std::vector<cl::Device> all_devices;
    default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
    if (all_devices.size() == 0){ // Check for issues
        std::cout << " No devices found. Check OpenCL installation!n";
        exit(1);
    }
    cl::Device default_device = all_devices[0];
    std::cout << "Using device: " << default_device.getInfo<CL_DEVICE_NAME>() << "n";
    // Create an OpenCL context
    cl::Context context({ default_device });
    cl::Program program(context, util::loadProgram("saxy_kernel.cl"), true);
    if (program.build({ default_device }) != CL_SUCCESS){
        std::cout << " Error building: " << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device) << "n";
        getchar();
        exit(1);
    }
    // create buffers on the device
    cl::Buffer buffer_X(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_Y(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_S(context, CL_MEM_READ_WRITE, sizeof(float)* ROWS*COLUMNS);
    cl::Buffer buffer_A(context, CL_MEM_READ_WRITE, sizeof(int));
    //create queue to which we will push commands for the device.
    cl::CommandQueue queue(context, default_device);
    //write arrays A and B to the device
    queue.enqueueWriteBuffer(buffer_X, CL_TRUE, 0, sizeof(float)* ROWS*COLUMNS, &h_x[0]);
    queue.enqueueWriteBuffer(buffer_Y, CL_TRUE, 0, sizeof(float)* ROWS*COLUMNS, &h_y[0]);
    queue.enqueueWriteBuffer(buffer_A, CL_TRUE, 0, sizeof(int), &A);
    StartCounter();
    //run the kernel
    cl::Kernel kernel_add = cl::Kernel(program, "simple_add");
    kernel_add.setArg(0, buffer_X);
    kernel_add.setArg(1, buffer_Y);
    kernel_add.setArg(2, buffer_S);
    kernel_add.setArg(3, buffer_A);
    cl::NDRange global(ROWS*COLUMNS);
    queue.enqueueNDRangeKernel(kernel_add, cl::NullRange, global, cl::NullRange);
    queue.finish();
    std::cout << "Kernel execution time: " << GetCounter() << "ms n";
    //read result C from the device to array C
    queue.enqueueReadBuffer(buffer_S, CL_TRUE, 0, sizeof(float)*ROWS*COLUMNS, &h_s[0]);

    /*Print vectors
    std::cout << "nMatrix #1: n";
    for (int i = 0; i<ROWS*COLUMNS; i++){

            std::cout << "" << h_x[i] << "t ";
    }
    std::cout << "nnMatrix #2: n";
    for (int i = 0; i<ROWS*COLUMNS; i++){

            std::cout << "" << h_y[i] << "t ";
    }
    std::cout << "nnResult: n";
    for (int i = 0; i<ROWS*COLUMNS; i++){

            std::cout << "" << h_s[i] << "t ";
    }*/
    getchar();
    return 0;
}

内核:

__kernel void kernel simple_add(
   __global float* X, 
   __global float* Y, 
   __global float* S, 
   __global int *A){
   S[get_global_id(0)] = X[get_global_id(0)] * Y[get_global_id(0)];
}

你能给我解释一下原因吗?我知道如果我执行一些算法优化，我可以获得更好的性能，但我试图弄清楚这是"幼稚"实现的阈值，还是我做错了什么(不正确地将工作分配给组)。

EDIT:因为我在评论中被问到，我正在运行内核的GPU是AMD R9 270/2GB RAM。CPU为i7-4771，内存为8GB。

#define NUM_PER_THREAD 16
__kernel void kernel simple_add(
 __global float* X, 
 __global float* Y, 
 __global float* S, 
 __global int *A)
{
   for(i = 0; i < NUM_PER_THREAD; i++)
   {
      size_t index = get_global_id(0)*NUM_PER_THREAD + i;
      S[index] = X[index] * Y[index];
   }
}