Блог

Я довольно новичок в OpenCL. Я немного научился в колледже, и под «немного» я подразумеваю, что мой преподаватель графики рассказывал нам о GPGPU и OpenCL в течение одного дня (в то время как остальная часть класса была сосредоточена на шейдерах, OpenGL и так далее).

Я взял пример программы и изменил ее для работы с вычислениями, которые я хочу, чтобы она выполнялась. Однако моя программа работает значительно быстрее на процессоре, чем на моем графическом процессоре, и я пытаюсь понять, почему.

Моя программа принимает один входной массив с плавающей запятой и имеет два выходных массива. В однопоточных условиях он имеет три аргумента. Размер входного массива: samplesPerTrace tracesIn sizeof(float) , а размер выходного массива: samplesPerTrace tracesOut sizeof(float) .

В моих тестовых примерах использовались параметры 25000 2500 250, потому что это в среднем размер массивов, которые я буду использовать (возможно, немного выше среднего). Значения заполняются случайным образом.

Вот исходный код, который OpenCL создает и запускает в ядре;

 const char* M_AND_S_OPENCL_SOURCE_TEXT =
"__kernel void sumAllCL(__global const float prestackTraces[],n"
"  __global float stackTracesOut[],n"
"  __global float powerTracesOut[], const unsigned int nTracesOut, const unsigned int nTracesIn,n"
"  const unsigned int samplesPerTrace) {n"
"n"
"  unsigned int k = get_global_id(0);n" // Thread ID
"n"
"  unsigned int kTimesIn = k * nTracesIn;n" // Store repeat ints
"  unsigned int kTimesSamples = k * samplesPerTrace;n"
"n"
"  for (int j = 0; j < ?       ; j  ) {n" // ? position to be replaced (nTracesOut)"
"n"
"    int jTimesSamplesPT = j * samplesPerTrace;n"
"n"
"    for (int i = 0; i < #       ; i  ) {n" // # position to be replaced ()
"n"
"      int valueIndex = i   jTimesSamplesPT;n"
"      float value = prestackTraces[valueIndex];n"
"n"
"      stackTracesOut[i   kTimesSamples]  = value;n"
"      powerTracesOut[i   kTimesSamples]  = (value * value);n"
"n"
"    }n"
"  }n"
"}n";
  

Обратите внимание, что ? и # заменяются во время выполнения фиксированными числами, я делаю это, потому что думал, что это поможет компилятору развернуть rl

С указанными выше параметрами (25000 2500 250 ~10 <1 или 2>) для завершения программы моему процессору требуется около 0,6 секунды, а графическому процессору — около 40 секунд. Это большая разница. К вашему сведению, я возился с 4-м параметром, чтобы увидеть, какое значение выполняется быстрее, что и подразумевается под ~ 10.

Моя видеокарта — MSI Radeon R9 390X 8GB, получившая название Hawaii. Когда я распечатываю информацию OpenCL об обоих моих устройствах, это то, что я получаю:

 OpenCL Platform 0: AMD Accelerated Parallel Processing
 ----- OpenCL Device # 0: Hawaii-----
Gflops: 47.520000
Max Clock Frequency: 1080
Max Compute Units: 44
Max Work Group Size: 256
   MEMORY...
Total Memory of Device: 8.000G   (CL_DEVICE_GLOBAL_MEM_SIZE)
Local Memory of Device: 32.000K   (CL_DEVICE_LOCAL_MEM_SIZE)
Max Memory Object Allocation: 3.999G   (CL_DEVICE_MAX_MEM_ALLOC_SIZE)
Cache Size: 16.000K   (CL_DEVICE_GLOBAL_MEM_CACHE_SIZE)
Cacheline Size: 64 bytes   (CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE)
   VERSIONS...
