-
Notifications
You must be signed in to change notification settings - Fork 205
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
support sparse opnorm #1466
support sparse opnorm #1466
Conversation
A side question: maybe GPUArrays should have some kind of sparse abstraction so that this code can be generic? |
I realize somehow this kernel doesn't work on large sparse array, and I got the following n = 15
A = sprand(1<<n, 1<<n, 0.01)
dA = cu(A) Any idea why it errors in this way? julia> sum_dim1(CuSparseMatrixCSR(dA))
ERROR: CUDA error: invalid argument (code 1, ERROR_INVALID_VALUE)
Stacktrace:
[1] throw_api_error(res::CUDA.cudaError_enum)
@ CUDA ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/error.jl:91
[2] macro expansion
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/error.jl:101 [inlined]
[3] cuLaunchKernel(f::CuFunction, gridDimX::UInt32, gridDimY::UInt32, gridDimZ::UInt32, blockDimX::UInt32, blockDimY::UInt32, blockDimZ::UInt32, sharedMemBytes::Int64, hStream::CuStream, kernelParams::Vector{Ptr{Nothing}}, extra::Ptr{Nothing})
@ CUDA ~/.julia/packages/CUDA/5jdFl/lib/utils/call.jl:26
[4] #39
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:69 [inlined]
[5] macro expansion
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:33 [inlined]
[6] macro expansion
@ ./none:0 [inlined]
[7] pack_arguments(::CUDA.var"#39#40"{Bool, Int64, CuStream, CuFunction, CuDim3, CuDim3}, ::CUDA.KernelState, ::CuDeviceVector{Float32, 1}, ::CuSparseDeviceMatrixCSR{Float32, Int32, 1})
@ CUDA ./none:0
[8] #launch#38
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:62 [inlined]
[9] #44
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:136 [inlined]
[10] macro expansion
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:95 [inlined]
[11] macro expansion
@ ./none:0 [inlined]
[12] convert_arguments
@ ./none:0 [inlined]
[13] #cudacall#43
@ ~/.julia/packages/CUDA/5jdFl/lib/cudadrv/execution.jl:135 [inlined]
[14] macro expansion
@ ~/.julia/packages/CUDA/5jdFl/src/compiler/execution.jl:204 [inlined]
[15] macro expansion
@ ./none:0 [inlined]
[16] #call#241
@ ./none:0 [inlined]
[17] #_#262
@ ~/.julia/packages/CUDA/5jdFl/src/compiler/execution.jl:462 [inlined]
[18] macro expansion
@ ~/.julia/packages/CUDA/5jdFl/src/compiler/execution.jl:104 [inlined]
[19] sum_dim1(A::CuSparseMatrixCSR{Float32, Int32})
@ EaRydCUDA ~/EaRyd/lib/EaRydCUDA/src/opnorm.jl:17
[20] top-level scope
@ ~/EaRyd/lib/EaRydCUDA/benchmark/expmv.jl:29 |
Ah nvm, I should div the threads |
Codecov Report
@@ Coverage Diff @@
## master #1466 +/- ##
==========================================
- Coverage 77.31% 76.57% -0.74%
==========================================
Files 120 121 +1
Lines 9265 9290 +25
==========================================
- Hits 7163 7114 -49
- Misses 2102 2176 +74
Continue to review full report at Codecov.
|
Thanks! |
I only supported Inf opnorm, which is relatively simple, and is used in Krylov subspace methods. I guess it's fine to just error on other norms so if anyone hit it they could submit PR for it.