From 3bf556dedcbe37af9e24b02bc18ebcb4e5c0283c Mon Sep 17 00:00:00 2001 From: Joe Evans Date: Fri, 13 Nov 2020 22:45:56 -0800 Subject: [PATCH] Cuda 11 build fixes (#19530) * Don't use namespace for pow() function, since it is built into cuda math library, and cast the second argument so it will find an acceptable form. * Properly case exponent. * Remove thrust library override and use default from cuda 11.0. * Fix lint. * Fix lint. Co-authored-by: Joe Evans --- ci/build_windows.py | 17 +---------------- src/operator/contrib/multi_lamb.cu | 8 ++++---- src/operator/contrib/multi_lans.cu | 8 ++++---- 3 files changed, 9 insertions(+), 24 deletions(-) diff --git a/ci/build_windows.py b/ci/build_windows.py index f6626d6629af..f18492298cc9 100755 --- a/ci/build_windows.py +++ b/ci/build_windows.py @@ -151,20 +151,6 @@ def windows_build(args): mxnet_root = get_mxnet_root() logging.info("Found MXNet root: {}".format(mxnet_root)) - if 'GPU' in args.flavour: - # Get Thrust version to be shipped in Cuda 11, due to flakyness of - # older Thrust versions with MSVC 19 compiler - with remember_cwd(): - tmpdirname = tempfile.mkdtemp() - os.chdir(tmpdirname) - r = requests.get('https://github.com/thrust/thrust/archive/1.9.8.zip', allow_redirects=True) - with open('thrust.zip', 'wb') as f: - f.write(r.content) - with zipfile.ZipFile('thrust.zip', 'r') as zip_ref: - zip_ref.extractall('.') - thrust_path = os.path.join(tmpdirname, "thrust-1.9.8") - - # cuda thrust / CUB + VS 2019 is flaky: try multiple times if fail MAXIMUM_TRY = 5 build_try = 0 @@ -178,8 +164,7 @@ def windows_build(args): os.chdir(path) env = os.environ.copy() if 'GPU' in args.flavour: - env["CXXFLAGS"] = '/FS /MD /O2 /Ob2 /I {}'.format(thrust_path) - env["CUDAFLAGS"] = '-I {}'.format(thrust_path) + env["CXXFLAGS"] = '/FS /MD /O2 /Ob2' cmd = "\"{}\" && cmake -GNinja {} {}".format(args.vcvars, CMAKE_FLAGS[args.flavour], mxnet_root) diff --git a/src/operator/contrib/multi_lamb.cu b/src/operator/contrib/multi_lamb.cu index 6415bfbda015..0a55b89bc501 100644 --- a/src/operator/contrib/multi_lamb.cu +++ b/src/operator/contrib/multi_lamb.cu @@ -50,10 +50,10 @@ __global__ void KernelStep1(const MultiLAMBKernelParam kernel_pa MPDType biascorrection1, biascorrection2; if (bias_correction) { - biascorrection1 = 1.0 - - static_cast(std::pow(beta1, kernel_params.step_count[tensor_id])); - biascorrection2 = 1.0 - - static_cast(std::pow(beta2, kernel_params.step_count[tensor_id])); + biascorrection1 = 1.0 - static_cast( + pow(beta1, static_cast(kernel_params.step_count[tensor_id]))); + biascorrection2 = 1.0 - static_cast( + pow(beta2, static_cast(kernel_params.step_count[tensor_id]))); } else { biascorrection1 = static_cast(1.0); biascorrection2 = static_cast(1.0); diff --git a/src/operator/contrib/multi_lans.cu b/src/operator/contrib/multi_lans.cu index 64de72116514..2a7acb6bcaa9 100644 --- a/src/operator/contrib/multi_lans.cu +++ b/src/operator/contrib/multi_lans.cu @@ -53,10 +53,10 @@ __global__ void KernelStep1(const MultiLANSKernelParam kernel_pa MPDType biascorrection1, biascorrection2; - biascorrection1 = 1.0 - - static_cast(std::pow(beta1, kernel_params.step_count[tensor_id])); - biascorrection2 = 1.0 - - static_cast(std::pow(beta2, kernel_params.step_count[tensor_id])); + biascorrection1 = 1.0 - static_cast( + pow(beta1, static_cast(kernel_params.step_count[tensor_id]))); + biascorrection2 = 1.0 - static_cast( + pow(beta2, static_cast(kernel_params.step_count[tensor_id]))); MPDType r_weight[ILP_LAMB]; MPDType r_grad[ILP_LAMB];