Device Vendor: Advanced Micro Devices, Inc.
Device Version: OpenCL 2.0 AMD-APP (2117.13)
Driver Version: 2117.13 (VM)
Device OpenCL Version: OpenCL C 2.0
 ----- OpenCL Device # 1: Intel(R) Core(TM) i7-6700K CPU ? 4.00GHz-----
Gflops: 32.064000
Max Clock Frequency: 4008
Max Compute Units: 8
Max Work Group Size: 1024
   MEMORY...
Total Memory of Device: 15.967G   (CL_DEVICE_GLOBAL_MEM_SIZE)
Local Memory of Device: 32.000K   (CL_DEVICE_LOCAL_MEM_SIZE)
Max Memory Object Allocation: 3.1028G   (CL_DEVICE_MAX_MEM_ALLOC_SIZE)
Cache Size: 32.000K   (CL_DEVICE_GLOBAL_MEM_CACHE_SIZE)
Cacheline Size: 64 bytes   (CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE)
   VERSIONS...
Device Vendor: GenuineIntel
Device Version: OpenCL 1.2 AMD-APP (2117.13)
Driver Version: 2117.13 (sse2,avx)
Device OpenCL Version: OpenCL C 1.2
  

Вот код, относящийся к OpenCL. Я бы опубликовал полный минимальный проверяемый полный пример, однако это ставит меня выше предела символов.

 /*
 * Prints the given int (numToInsert) at location inside chars.
 */
void PrintIntInStr(char* chars, int location, int numToInsert) {

  std::stringstream strs;
  strs << numToInsert;
  std::string temp_str = strs.str();
  char const* numToChars = temp_str.c_str();

  int numberLength = strlen(numToChars);

  int w;
  for (w = 0; w < numberLength; w  ) {
    chars[location   w] = numToChars[w];
  }
}

/*
 * Initialize fastest OpenCL device.
 */
