     // Device code with the CUDA runtime.
 #if defined(CUB_INCLUDE_DEVICE_CODE) && defined(CUB_RUNTIME_ENABLED)
 
-    // Re-enable when the new API is host/device callable (NVBug 3386786)
-#if 0 && defined(__CUDACC__) &&                                                \
+#if defined(__CUDACC__) &&                                                     \
   ((__CUDACC_VER_MAJOR__ > 11) ||                                              \
    ((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 6)))
     // CUDA >= 11.6

  if (CUB_IS_HOST_CODE)
  {
#if CUB_INCLUDE_HOST_CODE

    // Re-enable when the new API is host/device callable (NVBug 3386786)
#if 0 && defined(__CUDACC__) &&                                                \
  ((__CUDACC_VER_MAJOR__ > 11) ||                                              \
   ((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 6)))
    // CUDA >= 11.6
    result = __cudaDeviceSynchronizeDeprecationAvoidance();
#else // CUDA < 11.6
    result = cudaDeviceSynchronize();
#endif

#endif
  }
  else

 * Call `cudaDeviceSynchronize()` using the proper API for the current CUB and
 * CUDA configuration.
 */
#pragma nv_exec_check_disable
CUB_RUNTIME_FUNCTION inline cudaError_t device_synchronize()
{
  cudaError_t result = cudaErrorUnknown;
