OpenCL使用子緩衝區的父項作爲常量參數

Question

我有一個OpenCL（1.2）內核，其參數爲constant，它是一個子緩衝區。當我運行這個內核時，似乎使用了父緩衝區。如果我使用global const參數，它按預期工作。

我會把它放在一個驅動程序錯誤，除了我可以在不同的機器上的英特爾（Linux，beignet git）和nVidia（Linux，367.44-3）實現上重現它，這讓我覺得我已經做出了某處出現錯誤。

下面是一個工作示例。預期的輸出是1, 1025, 1, 1025,，但打印1, 1, 1, 1025,。

#include <CL/cl.h> 
#include <stdio.h> 
#include <assert.h> 
#include <string.h> 

#define NELEMS(x) (sizeof(x)/sizeof(*x)) 

#define PLATFORM 0 
#define DEVICE 0 

const char src[] = 
    "kernel void test1(constant int * const a) {\n" 
    " size_t i = get_global_id(0);\n" 
    " if (i == 1)\n" 
    "  printf(\"%i, \", a[i]);\n" 
    "}\n" 
    "\n" 
    "kernel void test2(global const int * const a) {\n" 
    " size_t i = get_global_id(0);\n" 
    " if (i == 1)\n" 
    "  printf(\"%i, \", a[i]);\n" 
    "}\n"; 
const size_t src_len = sizeof(src); 
const char * const kernels[] = {"test1", "test2"}; 

int main(void) { 
    cl_int err = -1; 

    cl_uint num_platforms; 
    clGetPlatformIDs(0, NULL, &num_platforms); 
    assert(num_platforms > PLATFORM); 

    cl_platform_id * platforms = malloc(sizeof(*platforms) * num_platforms); 
    clGetPlatformIDs(num_platforms, platforms, NULL); 

    cl_uint num_devices; 
    clGetDeviceIDs(platforms[PLATFORM], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices); 
    assert(num_devices >= DEVICE); 

    cl_device_id * devices = malloc(sizeof(*devices) * num_devices); 
    clGetDeviceIDs(platforms[PLATFORM], CL_DEVICE_TYPE_ALL, num_devices, devices, NULL); 

    cl_context_properties context_properties[] = { 
     CL_CONTEXT_PLATFORM, (cl_context_properties) platforms[PLATFORM], 0 
    }; 

    cl_context context = clCreateContext(context_properties, 1, &devices[DEVICE], NULL, NULL, &err); 
    assert(err == CL_SUCCESS); 

#pragma GCC diagnostic push 
#pragma GCC diagnostic ignored "-Wdeprecated-declarations" 
    cl_command_queue queue = clCreateCommandQueue(context, devices[DEVICE], 0, &err); 
#pragma GCC diagnostic pop 
    assert(err == CL_SUCCESS); 

    cl_program program; 
    { 
     // Crashes if directly using src[] 
     char * source = malloc(src_len); 
     memcpy(source, src, src_len); 
     program = clCreateProgramWithSource(context, 1, (const char **) &source, &src_len, &err); 
     assert(err == CL_SUCCESS); 
     free(source); 
    } 

    err = clBuildProgram(program, 1, &devices[DEVICE], "", NULL, NULL); 
    assert(err == CL_SUCCESS); 

    size_t buffer_size = 8192; 
    size_t subbuffer_size = buffer_size/2; 
    { 
     cl_uint align; 
     err = clGetDeviceInfo(devices[DEVICE], CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(align), &align, NULL); 
     assert(err == CL_SUCCESS); 
     assert(subbuffer_size % align == 0); 

     cl_ulong constbuf_size; 
     err = clGetDeviceInfo(devices[DEVICE], CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(constbuf_size), &constbuf_size, NULL); 
     assert(err == CL_SUCCESS); 
     assert(constbuf_size > subbuffer_size); 
    } 
    cl_mem buffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_HOST_WRITE_ONLY, buffer_size, NULL, &err); 
    assert(err == CL_SUCCESS); 

    cl_mem sub_buffers[2]; 
    for (size_t i = 0; i < NELEMS(sub_buffers); i++){ 
     cl_buffer_region region = { 
      .origin = i * subbuffer_size, 
      .size = subbuffer_size, 
     }; 
     sub_buffers[i] = clCreateSubBuffer(buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err); 
     assert(err == CL_SUCCESS); 
    } 

    { 
     cl_int * data = clEnqueueMapBuffer(queue, buffer, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0, buffer_size, 0, NULL, NULL, &err); 
     assert(err == CL_SUCCESS); 
     for (size_t i = 0; i < buffer_size/sizeof(cl_int); i++) 
      data[i] = i; 
     cl_event unmap_event; 
     err = clEnqueueUnmapMemObject(queue, buffer, data, 0, NULL, &unmap_event); 
     assert(err == CL_SUCCESS); 
     err = clWaitForEvents(1, &unmap_event); 
     assert(err == CL_SUCCESS); 
    } 

    for (size_t k = 0; k < NELEMS(kernels); k++) { 
     cl_kernel kernel = clCreateKernel(program, kernels[k], &err); 
     assert(err == CL_SUCCESS); 

     cl_event run_event; 
     for (size_t i = 0; i < NELEMS(sub_buffers); i++){ 
      err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &sub_buffers[i]); 
      assert(err == CL_SUCCESS); 
      size_t work_size[] = {subbuffer_size/sizeof(cl_int)}; 
      err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, work_size, NULL, 0, NULL, &run_event); 
      assert(err == CL_SUCCESS); 
      err = clWaitForEvents(1, &run_event); 
      assert(err == CL_SUCCESS); 
      err = clFinish(queue); 
      assert(err == CL_SUCCESS); 
     } 

     clReleaseKernel(kernel); 
    } 
    puts(""); 

    for (size_t i = 0; i < NELEMS(sub_buffers); i++) 
     clReleaseMemObject(sub_buffers[i]); 
    clReleaseMemObject(buffer); 
    clReleaseProgram(program); 
    clReleaseCommandQueue(queue); 
    clReleaseContext(context); 

    free(devices); 
    free(platforms); 

    return 0; 
} 

Answer 1

這很有趣。我在差異設備上試用它，在MacBookPro中有3個設備包括Nvidia IRIS和Intel，都可以獲得正確的輸出。在使用Nvidia驅動程序的MBP中的窗口10中，輸出完全相同的錯誤。 我認爲這是一個Nvidia錯誤，但不限於Nvidia。