int InitOpenCL(int verbose, cl_int deviceType) {

  cl_uint Nplat;
  cl_int  err;
  char name[1024];
  int  MaxGflops = -1;

  cl_platform_id winnerPlatform = 0;

  // Reset (TODO)
  _deviceID = NULL;
  _context = NULL;
  _queue = NULL;

  // Get platforms
  cl_platform_id platforms[4];
  if (clGetPlatformIDs(4, platforms, amp;Nplat)) Fatal("Cannot get number of OpenCL platformsn");
  else if (Nplat<1) Fatal("No OpenCL platforms foundn");

  // Loop over platforms
  for (unsigned int platform = 0; platform < Nplat; platform  ) {

    if (clGetPlatformInfo(platforms[platform], CL_PLATFORM_NAME, sizeof(name), name, NULL)) Fatal("Cannot get OpenCL platform namen");
    if (verbose) printf("OpenCL Platform %d: %sn", platform, name);

    // Get GPU device IDs
    cl_uint Ndev;
    cl_device_id id[4];
    if (clGetDeviceIDs(platforms[platform], deviceType, 4, id, amp;Ndev))
      Fatal("Cannot get number of OpenCL devices: %dn", platform);
    else if (Ndev < 1) Fatal("No OpenCL devices found.n");

    // Find the fastest device
    for (unsigned int devId = 0; devId < Ndev; devId  ) {

      // Print informatio about the device
      cl_uint compUnits, freq, cacheLineSize;
      cl_ulong memSize, maxAlloc, localMemSize, globalCacheSize;
      size_t maxWorkGrps;
      char deviceVendor[50];
      char deviceVersion[50];
      char driverVersion[50];
      char deviceOpenCLVersion[50];

      // Computing Power...
      if (clGetDeviceInfo(id[devId], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compUnits), amp;compUnits, NULL)) Fatal("Cannot get OpenCL device unitsn");
      if (clGetDeviceInfo(id[devId], CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(freq), amp;freq, NULL)) Fatal("Cannot get OpenCL device frequencyn");
      if (clGetDeviceInfo(id[devId], CL_DEVICE_NAME, sizeof(name), name, NULL)) Fatal("Cannot get OpenCL device namen");
      if (clGetDeviceInfo(id[devId], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(maxWorkGrps), amp;maxWorkGrps, NULL)) Fatal("Cannot get OpenCL max work group sizen");
      // Memory...
      if (clGetDeviceInfo(id[devId], CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(memSize), amp;memSize, NULL)) Fatal("Cannot get OpenCL memory size.n");
      if (clGetDeviceInfo(id[devId], CL_DEVICE_LOCAL_MEM_SIZE, sizeof(localMemSize), amp;localMemSize, NULL)) localMemSize = 0;
      if (clGetDeviceInfo(id[devId], CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(maxAlloc), amp;maxAlloc, NULL)) Fatal("Cannot get OpenCL memory size.n");
      if (clGetDeviceInfo(id[devId], CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, sizeof(globalCacheSize), amp;globalCacheSize, NULL)) globalCacheSize = 0;
      if (clGetDeviceInfo(id[devId], CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, sizeof(cacheLineSize), amp;cacheLineSize, NULL)) cacheLineSize = 0;
      // Versions...
      clGetDeviceInfo(id[devId], CL_DEVICE_VENDOR, sizeof(deviceVendor), deviceVendor, NULL);
      clGetDeviceInfo(id[devId], CL_DEVICE_VERSION, sizeof(deviceVersion), deviceVersion, NULL);
      clGetDeviceInfo(id[devId], CL_DRIVER_VERSION, sizeof(driverVersion), driverVersion, NULL);
      clGetDeviceInfo(id[devId], CL_DEVICE_OPENCL_C_VERSION, sizeof(deviceOpenCLVersion), deviceOpenCLVersion, NULL);

      int Gflops = compUnits * freq;

      if (verbose) printf(" ----- OpenCL Device # %d: %s-----n"
        "Gflops: %fn"
        "Max Clock Frequency: %dn"
        "Max Compute Units: %dn"
        "Max Work Group Size: %zun"
        "   MEMORY...n"
        "Total Memory of Device: %s   (CL_DEVICE_GLOBAL_MEM_SIZE)n"
        "Local Memory of Device: %s   (CL_DEVICE_LOCAL_MEM_SIZE)n"
        "Max Memory Object Allocation: %s   (CL_DEVICE_MAX_MEM_ALLOC_SIZE)n"
        "Cache Size: %s   (CL_DEVICE_GLOBAL_MEM_CACHE_SIZE)n"
        "Cacheline Size: %s   (CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE)n"
        "   VERSIONS...n"
        "Device Vendor: %sn"
        "Device Version: %sn"
        "Driver Version: %sn"
        "Device OpenCL Version: %sn",
        devId,
        name,
        (1e-3 * Gflops),
        freq,
        compUnits,
        maxWorkGrps,
        byteConverter((unsigned long)memSize),
        byteConverter((unsigned long)localMemSize),
        byteConverter((unsigned long)maxAlloc),
        byteConverter((unsigned long)globalCacheSize),
        byteConverter((unsigned long)cacheLineSize),
        deviceVendor,
        deviceVersion,
        driverVersion,
        deviceOpenCLVersion);

      if(Gflops > MaxGflops)
      {
        _deviceID = id[devId];
        MaxGflops = Gflops;

        winnerPlatform = platforms[platform];
      }
    }
  }

  // Print fastest device info (TODO: don't get name twice)
  if (clGetDeviceInfo(_deviceID, CL_DEVICE_NAME, sizeof(name), name, NULL)) Fatal("Cannot get OpenCL device namen");
  printf("n   Selected Fastest Open CL Device: %s (#%lu)n", name, (unsigned long)_deviceID);

  // Check thread count
  size_t mwgs;
  if (clGetDeviceInfo(_deviceID, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(mwgs), amp;mwgs, NULL))
    Fatal("Cannot get OpenCL max work group sizen");

  // Create OpenCL context for fastest device
  cl_context_properties cps[3] =
  {
    CL_CONTEXT_PLATFORM,
    (cl_context_properties)winnerPlatform,
    (cl_context_properties)0
  };
  _context = clCreateContextFromType(cps, deviceType, NULL, NULL, amp;err);
  if (!_context || err) Fatal("Cannot create OpenCL Contextn");

  // Properties for create command queue; currently nothing
  // cl_command_queue_properties *propers;
  cl_command_queue_properties prop = 0;
  //prop |= CL_QUEUE_PROFILING_ENABLE;
  //prop |= CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
  // propers = amp;prop;

   _queue = clCreateCommandQueueWithProperties(_context, _deviceID, amp;prop, amp;err); //  Create OpenCL command queue for fastest device
  // _queue = clCreateCommandQueue(_context, _deviceID, amp;prop, amp;err);
  if (!_queue || err) {
    if (err == CL_INVALID_CONTEXT) Fatal("Cannot create OpenCL command cue: CL_INVALID_CONTEXTn");
    else if (err == CL_INVALID_DEVICE) Fatal("Cannot create OpenCL command cue: CL_INVALID_DEVICEn");
    else if (err == CL_INVALID_VALUE) Fatal("Cannot create OpenCL command cue: CL_INVALID_VALUEn");
    else if (err == CL_INVALID_QUEUE_PROPERTIES) Fatal("Cannot create OpenCL command cue: CL_INVALID_QUEUE_PROPERTIESn");
    else if (err == CL_OUT_OF_RESOURCES) Fatal("Cannot create OpenCL command cue: CL_OUT_OF_RESOURCESn");
    else if (err == CL_OUT_OF_HOST_MEMORY) Fatal("Cannot create OpenCL command cue: CL_OUT_OF_HOST_MEMORYn");
    else if (!_queue) Fatal("Cannot create OpenCL command cue: !queuen");
    else Fatal("Cannot create OpenCL command cue: ?????n");
  }

  if (_VERBOSE) printf("Init complete.n");

  return mwgs;
}

