From e6c5f119c21f2ccd08aaa6a1b38cbdd505d0359d Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Sat, 17 Jul 2021 10:15:55 +0530 Subject: [PATCH 1/7] add CUDA overloads for depthwise_conv* --- src/cudnn/conv.jl | 35 +++++++++++++++++++++++++---------- 1 file changed, 25 insertions(+), 10 deletions(-) diff --git a/src/cudnn/conv.jl b/src/cudnn/conv.jl index 027d594..427b50c 100644 --- a/src/cudnn/conv.jl +++ b/src/cudnn/conv.jl @@ -1,6 +1,6 @@ - -using NNlib: DenseConvDims +using NNlib: DenseConvDims, DepthwiseConvDims import NNlib: conv!, ∇conv_filter!, ∇conv_data!, conv_bias_act! +import NNlib: depthwise_conv!, ∇depthwise_conv_filter!, ∇depthwise_conv_data! using CUDA.CUDNN: scalingParameter, CUDNN_CONVOLUTION, convdims, cudnnConvolutionDescriptor, cudnnConvolutionBwdDataAlgoPerf, @@ -10,8 +10,8 @@ using CUDA.CUDNN: scalingParameter, CUDNN_CONVOLUTION, convdims, const CUDNNFloat = Union{Float16,Float32,Float64} -function cudnnConvolutionDescriptor(cdims::DenseConvDims, x::DenseCuArray{T}) where T - mode=(NNlib.flipkernel(cdims) ? CUDNN_CROSS_CORRELATION : CUDNN_CONVOLUTION) +function cudnnConvolutionDescriptor(cdims::ConvDims, x::DenseCuArray{T}) where T + mode = (NNlib.flipkernel(cdims) ? CUDNN_CROSS_CORRELATION : CUDNN_CONVOLUTION) cudnnConvolutionDescriptor(convdims(nnlibPadding(cdims),size(x),0), convdims(NNlib.stride(cdims),size(x),1), convdims(NNlib.dilation(cdims),size(x),1), @@ -22,8 +22,8 @@ function cudnnConvolutionDescriptor(cdims::DenseConvDims, x::DenseCuArray{T}) wh Cint(1)) end -function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::DenseConvDims; - alpha=1, beta=0, algo=-1) where T<:CUDNNFloat +function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::ConvDims; + alpha = 1, beta = 0, algo = -1) where T<:CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -35,8 +35,8 @@ function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims end function conv_bias_act!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, - cdims::DenseConvDims, bias::DenseCuArray{T}, σ=identity; - z::DenseCuArray{T}=y, alpha=1, beta=0, algo=-1) where T<:CUDNNFloat + cdims::ConvDims, bias::DenseCuArray{T}, σ = identity; + z::DenseCuArray{T} = y, alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -54,7 +54,7 @@ function conv_bias_act!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{ end function ∇conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray{T}, - cdims::DenseConvDims; alpha=1, beta=0, algo=-1) where T<:CUDNNFloat + cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -72,7 +72,7 @@ function ∇conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray end function ∇conv_filter!(dw::DenseCuArray{T}, x::DenseCuArray{T}, dy::DenseCuArray{T}, - cdims::DenseConvDims; alpha=1, beta=0, algo=-1) where T<:CUDNNFloat + cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -95,3 +95,18 @@ function ∇conv_bias!(db::DenseCuArray{T}, dy::DenseCuArray{T}; alpha=1, beta=0 cudnnConvolutionBackwardBias(handle(), alpha, yDesc, dy, beta, bDesc, db) return db end + +function depthwise_conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::DepthwiseConvDims; + alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat + conv!(y, x, w, cims; alpha, beta, algo) +end + +function ∇depthwise_conv_filter!(dw::DenseCuArray{T}, x::DenseCuArray{T}, dy::DenseCuArray{T}, + cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat + ∇conv_filter!(dw, x, dy, cdims; alpha, beta, algo) +end + +function ∇depthwise_conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray{T}, + cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat + ∇conv_data!(dx, dy, w, cdims; alpha, beta, algo) +end From 28b90489d00fda9d37513e168ea02db2f979faca Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Tue, 17 Aug 2021 20:51:24 +0530 Subject: [PATCH 2/7] bump patch --- .buildkite/pipeline.yml | 4 ++++ Project.toml | 6 +++--- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index 5d446bd..1bcba73 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -4,6 +4,8 @@ steps: - JuliaCI/julia#v1: version: "1.6" - JuliaCI/julia-test#v1: ~ + env: + JULIA_PKG_SERVER: "" agents: queue: "juliagpu" cuda: "*" @@ -14,6 +16,8 @@ steps: - JuliaCI/julia#v1: version: "1.7-nightly" - JuliaCI/julia-test#v1: ~ + env: + JULIA_PKG_SERVER: "" agents: queue: "juliagpu" cuda: "*" diff --git a/Project.toml b/Project.toml index de5f246..233aaac 100644 --- a/Project.toml +++ b/Project.toml @@ -1,6 +1,6 @@ name = "NNlibCUDA" uuid = "a00861dc-f156-4864-bf3c-e6376f28a68d" -version = "0.1.7" +version = "0.1.8" [deps] CUDA = "052768ef-5323-5732-b1bb-66c8b64840ba" @@ -10,8 +10,8 @@ Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c" Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2" [compat] -CUDA = "3.3.1" -NNlib = "0.7.23" +CUDA = "3.3.1, 3.4.1" +NNlib = "0.7.25" julia = "1.6" [extras] From ed4c43618600e494747ea3e60d66c66a450b7603 Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Tue, 17 Aug 2021 20:55:33 +0530 Subject: [PATCH 3/7] revert an extra commit --- src/cudnn/conv.jl | 37 ++++++++----------------------------- 1 file changed, 8 insertions(+), 29 deletions(-) diff --git a/src/cudnn/conv.jl b/src/cudnn/conv.jl index 0240926..a4f23c1 100644 --- a/src/cudnn/conv.jl +++ b/src/cudnn/conv.jl @@ -1,6 +1,6 @@ -using NNlib: DenseConvDims, DepthwiseConvDims + +using NNlib: DenseConvDims import NNlib: conv!, ∇conv_filter!, ∇conv_data!, conv_bias_act! -import NNlib: depthwise_conv!, ∇depthwise_conv_filter!, ∇depthwise_conv_data! using CUDA.CUDNN: scalingParameter, CUDNN_CONVOLUTION, convdims, cudnnConvolutionDescriptor, cudnnConvolutionBwdDataAlgoPerf, @@ -10,8 +10,8 @@ using CUDA.CUDNN: scalingParameter, CUDNN_CONVOLUTION, convdims, const CUDNNFloat = Union{Float16,Float32,Float64} -function cudnnConvolutionDescriptor(cdims::ConvDims, x::DenseCuArray{T}) where T - mode = (NNlib.flipkernel(cdims) ? CUDNN_CROSS_CORRELATION : CUDNN_CONVOLUTION) +function cudnnConvolutionDescriptor(cdims::DenseConvDims, x::DenseCuArray{T}) where T + mode=(NNlib.flipkernel(cdims) ? CUDNN_CROSS_CORRELATION : CUDNN_CONVOLUTION) cudnnConvolutionDescriptor(convdims(nnlibPadding(cdims),size(x),0), convdims(NNlib.stride(cdims),size(x),1), convdims(NNlib.dilation(cdims),size(x),1), @@ -22,8 +22,8 @@ function cudnnConvolutionDescriptor(cdims::ConvDims, x::DenseCuArray{T}) where T Cint(NNlib.groupcount(cdims))) end -function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::ConvDims; - alpha = 1, beta = 0, algo = -1) where T<:CUDNNFloat +function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::DenseConvDims; + alpha=1, beta=0, algo=-1) where T<:CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -34,15 +34,9 @@ function conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims cudnnConvolutionForward!(y, w, x, d; alpha, beta, z=y) end -<<<<<<< HEAD -function conv_bias_act!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, - cdims::ConvDims, bias::DenseCuArray{T}, σ = identity; - z::DenseCuArray{T} = y, alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat -======= function conv_bias_act!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::DenseConvDims, bias::DenseCuArray{T}, σ=identity; z::DenseCuArray{T}=y, alpha=1, beta=0, algo=-1) where T<:CUDNNFloat ->>>>>>> 4092e7f68ddd00b60c9dc960be14dbd176e80ef8 if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -60,7 +54,7 @@ function conv_bias_act!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{ end function ∇conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray{T}, - cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat + cdims::DenseConvDims; alpha=1, beta=0, algo=-1) where T<:CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -78,7 +72,7 @@ function ∇conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray end function ∇conv_filter!(dw::DenseCuArray{T}, x::DenseCuArray{T}, dy::DenseCuArray{T}, - cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat + cdims::DenseConvDims; alpha=1, beta=0, algo=-1) where T<:CUDNNFloat if cudnnversion() < v"6" all(x -> x == 1, dilation(cdims)) || error("Only dilation = 1 is supported in cuDNN version < 6") end @@ -101,18 +95,3 @@ function ∇conv_bias!(db::DenseCuArray{T}, dy::DenseCuArray{T}; alpha=1, beta=0 cudnnConvolutionBackwardBias(handle(), alpha, yDesc, dy, beta, bDesc, db) return db end - -function depthwise_conv!(y::DenseCuArray{T}, x::DenseCuArray{T}, w::DenseCuArray{T}, cdims::DepthwiseConvDims; - alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat - conv!(y, x, w, cims; alpha, beta, algo) -end - -function ∇depthwise_conv_filter!(dw::DenseCuArray{T}, x::DenseCuArray{T}, dy::DenseCuArray{T}, - cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat - ∇conv_filter!(dw, x, dy, cdims; alpha, beta, algo) -end - -function ∇depthwise_conv_data!(dx::DenseCuArray{T}, dy::DenseCuArray{T}, w::DenseCuArray{T}, - cdims::ConvDims; alpha = 1, beta = 0, algo = -1) where T <: CUDNNFloat - ∇conv_data!(dx, dy, w, cdims; alpha, beta, algo) -end From 74348db8fdcd3acf5011534b6208a5548f8ee3e7 Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Tue, 17 Aug 2021 20:57:20 +0530 Subject: [PATCH 4/7] ditto with tests commit --- test/conv.jl | 38 ++++++++++++++++---------------------- 1 file changed, 16 insertions(+), 22 deletions(-) diff --git a/test/conv.jl b/test/conv.jl index e4abd30..f605519 100644 --- a/test/conv.jl +++ b/test/conv.jl @@ -9,21 +9,23 @@ using NNlib: DenseConvDims @test ∇conv_filter(a, c, cdims) ≈ collect(∇conv_filter(da, dc, cdims)) # Test for agreement between CPU NNlib and CuDNN versions, across a variety of kwargs - for num_spatial_dims in (1, 2, 3) + options = Dict{Any, Any}.(( + (), (:dilation => 2), (:flipkernel => true), (:stride => 2), + (:padding => 1), + )) + C_in_ = 3 + C_out = 4 + batch_size = 1 + + for groups in (1, 2, 4), num_spatial_dims in (1, 2, 3) + # Make `C_in = C_out` when using grouped convolution. + C_in = groups == 1 ? C_in_ : C_out # Initialize data we'll run our tests over - C_in = 3 - C_out = 4 - batch_size = 1 x = rand(Float64, fill(8, num_spatial_dims)..., C_in, batch_size) - w = rand(Float64, fill(2, num_spatial_dims)..., C_in, C_out) - b = rand(Float64, fill(1, num_spatial_dims)..., C_in, C_out) - options = (Dict(), Dict(:dilation => 2), Dict(:flipkernel => true), Dict(:stride => 2), Dict(:padding => 1)) - - # @denizyuret: algo option deprecated for nnlib, handling in cudnn - # algos = (1, 0, 1, 1,) - # for (opts, algo) in zip(options, algos) + w = rand(Float64, fill(2, num_spatial_dims)..., C_in ÷ groups, C_out) - for opts in options + for opts in options + opts[:groups] = groups cdims = DenseConvDims(x, w; opts...) y = NNlib.conv(x, w, cdims) @@ -36,19 +38,11 @@ using NNlib: DenseConvDims gputest((x, w) -> NNlib.conv(x, w, cdims; alpha=2.0), x, w, checkgrad=false) # TODO gputest((y, w) -> NNlib.∇conv_data(y, w, cdims; alpha=2.0), y, w, checkgrad=false) # TODO gputest((x, y) -> NNlib.∇conv_filter(x, y, cdims; alpha=2.0), x, y, checkgrad=false) # TODO - + gputest((y, x, w) -> NNlib.conv!(copy(y), x, w, cdims; beta=2.0), y, x, w, checkgrad=false) # TODO # @test_broken gputest((x, y, w) -> NNlib.∇conv_data!(copy(x), y, w, cdims; beta=2.0), x, y, w, checkgrad=false) #TODO gputest((w, x, y) -> NNlib.∇conv_filter!(copy(w), x, y, cdims; beta=2.0), w, x, y, checkgrad=false) # TODO end - - # CPU implementation of ∇conv_bias! - db = zeros(Float64, 1, 1, 3, 1) - dy = randn(Float64, 8, 8, 3, 1) - function NNlibCUDA.∇conv_bias!(db, dy) - db .= sum(dy, dims=1:(ndims(dy)-2)) - return db - end - gputest(NNlibCUDA.∇conv_bias!, db, dy, checkgrad=false) end + end From 8d8b2d8d4c20604ea2f3290a59d0ba5b37754528 Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Thu, 19 Aug 2021 23:55:20 +0530 Subject: [PATCH 5/7] force newest CUDA version --- Project.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Project.toml b/Project.toml index 233aaac..c9f096f 100644 --- a/Project.toml +++ b/Project.toml @@ -10,7 +10,7 @@ Random = "9a3f8284-a2c9-5f02-9a11-845980a1fd5c" Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2" [compat] -CUDA = "3.3.1, 3.4.1" +CUDA = "3.4.1" NNlib = "0.7.25" julia = "1.6" From 1d5f2c3a5aeb6faa7a72f49600593e2361f0070b Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Fri, 20 Aug 2021 00:04:41 +0530 Subject: [PATCH 6/7] force a second registry update --- .buildkite/pipeline.yml | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index 1bcba73..4cbcf34 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -1,8 +1,9 @@ steps: - label: "GPU integration with julia v1" + - command: julia --project -e "import Pkg; Pkg.Registry.update()" plugins: - JuliaCI/julia#v1: - version: "1.6" + version: "1.6" - JuliaCI/julia-test#v1: ~ env: JULIA_PKG_SERVER: "" From 611b05995a2e9b6865f07a6e461383d34511b38d Mon Sep 17 00:00:00 2001 From: Dhairya Gandhi Date: Fri, 20 Aug 2021 00:06:42 +0530 Subject: [PATCH 7/7] revert the forced update --- .buildkite/pipeline.yml | 1 - 1 file changed, 1 deletion(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index 4cbcf34..aa1eb1c 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -1,6 +1,5 @@ steps: - label: "GPU integration with julia v1" - - command: julia --project -e "import Pkg; Pkg.Registry.update()" plugins: - JuliaCI/julia#v1: version: "1.6"