OpenCL non-diverging if() followed by printf() causing massive slowdown and kernel code interfering on maximum workgroup size

user_16601442 0
Viewed 81

I am new to OpenCL and I have read in several places that if/else structures should be avoided, mainly because when the evaluation of threads is different (divergent branching) there is a significant slowdown.

Nonetheless, I used one if(cond) followed by some prints to guarantee that when a forbidden condition is met I can debug what caused it. The point is, whenever the if() is true for a single thread I will kill the process, therefore I am not worried about different threads evaluating the condition differently.

Yet, I found out that even when all threads evaluate false in this if(), there is a massive slowdown when compared to not using the if() -- I commented out the conditional and its body statements to verify.

An observation: I have a kernel (128 workgroups, each one with 128 workitems) that calls a function foo(), and the if/prints are inside foo(). The relevant part of foo() is as follows:

foo(){
     bool leftCorrect, rightCorrect, topCorrect, bottomCorrect, topLeftCorrect, topRightCorrect, bottomLeftCorrect, bottomRightCorrect;
     for(i=0;i<11;i++){
          for(j=0;j<11;j++){
               // Initial assignments. Using select() to avoid branch divergence
               leftCorrect = select(true,false,condition); 
               rightCorrect = select(true,false,condition); 
               topCorrect = select(true,false,condition); 
               bottomCorrect = select(true,false, condition);
            
               // Use the previous variables to compute the other bool variables and update them
               // Some boolean operations ...
           
               // This if basically tests if more than one is true
               if(leftCorrect + rightCorrect + topCorrect + bottomCorrect + topLeftCorrect + topRightCorrect + bottomLeftCorrect + bottomRightCorrect > 1){
                    printf("@@@@@\nFATAL ERROR: Multiple corrections in gid=%d\n",get_local_id(0));
                    printf("L  %d\n", leftCorrect);
                    printf("R  %d\n", rightCorrect);
                    printf("T  %d\n", topCorrect);
                    printf("B  %d\n", bottomCorrect);

                    printf("TL %d\n", topLeftCorrect);
                    printf("TR %d\n", topRightCorrect);
                    printf("BL %d\n", bottomLeftCorrect);
                    printf("BR %d\n", bottomRightCorrect);
              } 
              // Do something with the boolean variables and select() statements before return
         }
    }
}

Relevant information:

  1. With the code I showed, the entire program (main host + kernel + foo) takes around 90 seconds to run
  2. If I comment out the entire if() statement (condition + prints), the program takes around 2 (two) seconds to run
  3. If I remove all the prints and use a single dummy assignment (z=0, for instance) on the if(), the program takes around 2 seconds to run
  4. If I remove all the prints and use a single assignment involving get_local_id(0) (z=get_local_id(0), for instance) on the if(), the program takes around 2 seconds to run
  5. If I remove all the prints and use a single assignment involving get_local_id(0) and print it afterward, the program takes around 7 seconds to run
  6. By removing all printf() statements from the original code and adding one at a time (only the first print, then the 1st and 2nd, then from 1st to 3rd, ...), I found out that up to the fourth print the running time is around 7 seconds, but after adding the 5th print the running time jumps to 90 seconds

Another relevant information, discovered in the previous tests: I am printing the maximum workgroup size in the host code to know the differences in different platforms. In the cases with a small running time (cases 2-5, up to 7 seconds), the host outputs that the maximum WG size is 8192. However, in the cases of long running time, it outputs a maximum WG size of 4096.

The clinfo command shows me the following information:

Platform Name                                   Intel(R) CPU Runtime for OpenCL(TM) 
Device Name                                     Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz
Max work item sizes                             8192x8192x8192
Max work group size                             8192

And the host code to show the maximum WG size is:

size_t size_ret;
cl_uint maximum_size;
error = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, 
size_ret, &maximum_size, NULL);
cout << "-- Maximum WG size " << maximum_size << endl;

All this information is summarized in the two questions:

  1. Why is this tiny modification in the code causing such a massive slowdown, when the if() code is never executed?
  2. How can my code interfere in the maximum WG size, which is a property of the device?
opencl c
2 Answers

The way printf is implemented from OpenCL kernels must invariably include some kind of buffer for sending the written text to the host. But there must be coordination between work items so they don't write to the same part of the buffer. This means using local memory, atomics, etc. The compiler can't statically determine that your printf will never be called, so it must still allocate enough memory for all possible debug output to be returned. If the coordination involves barriers, those barriers will need to be inserted on the else path as well because all work items in a group must reach barriers in lockstep. As your maximum work group size changes depending on whether you have print statements or not, that's a good indication that your platform's printf implementation uses local memory to some extent, and that usually means it uses barriers. So the invisible barriers on the else path could be what's slowing things down so badly for you.

So by its nature, generic debug output from OpenCL kernels will always slow things down. You shouldn't have debug output in kernels you are profiling or deploying, because they'll wreck whatever optimisations you might be trying to implement.

If you're trying to debug or develop a kernel and are finding the necessary debug output to be slowing things down so much it's interfering with your work, you may want to get rid of printf entirely and use output structured diagnostics instead: a global buffer with space for all the different diagnostic values you want to return from your work items. So perhaps a struct with arrays for each returned value, one array entry per work item. Then convert that to a human-readable form (i.e. printf) on the host side. The advantages of this are:

  • printf itself is actually a fairly involved function, particularly for a GPU, so it itself will slow things down. Even if the text formatting itself is done on the host, with only the data arguments written to a shared buffer by the device, access to that buffer needs to be synchronised in some way, and unless the required size of the buffer can be and is statically proven by the compiler, resize or flush operations or similar will be required too.
  • Managing the diagnostic output memory buffer yourself means no running into output buffer size limits, reallocating, etc.
  • If no two work-items write to the same diagnostic value memory cell, you don't need any synchronisation or atomics.
  • Simply storing diagnostic output as arrays of floats, integers, etc. is more memory efficient, so this won't put so much pressure on GPU memory bandwidth.
user_486600

it seems that you got a problem with the debug. Iknow that using ARM mali you need to add in you CCP

        /* Enable a printf callback function for this context. */
        CL_PRINTF_CALLBACK_ARM,   (cl_context_properties) printf_callback,

        /* Request a minimum printf buffer size of 4MiB for devices in the
         context that support this extension. */
        CL_PRINTF_BUFFERSIZE_ARM, (cl_context_properties) 0x1000,

and in your kernel

        pragma OPENCL EXTENSION cl_arm_printf : enable

so check you GPU and try to find if you do not need some spécial extensions ;))

user_66367870
Related