/*
 * Modify the source text to fit this run.
 */
char* ModifySourceText(unsigned int nTracesIn, unsigned int samplesPerT) {

  size_t sourceSize = strlen(M_AND_S_OPENCL_SOURCE_TEXT)   1;
  char* moveStackSourceCode = new char[sourceSize];
  strncpy(moveStackSourceCode, M_AND_S_OPENCL_SOURCE_TEXT, sourceSize);
  moveStackSourceCode[sourceSize] = '';

  // Print out the locations of the characters where we should insert other text if asked to do so
  if (_FIND_INSERT_LOCATIONS) {
    size_t z;
    for (z = 0; z < sourceSize; z  ) {
      if (moveStackSourceCode[z] == '@') {
        printf("Found @ at position %zun", z);
        break;
      }
    }
    for (z = 0; z < sourceSize; z  ) {
      if (moveStackSourceCode[z] == '#') {
        printf("Found # at position %zun", z);
        break;
      }
    }
  }

  // Insert the digit that for loops go to inside of the source
  PrintIntInStr(moveStackSourceCode, INSERT_LOCATION_1, nTracesIn);
  PrintIntInStr(moveStackSourceCode, INSERT_LOCATION_2, samplesPerT);

  // Print the modified source code if verbose
  if (_FIND_INSERT_LOCATIONS) {
    printf("n   GPU Source Code: n");
    printf("%sn", moveStackSourceCode);
  }

  return moveStackSourceCode;
}

/*
 * Wait for event and then release it.
 */
static void WaitForEventAndRelease(cl_event *event) {

  printf("WaitForEventAndRelease()n");

  cl_int status = CL_SUCCESS;

  status = clWaitForEvents(1, event);
  if (status) Fatal("clWaitForEvents Failed with Error Code");

  printf("About to release event...n");

  status = clReleaseEvent(*event);
  if (status) Fatal("clReleaseEvent Failed with Error Code");
}


