Flux.jl: DepthwiseConv does not run on GPU

[terasakisatoshi]

I'm happy to hear Adds support for Depthwise Convolutions #279 is merged into master.

I updated the latest flux i.e. pkg> add Flux#master.

Here is a sample code I tested

using Flux:DepthwiseConv
num_ch=10
test_input = rand(224,224,num_ch,2) #WHCB style
depthwiseconv = DepthwiseConv((3,3),num_ch,stride=1,pad=1)

@show size(depthwiseconv(test_input)) # Works fine

This works fine, but once I copy depthwiseconv layer into GPU via gpu function

using CuArrays
depthwiseconv = depthwiseconv |> gpu
test_input = test_input |> gpu
@show size(depthwiseconv(test_input)) # Oops

the code above occurs error with the following message:

ERROR: LoadError: UndefVarError: gpu not defined
Stacktrace:
 [1] top-level scope at none:0
 [2] include_string(::Module, ::String, ::String) at ./loading.jl:1002
 [3] (::getfield(Atom, Symbol("##129#135")){String,String,Module})() at /home/terasaki/.julia/packages/Atom/7rQ1O/src/eval.jl:125
 [4] withpath(::getfield(Atom, Symbol("##129#135")){String,String,Module}, ::String) at /home/terasaki/.julia/packages/CodeTools/hB4Hy/src/utils.jl:30
 [5] withpath at /home/terasaki/.julia/packages/Atom/7rQ1O/src/eval.jl:46 [inlined]
 [6] #128 at /home/terasaki/.julia/packages/Atom/7rQ1O/src/eval.jl:122 [inlined]
 [7] with_logstate(::getfield(Atom, Symbol("##128#134")){String,String,Module}, ::Base.CoreLogging.LogState) at ./logging.jl:397
 [8] with_logger at ./logging.jl:493 [inlined]
 [9] #127 at /home/terasaki/.julia/packages/Atom/7rQ1O/src/eval.jl:121 [inlined]
 [10] hideprompt(::getfield(Atom, Symbol("##127#133")){String,String,Module}) at /home/terasaki/.julia/packages/Atom/7rQ1O/src/repl.jl:85
 [11] macro expansion at /home/terasaki/.julia/packages/Atom/7rQ1O/src/eval.jl:120 [inlined]
 [12] (::getfield(Atom, Symbol("##126#132")){Dict{String,Any}})() at ./task.jl:85
in expression starting at /home/terasaki/work/juliaExer/neuralNetwork/fluxExer/layers/depthwiseconv.jl:10

[MikeInnes]

Try using Flux: gpu. If you write using Flux: depthwiseconv then you'll only get that specific name.

You can also use cu instead of gpu for single arrays.

[terasakisatoshi]

I'm sorry I should have write using Flux.

The following code is what I wanted to say ......

using Flux:DepthwiseConv

num_ch=10
test_input = rand(224,224,num_ch,2) #WHCB style
depthwiseconv = DepthwiseConv((3,3),num_ch,stride=1,pad=1)

@show size(depthwiseconv(test_input)) # Works fine

using Flux
using CuArrays
depthwiseconv = depthwiseconv |> gpu
test_input = test_input |> gpu
@show size(depthwiseconv(test_input)) # Oops

EDITED this occurs the following error.

ERROR: LoadError: conversion to pointer not defined for CuArray{Float32,2}
Stacktrace:
 [1] error(::String) at ./error.jl:33
 [2] unsafe_convert(::Type{Ptr{Float32}}, ::CuArray{Float32,2}) at ./pointer.jl:67
 [3] pointer at ./abstractarray.jl:861 [inlined]
 [4] #depthwiseconv2d!#37(::Tuple{Int64,Int64}, ::Tuple{Int64,Int64}, ::Int64, ::Float32, ::Function, ::CuArray{Float32,4}, ::CuArray{Float32,4}, ::CuArray{Float32,4}) at /home/terasaki/.julia/packages/NNlib/0EAe7/src/impl/conv.jl:200
 [5] #depthwiseconv2d! at ./none:0 [inlined]
 [6] #depthwiseconv!#69 at /home/terasaki/.julia/packages/NNlib/0EAe7/src/conv.jl:100 [inlined]
 [7] #depthwiseconv! at ./none:0 [inlined]
 [8] #depthwiseconv#68(::Tuple{Int64,Int64}, ::Tuple{Int64,Int64}, ::Function, ::CuArray{Float32,4}, ::CuArray{Float32,4}) at /home/terasaki/.julia/packages/NNlib/0EAe7/src/conv.jl:96

[MikeInnes]

I guess we may not have an implementation in CuArrays yet; cc @avik-pal

[avik-pal]

Yes that is the case. Since the CPU code is by calling the im2col (thats were the conversion to pointer error is coming from) and corresponding functions, the implementation does not work on GPU. However, this can be fixed easily by using CUDNN (I recently came to know about this feature).
This function needs to be added and group count should be equal to the no of channels.
Also this cudnn function will allow us to support grouped convolutions. :)

[Sleort]

What's the status on this issue, @avik-pal ? My current research depends on DepthwiseConv, thus I'm currently limited to the CPU...

Unfortunately, I'm not really familiar with CUDNN/CuArrays.jl "under the hood", but if there's anything I'm able to assist with, I'm happy to help...

[dbadrian]

I second Sleort's comment. I'm highly interested in getting this available on the GPU. Just CPU performance is prohibitively slow except for the tiniest of problems.

[avik-pal]

@Sleort @dbadrian I am currently working on some other problems so won't be able to get to this immediately. I had linked the function which needs to be wrapped for this functionality.

You could look into the functions in libcudnn.jl to see how these are wrapped.

[dbadrian]

Admittedly, at this point, this is out of my depth to fully grasp what to do.

I gather you are referring to https://github.com/JuliaGPU/CuArrays.jl/blob/master/src/dnn/libcudnn.jl ?

Is the linked function the only thing you require to be wrapped? Where do you think this should go? Essentially libcudnn would be the appropriate place, right? Tell me where, Ill do the PR if so desired.

function cudnnSetConvolutionGroupCount(convDesc,groupCount)
    @check ccall((:cudnnSetConvolutionGroupCount,libcudnn),
                 cudnnStatus_t,
                 (cudnnConvolutionDescriptor_t,Cint),
                 convDesc,groupCount)
end

function cudnnGetConvolutionGroupCount(convDesc,groupCount)
    @check ccall((:cudnnGetConvolutionGroupCount,libcudnn),
                 cudnnStatus_t,
                 (cudnnConvolutionDescriptor_t,Ptr{Cint}),
                 convDesc,groupCount)
end

Would be great to see this supported in Flux, as its an extremely useful feature for various architectures.

Would then, as you said regarding group conv, also resolve https://github.com/FluxML/Flux.jl/issues/330.

[avik-pal]

Yes this is the only function to be wrapped. You should open the pr in CuArrays.jl and we can discuss there.

The next thing would be to modify the ConvDesc call here to call the wrapped function first.

[dbadrian]

Great, opened the PR as you probably been notified about. Happy to have a crack at other changes required as well (albeit with some help probably :)).