   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.
 */
CUB_RUNTIME_FUNCTION inline cudaError_t device_synchronize()
{
  cudaError_t result = cudaErrorUnknown;

    // Device code with the CUDA runtime.
#if defined(CUB_INCLUDE_DEVICE_CODE) && defined(CUB_RUNTIME_ENABLED)

#if defined(__CUDACC__) &&                                                     \
  ((__CUDACC_VER_MAJOR__ > 11) ||                                              \
   ((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 6)))
    // CUDA >= 11.6