// Runs the program via open CL
static double RunOpenCL(float prestackTracesArray[], float stackTracesOut1DArray[], float powerTracesOut1DArray[],
  unsigned int nTracesOut, unsigned int nTracesIn, unsigned int samplesPerT,
  size_t inXsamples, size_t outXsamples,
  unsigned int localThreadCount)
{

  cl_int err;

  // Get the source code
  char* modifiedGpuSource = ModifySourceText(nTracesIn, samplesPerT);

  // Allocate device memory
  // CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR | CL_MEM_USE_PERSISTENT_MEM_AMD (?)
  // Input...
  cl_mem prestackTracesCL = clCreateBuffer(_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    inXsamples * sizeof(cl_float), prestackTracesArray, amp;err);
  if (err) FatalBufferCreation("Prestack traces", err);
  // Output... TODO: How do we know that the output is zeroed out?
  cl_mem stackTracesOutCL = clCreateBuffer(_context, CL_MEM_WRITE_ONLY,
    outXsamples * sizeof(cl_float), NULL, amp;err);
  if (err) FatalBufferCreation("Stack traces", err);
  cl_mem powerTracesOutCL = clCreateBuffer(_context, CL_MEM_WRITE_ONLY,
    outXsamples * sizeof(cl_float), NULL, amp;err);
  if (err) FatalBufferCreation("Power traces", err);

  // Compile the source code
  char* gpuSourceText[1];
  gpuSourceText[0] = modifiedGpuSource;
  size_t sourceLength[1];
  sourceLength[0] = strlen(modifiedGpuSource);
  cl_program moveoutAndStackCLProgram = clCreateProgramWithSource(_context, 1, (const char**)gpuSourceText,
    (const size_t*)sourceLength, amp;err);
  if (err != CL_SUCCESS) {
    if (err == CL_INVALID_CONTEXT) Fatal("Cannot create program: CL_INVALID_CONTEXTn");
    else if (err == CL_INVALID_VALUE) Fatal("Cannot create program: CL_INVALID_VALUEn");
    else if (err == CL_OUT_OF_HOST_MEMORY) Fatal("Cannot create program: CL_OUT_OF_HOST_MEMORYn");
    else Fatal("Cannot create program_S %dn", err);
  }

  // Build the program
  cl_int buildCode = clBuildProgram(moveoutAndStackCLProgram, 0, NULL, NULL, NULL, NULL);
  if (buildCode != CL_SUCCESS) {
    // Attempt to get compile errors
    char log[1048576];
    if (clGetProgramBuildInfo(moveoutAndStackCLProgram, _deviceID, CL_PROGRAM_BUILD_LOG, sizeof(log), log, NULL)) {
      log[0] = ''; // Failed to get the log file
    }

    if (buildCode == CL_INVALID_PROGRAM) Fatal("Cannot build program: CL_INVALID_PROGRAMn%s", log);
    else if (buildCode == CL_INVALID_VALUE) Fatal("Cannot build program: CL_INVALID_VALUEn%s", log);
    else if (buildCode == CL_INVALID_DEVICE) Fatal("Cannot build program: CL_INVALID_DEVICEn%s", log);
    else if (buildCode == CL_INVALID_BINARY) Fatal("Cannot build program: CL_INVALID_BINARYn%s", log);
    else if (buildCode == CL_INVALID_BUILD_OPTIONS) Fatal("Cannot build program: CL_INVALID_BUILDn_OPTIONSn%s", log);
    else if (buildCode == CL_INVALID_OPERATION) Fatal("Cannot build program: CL_INVALID_OPERATIONn%s", log);
    else if (buildCode == CL_COMPILER_NOT_AVAILABLE) Fatal("Cannot build program: CL_COMPILER_NOT_AVAILABLEn%s", log);
    else if (buildCode == CL_BUILD_PROGRAM_FAILURE) Fatal("Cannot build program: CL_BUILD_PROGRAM_FAILUREn%s", log);
    else if (buildCode == CL_INVALID_OPERATION) Fatal("Cannot build program: CL_INVALID_OPERATIONn%s", log);
    else if (buildCode == CL_OUT_OF_HOST_MEMORY) Fatal("Cannot build program: CL_OUT_OF_HOST_MEMORYn%s", log);
    else Fatal("Cannot build program: %dn%s", buildCode, log);
  }

  // Compile the source code amp; build the kernel
  cl_kernel kernel = clCreateKernel(moveoutAndStackCLProgram, "sumAllCL", amp;err);
  if (err) {
    if (err == CL_INVALID_PROGRAM) Fatal("Cannot create kernel: CL_INVALID_PROGRAMn");
    else if (err == CL_INVALID_PROGRAM_EXECUTABLE) Fatal("Cannot create kernel: CL_INVALID_PROGRAM_EXECUTABLEn");
    else if (err == CL_INVALID_KERNEL_NAME) Fatal("Cannot create kernel: CL_INVALID_KERNEL_NAMEn");
    else if (err == CL_INVALID_KERNEL_DEFINITION) Fatal("Cannot create kernel: CL_INVALID_KERNEL_DEFINITIONn");
    else if (err == CL_INVALID_VALUE) Fatal("Cannot create kernel: CL_INVALID_VALUEn");
    else if (err == CL_OUT_OF_HOST_MEMORY) Fatal("Cannot create kernel: CL_OUT_OF_HOST_MEMORn");
    else Fatal("Cannot create kernel: %dn", err);
  }

  // Set program parameters
  cl_int returnValArgSet;
  returnValArgSet = clSetKernelArg(kernel, 0, sizeof(cl_mem), amp;prestackTracesCL);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("prestackTracesCL", returnValArgSet);
  returnValArgSet = clSetKernelArg(kernel, 1, sizeof(cl_mem), amp;stackTracesOutCL);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("stackTracesOutCL", returnValArgSet);
  returnValArgSet = clSetKernelArg(kernel, 2, sizeof(cl_mem), amp;powerTracesOutCL);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("powerTracesOutCL", returnValArgSet);
  returnValArgSet = clSetKernelArg(kernel, 3, sizeof(unsigned int), amp;nTracesOut);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("nTracesOut", returnValArgSet);
  returnValArgSet = clSetKernelArg(kernel, 4, sizeof(unsigned int), amp;nTracesIn);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("nTracesIn", returnValArgSet);
  returnValArgSet = clSetKernelArg(kernel, 5, sizeof(unsigned int), amp;samplesPerT);
  if (returnValArgSet != CL_SUCCESS) FatalSetArgs("samplesPerT", returnValArgSet);

  // TODO: verbose
  printf("About to run Kernel...n");

  // Start timer TODO: move?
  double runTime = GetTime();

  // Run the kernel (amp; also set the number of threads)
  cl_event runEvent;
  size_t Global[1] = { nTracesOut };
  size_t Local[1]  = { localThreadCount };
  if (localThreadCount > 0) err = clEnqueueNDRangeKernel(_queue, kernel, 1, NULL, Global, Local, 0, NULL, amp;runEvent);
  else err = clEnqueueNDRangeKernel(_queue, kernel, 1, NULL, Global, NULL, 0, NULL, amp;runEvent);
  if (err) {
    if (err == CL_INVALID_PROGRAM_EXECUTABLE) {
      Fatal("Cannot run Kernel: No successfully built program executable available.n");
    } else if (err == CL_INVALID_COMMAND_QUEUE) {
      Fatal("Cannot run Kernel: Command_queue is not a valid command-queue.n");
    } else if (err == CL_INVALID_KERNEL) {
      Fatal("Cannot run Kernel: Kernel is not a valid kernel object.n");
    } else if (err == CL_INVALID_CONTEXT) {
      Fatal("Cannot run Kernel: Context associated with command_queue and kernel is not the same or if "
        "the context associated with command_queue and events in event_wait_list are not the same.n");
    } else if (err == CL_INVALID_KERNEL_ARGS) {
      Fatal("Cannot run Kernel: Kernel argument values have not been specified.n");
    } else if (err == CL_INVALID_WORK_DIMENSION) {
      Fatal("Cannot run Kernel: work_dim is not a valid value (must be between 1 and 3).n");
    } else if (err == CL_INVALID_WORK_GROUP_SIZE) {
      Fatal("Cannot run Kernel: local_work_size is specified and number of work-items specified by global_work_size "
        "is not evenly divisable by size of work-group given by local_work_size or does not match the "
        "work-group size specified for kernel using the __attribute__((reqd_work_group_size(X, Y, Z))) "
        "qualifier in program source.n");
    } else if (err == CL_INVALID_WORK_ITEM_SIZE) {
      Fatal("Cannot run Kernel: If the number of work-items specified in any of local_work_size[0], ... "
        "local_work_size[work_dim - 1] is greater than the corresponding values specified "
        "by CL_DEVICE_MAX_WORK_ITEM_SIZES[0], .... CL_DEVICE_MAX_WORK_ITEM_SIZES[work_dim - 1]. .n");
    } else if (err == CL_INVALID_GLOBAL_OFFSET) {
      Fatal("Cannot run Kernel: Global_work_offset is not NULL.n");
    } else if (err == CL_OUT_OF_RESOURCES) {
      Fatal("Cannot run Kernel: CL_OUT_OF_RESOURCES.n");
    } else if (err == CL_MEM_OBJECT_ALLOCATION_FAILURE) {
      Fatal("Cannot run Kernel: Failure to allocate memory for data store associated with image or buffer "
        "objects specified as arguments to kernel.n");
    } else if (err == CL_INVALID_EVENT_WAIT_LIST) {
      Fatal("Cannot run Kernel: event_wait_list is NULL and num_events_in_wait_list > 0, or event_wait_list "
        "is not NULL and num_events_in_wait_list is 0, or if event objects in event_wait_list "
        "are not valid events..n");
    } else if (err == CL_OUT_OF_HOST_MEMORY) {
      Fatal("Cannot run Kernel: Failure to allocate resources required by the OpenCL implementation on the host.n");
    } else {
      Fatal("Cannot run Kernel: Unknown Error. (clEnqueueNDRangeKernel)");
    }
  }

  // Flush the program amp; wait for the program to finish executing
  if (clFlush(_queue)) printf("Flush Fail (Run)");
  WaitForEventAndRelease(amp;runEvent);

  // Copy the end result back to CPU memory side
  if (clEnqueueReadBuffer(_queue, stackTracesOutCL, CL_TRUE, 0, outXsamples * sizeof(cl_float), stackTracesOut1DArray, 0, NULL, NULL))
    Fatal("Cannot copy stackTracesOutCL from device to hostn");
  if (clEnqueueReadBuffer(_queue, powerTracesOutCL, CL_TRUE, 0, outXsamples * sizeof(cl_float), powerTracesOut1DArray, 0, NULL, NULL))
    Fatal("Cannot copy powerTracesOutCL from device to hostn");

  // Release kernel and program
  if (clReleaseKernel(kernel)) Fatal("Cannot release kerneln");
  if (clReleaseProgram(moveoutAndStackCLProgram)) Fatal("Cannot release programn");

  // Free device memory
  clReleaseMemObject(prestackTracesCL);
  clReleaseMemObject(stackTracesOutCL);
  clReleaseMemObject(powerTracesOutCL);

  // Release the context and queue
  clReleaseCommandQueue(_queue);
  clReleaseContext(_context);

  // Return the time it took to run this program
  return runTime;
}

