内核运行需要调用 clGetDeviceInfo

Question

我有一些运行良好的 OpenCL 代码，如下所示:

# ocl_helpers.h
typedef struct {
  cl_context context;
  cl_command_queue queue;
  cl_device_id device;
  cl_platform_id platform;
} OpenCLEnvironment;

OpenCLEnvironment * get_ocl_env() {
  cl_int status;
  cl_uint num_platforms = 0;
  cl_platform_id *platforms = NULL;

  status = clGetPlatformIDs(0, NULL, &num_platforms);
  platforms = (cl_platform_id*)malloc(sizeof(cl_platform_id) * num_platforms);
  if(platforms == NULL) {
    fprintf(stdout, "Malloc failed for platforms...\n");
    exit(1);
  }
  status = clGetPlatformIDs(num_platforms, platforms, NULL);

  cl_uint num_devices = 0;
  cl_device_id *devices = NULL;

  status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);
  devices = (cl_device_id*)malloc(sizeof(cl_device_id) * num_devices);
  if(devices == NULL) {
    fprintf(stdout, "Malloc failed for devices...\n");
    exit(1);
  }
  status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL);

  cl_context context = CHECK_WITH_ERR_ARG(clCreateContext, status, NULL, num_devices, devices, NULL, NULL);

  cl_command_queue command_queue = CHECK_WITH_ERR_ARG(clCreateCommandQueue, status, context, devices[0], 0);

  OpenCLEnvironment *ocl_env = (OpenCLEnvironment *)malloc(sizeof(OpenCLEnvironment));

  if(ocl_env == NULL) {
    fprintf(stdout, "Malloc failed for OCL env...\n");
    exit(1);
  }

  ocl_env->context = context;
  ocl_env->queue = command_queue;
  ocl_env->device = devices[0];
  ocl_env->platform = platforms[0];

  // Free the devices list here - no need to keep it since we copied over device
  free(devices);
  // Likewise for platforms
  free(platforms);

  return ocl_env;
}

# matmul.c
// System includes
#include <stdbool.h>
#include <stdio.h>
#include <stdlib.h>

// Single header for OpenCL and associated helpers
#include "ocl_helpers.h"

const char* program_source = 
"__kernel                                           \n"
"void vecmad(__global float *A,                     \n"
"            __global float *B,                     \n"
"            __global float *C,                     \n"
"            int M, int N, int K) {                 \n"
" int idx = get_global_id(0);                       \n"
" int col = idx % M;                                \n"
" int row = (idx - col);                            \n"
" C[row + col] = 0.0f;                              \n"
" for(int k = 0; k < K; ++k) {                      \n"
"   C[row + col] += A[row + k] * B[k*N + col];      \n"
" }                                                 \n"
"}                                                  \n";

int main(int argc, char **args) {
  // Initialise host-side memory
  int N = 16;
  int n_elements = N*N;
  int data_size = n_elements * sizeof(float);
  printf("%d, %d\n", n_elements, data_size);
  float *A = get_ptr(n_elements), *B = get_ptr(n_elements), *C = get_ptr(n_elements);

  for(int i = 0; i < n_elements; ++i) {
    float val = (float)(i % N);
    A[i] = val;
    B[i] = val;
    C[i] = 0.0f;
  }

  print_sq_matrix(A, N);
  print_sq_matrix(B, N);

  printf("Host-memory initialised...\n");

  cl_int status;
  OpenCLEnvironment *env = get_ocl_env();

  // Get max sizes                                                              // Offending block
  size_t store[3] = { 0 };                                                      // Offending block
  CHECK_WITH_ERR_RES(clGetDeviceInfo, status, env->device,                      // Offending block
                     CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(store), store, NULL);// Offending block

  cl_mem buf_A = CHECK_WITH_ERR_ARG(clCreateBuffer, status,
                                    env->context, CL_MEM_READ_ONLY,
                                    data_size, NULL);
  cl_mem buf_B = CHECK_WITH_ERR_ARG(clCreateBuffer, status,
                                    env->context, CL_MEM_READ_ONLY,
                                    data_size, NULL);
  cl_mem buf_C = CHECK_WITH_ERR_ARG(clCreateBuffer, status,
                                    env->context, CL_MEM_WRITE_ONLY,
                                    data_size, NULL);

  CHECK_WITH_ERR_RES(clEnqueueWriteBuffer, status, env->queue,
                     buf_A, CL_FALSE, 0, data_size, A, 0, NULL, NULL);
  CHECK_WITH_ERR_RES(clEnqueueWriteBuffer, status, env->queue,
                     buf_B, CL_FALSE, 0, data_size, B, 0, NULL, NULL);

  printf("OpenCL machinery and buffers initialised, buffers written to...\n");

  cl_program program = CHECK_WITH_ERR_ARG(clCreateProgramWithSource, status,
                                          env->context, 1,
                                          (const char**)&program_source, NULL);
  CHECK_WITH_ERR_RES(clBuildProgram, status, program, 1, &(env->device),
                     NULL, NULL, NULL);

  printf("Program built...\n");

  cl_kernel kernel = CHECK_WITH_ERR_ARG(clCreateKernel, status, program,
                                        "vecmad");
  set_kernel_args(&kernel, buf_A, buf_B, buf_C, N, N, N);


  size_t globalws[1] = {256};

  status = clEnqueueNDRangeKernel(env->queue, kernel, 2, NULL, globalws, NULL, 0, NULL, NULL);
  if(status != CL_SUCCESS) {
    printf("%d\n", status);
    exit(1);
  }

  printf("Kernel queued...\n");

  // Wait on the kernel to finish
  CHECK_WITH_ERR_RES(clFinish, status, env->queue);

  // Read buffer and do some stuff
  ...
}

这段代码的问题是，如果我注释掉上面所示的整个“有问题的” block ，内核就会挂起，输出如下:

$ ./matmul
...
Host-memory initialised...
OpenCL machinery and buffers initialised, buffers written to...
Program built...
Kernel args set...
Kernel queued...
# hangs here

真正奇怪的部分是，如果我只注释掉对 clGetDeviceInfo 的调用，则会收到 -63 (CL_INVALID_GLOBAL_WORK_SIZE) 错误。我在 get_ocl_env() 中的设备和平台有问题吗？

Answer 1

NumPy 。在对 clEnqueueNDRangeKernel 的调用中，我将 2 列为工作项中的维度数，仅指定了一个维度:

  status = clEnqueueNDRangeKernel(env->queue, kernel, 2,                // <-- should be 1
                                  NULL, globalws, NULL, 0, NULL, NULL);

看这里: https://www.khronos.org/registry/OpenCL/sdk/1.2/docs/man/xhtml/clEnqueueNDRangeKernel.html