Why does clGetPlatformInfo get called in every clEnqueue function?

user_7082579 0
Viewed 139

We are profiling an OpenCL application running on an NVidia GPU on both the host and the device. We were surprised to find that (based on gperftools) the host was spending 44% of its time in clGetPlatformInfo, a method which is only called a single time in our own code. It is called by clEnqueueCopyBuffer_hid, clEnqueueWriteBuffer_hid, and clEnqueueNDRangeKernel_hid (and presumably all the other clEnqueue methods, but they are less commonly called in our code). Since this is taking so much of our host time, and we appear to be bound by the host speed right now, I need to know if there's a way to eliminate these extra calls.

Why is this being called by every OpenCL call? (Presumably it's static information that could be stored in the context?) Did we perhaps initialize our context incorrectly?

EDIT: I was asked for an MWE:

#include <CL/opencl.h>

#include <vector>
using namespace std;


int main ()
{
    cl_uint numPlatforms;
    clGetPlatformIDs (0, nullptr, &numPlatforms);

    vector<cl_platform_id> platformIdArray (numPlatforms);
    clGetPlatformIDs (numPlatforms, platformIdArray.data (), nullptr);

    // Assume the NVidia GPU is the first platform
    cl_platform_id platformId = platformIdArray[0];

    cl_uint numDevices;
    clGetDeviceIDs (platformId, CL_DEVICE_TYPE_GPU, 0, nullptr, &numDevices);

    vector<cl_device_id> deviceArray (numDevices);
    clGetDeviceIDs (platformId, CL_DEVICE_TYPE_GPU, numDevices, deviceArray.data (), nullptr);

    // Assume the NVidia GPU is the first device
    cl_device_id deviceId = deviceArray[0];

    cl_context context = clCreateContext (
        nullptr,
        1,
        &deviceId,
        nullptr,
        nullptr,
        nullptr);

    cl_command_queue commandQueue = clCreateCommandQueue (context, deviceId, {}, nullptr);

    cl_mem mem = clCreateBuffer (context, CL_MEM_READ_WRITE, sizeof(cl_int),
                                 nullptr, nullptr);

    cl_int i = 0;

    while (true)
    {
        clEnqueueWriteBuffer (
            commandQueue,
            mem,
            CL_TRUE,
            0,
            sizeof (i),
            &i,
            0,
            nullptr,
            nullptr);

        ++i;
    }
}

This MWE generates the following profile over the course of several seconds. Note that 99% of the time is spent in clGetPlatformInfo. MWE Profile Results

opencl profiling performance
2 Answers

Try to pass NULL as first parameter to clCreateContext. Device id is already being passed so that first parameter isn't probably required and may cause these extra calls to clGetPlatformInfo.

Another thing to try would be to link with non-Nvidia OpenCL library. It's not required to use GPU vendor OpenCL library, any should work as long as the features you are using are implemented in this other OpenCL library. With Nvidia there is no risk because as for today the latest supported version is OpenCL 1.2 which most if not all vendors already support. So you can try OpenCL lib from other vendors SDK's like Intel or AMD. If you are on Ubuntu you can use ocl-icd-opencl-dev.

====== UPDATE =========

Try to specify platform when creating the context:

const cl_context_properties properties[] = { CL_CONTEXT_PLATFORM, (cl_context_properties) platformId, 0}; 

cl_context context = clCreateContext (
    properties, // <-- here
    1,
    &deviceId,
    nullptr,
    nullptr,
    nullptr);

It may be that when the platform isn't specified when creating the context then it is queried each time it's needed but when is specified as above it won't be.

user_48993300

We figured out the problem: gproftools was having a difficult time giving the correct backtrace. The code was not actually calling clGetPlatformInfo thousands of times like gperftools said it was. According to a conversation with bashbaug on the Khronos forums:

When I run a test with gperftools using our GPU driver I see most of the time attributed to GTPin_Init, as you mentioned. I think this is because an OpenCL ICD has to export very few symbols, since calls into most OpenCL APIs occur through the ICD dispatch table.

We used his suggested profiling tools (the OpenCL Intercept Layer, found at https://github.com/intel/opencl-intercept-layer) to give us a better understanding of the runtime characteristics of the kernels and helped us find some memory leaks. It was the memory leaks that were actually causing the slowdown --- it seems as though kernels take a long time to start up if they are passed memory with high reference counts as arguments.

You can find the full conversation on the Khronos forums here: https://community.khronos.org/t/why-does-clgetplatforminfo-get-called-in-every-clenqueue-function/105756/5

user_70825790
Related