double RunProg(unsigned int samplesPerTrace, unsigned int nTracesIn, unsigned int nTracesOut,
  unsigned int localThreadCount, unsigned int deviceType) {

  // Stores sizes of the various arrays
  size_t tracesInxSample = nTracesIn * samplesPerTrace;
  size_t tracesOutxSample = nTracesOut * samplesPerTrace;

  // Allocate arrays
  float* prestackTraces1D = (float*)malloc(tracesInxSample * sizeof(float));
  float* stackTracesOut1Dgpu = (float*)calloc(tracesOutxSample, sizeof(float)); // output; zero-out
  float* powerTracesOut1Dgpu = (float*)calloc(tracesOutxSample, sizeof(float)); // output; zero-out

  // Count how much memory all of this is
  if (_VERBOSE)
  {
    // Make sure it is consistent with above allocation
    unsigned long allocatedMemory = 0;
    allocatedMemory  = tracesInxSample * sizeof(float);
    allocatedMemory  = tracesOutxSample * sizeof(float);
    allocatedMemory  = tracesOutxSample * sizeof(float);

    printf("TOTAL MEMORY ALLOCATED: %sn", byteConverter(allocatedMemory));
    printf("Input Array Sizes: %sn", byteConverter((unsigned int)(tracesInxSample * sizeof(float))));
    printf("Output Array Sizes: %sn", byteConverter((unsigned int)(tracesOutxSample * sizeof(float))));
  }

  // Fill in array with randoms
  RandomFillArray(prestackTraces1D, (unsigned int)tracesInxSample);

  // Init OpenCL using the desired device type
  double preInitTime = GetTime();
  int maxWorkGroupSize;
  if (deviceType == 0) maxWorkGroupSize = InitOpenCL(_VERBOSE, CL_DEVICE_TYPE_ALL);
  else if (deviceType == 1) maxWorkGroupSize = InitOpenCL(_VERBOSE, CL_DEVICE_TYPE_GPU);
  else maxWorkGroupSize = InitOpenCL(_VERBOSE, CL_DEVICE_TYPE_CPU);
  printf("Max work size for the device is: %dn", maxWorkGroupSize);

  // --- ACTUAL TEST ---
  // Run OpenCL
  double startTime = GetTime();
  double runTime = RunOpenCL(prestackTraces1D, stackTracesOut1Dgpu, powerTracesOut1Dgpu, // arrays
    nTracesOut, nTracesIn, samplesPerTrace, // ints
    tracesInxSample, tracesOutxSample,
    localThreadCount); // samples

  // Display run time
  double endTime = GetTime();
  printf("Elapsed Time:      %fsecsn", (endTime - runTime));
  printf("                   %fsecs (Before Function Call)n", (endTime - startTime));
  printf("                   %fsecs (Including Init)nn", (endTime - preInitTime));

  // Free the 1D arrays
  free(powerTracesOut1Dgpu);
  free(stackTracesOut1Dgpu);
  free(prestackTraces1D);

  return (endTime - startTime);
}
  

Моя первая мысль о том, почему он работает намного медленнее на моем графическом процессоре, чем на моем процессоре, заключается в том, что, возможно, это потому, что я передаю так много данных на видеокарту, прежде чем что-либо запустится. Возможно, лучшая реализация включала бы разделение рабочей нагрузки на несколько запусков, чтобы код мог выполняться во время обработки большего количества данных (я полагаю, это так).). Однако теперь, когда я думаю об этом, это почти наверняка неверно, потому что, как я уже сказал, я написал эту программу на основе примера, и в этом примере выполнялось умножение матриц, и этот пример выполняется намного быстрее на графическом процессоре, чем на моем процессоре. Я действительно не знаю, в чем разница.

Проблема заключалась в кэшировании; я много читал и записывал из массива. Итак, я создал версию, которая записывает в локальные переменные как можно больше перед записью в массив, и теперь она работает намного быстрее на графическом процессоре.