0

OpenCL を使用すると、Radeon 7970 から i5 デスクトップのメイン メモリに 7MB/秒を超えるデータを取り込めないようです。

#include <iostream>
#include <Windows.h>
#include <CL/cl.h>

int main(int argc, char ** argv)
{
    cl_platform_id platform;
    clGetPlatformIDs(1, &platform, NULL);
    cl_device_id device;
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, NULL);
    cl_command_queue queue = clCreateCommandQueue(context, device, 0, NULL);
    const char *source =
    "__kernel void copytest(__global short* dst) {\n"
    "    __local short buff[1024];\n"
    "    for (int i = 0; i < 1024; i++) {\n"
    "        for (int j = 0; j < 1024; j++)\n"
    "            buff[j] = j;\n"
    "        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);\n"
    "    }\n"
    "}\n";
    cl_program program = clCreateProgramWithSource(context, 1, &source, NULL, NULL);
    clBuildProgram( program, 1, &device, NULL, NULL, NULL);
    cl_kernel kernel = clCreateKernel( program, "copytest", NULL);
    cl_mem buf = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 1024 * 1024 * 2, NULL, NULL);
    const size_t global_work_size = 1;
    clSetKernelArg(kernel, 0, sizeof(buf), (void*)&buf);
    LARGE_INTEGER pcFreq = {}, pcStart = {}, pcEnd = {};
    QueryPerformanceFrequency(&pcFreq);
    QueryPerformanceCounter(&pcStart);
    clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);
    clFinish(queue);
    QueryPerformanceCounter(&pcEnd);
    std::cout << 2.0 * pcFreq.QuadPart / (pcEnd.QuadPart-pcStart.QuadPart) << "MB/sec";
}

ご覧のとおり、すべて単一の作業単位で動作しています。async_work_group_copy() を複数 (64) の作業単位に分散されたループに置き換えてみましたが、役に立ちませんでした。

Radeon から 7MB/秒よりも速くメモリを引き出す方法はありますか? 数百 MB/秒に興味があります。NVidiaの方が速いでしょうか？

4

3 に答える 3

1

問題は、GPU で 1 つのスレッドしか使用していないため、数千のスレッドがアイドル状態のままになっていることです。この場合、より高速な速度を実現するためにできることが 2 つあります。

まず、ワークグループでより多くのスレッドを使用してみてください:

__kernel void copytest(__global short* dst) {
    __local short buff[1024];
    for (int i = 0; i < 1024; i++) {
        for (int j = get_local_id(0); j < 1024; j+= get_local_size(0))
            buff[j] = j;
        barrier(CLK_LOCAL_MEM_FENCE);
        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
    }
}

次に、ワークグループのサイズを 256 程度に増やすことができます。

const size_t local_work_size = 256;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

次に、GPU を使用しているため、1 つのワーク グループのみを使用するべきではありません。次のように、より多くのワークグループを使用できます。

__kernel void copytest(__global short* dst) {
    __local short buff[1024];
    for (int i = get_group_id(0); i < 1024; i += get_num_groups(0)) {
        for (int j = get_local_id(0); j < 1024; j+= get_local_size(0))
            buff[j] = j;
        barrier(CLK_LOCAL_MEM_FENCE);
        (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
    }
}

次に、ワークグループの数を増やすことができます。

const size_t local_work_size = 256;
const size_t global_work_size = local_work_size * 32;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_work_size, &local_work_size, 0, NULL, NULL);

願わくば、これがアプリケーションの高速化に役立つことを願っています。

于 2012-10-17T17:28:29.187 に答える
0

ここでは「j」ループは必要ありません。async_work_group_copy は両方の方法で機能します。グローバル スペースとローカル スペースの間でコピーできます。

//kernel will copy 2MB of short* in memory
__kernel void copytest(__global short* dst) {
  __local short buff[1024];
  for (int i = get_group_id(0); i < 1024; i += get_num_groups(0)) {
    (void)async_work_group_copy(buff, &dst[i*1024], 1024, 0);
    (void)async_work_group_copy(&dst[i*1024], buff, 1024, 0);
  }
}

opencl 1.0 仕様では、16kb 以上のローカル メモリを使用できるようにする必要があります (ocl v1.1 以降では 32kb 以上)。多くのデバイスは実際には 32kb を持っています。システムをポーリングし、できる限り使用することをお勧めします。現実的には、他の目的のためにローカル メモリを保存する必要があります。clGetDeviceInfo (CL_DEVICE_LOCAL_MEM_SIZE)を参照してください。

//using 16kb local memory per work group to copy 2MB...
__kernel void copytest(__global short* dst) {
  __local short buff[16384];
  for (int i = get_group_id(0); i < 64; i += get_num_groups(0)) {
    (void)async_work_group_copy(buff, &dst[i*16384], 16384, 0);
    (void)async_work_group_copy(&dst[i*16384], buff, 16384, 0);
  }
}

実行するコピーの量が十分に少ない場合は、作業を完了するのに必要な数の作業グループを正確に使用することで、'i' ループを取り除くことができます。これにより、アセンブリ内の分岐が少なくなります。

//using 32kb and 16 work groups to copy 2MB...
__kernel void copytest(__global short* dst) {
  __local short buff[32768];
  int i = get_group_id(0);
  (void)async_work_group_copy(buff, &dst[i*32768], 32768, 0);
  (void)async_work_group_copy(&dst[i*32768], buff, 32768, 0);
}
于 2012-10-19T01:55:17.410 に答える
0

バッファが正しく割り当てられていることを確認してください: NVIDIA OpenCL プログラミング ガイド (見つけた場合) を読んで、固定メモリの割り当てを確認してください。6GB/秒を達成できる本格的な例があります.AMDにも同じ原則が適用されます. 特に、CL_MEM_ALLOC_HOST_PTR フラグ。

于 2012-10-18T21:03:19.143 に答える