From 224cac24f15449a338abebe1bc17ba74c07d9f5c Mon Sep 17 00:00:00 2001 From: krasow Date: Tue, 7 Apr 2026 22:15:40 -0500 Subject: [PATCH 1/2] complex type patches and fixes. --- lib/cunumeric_jl_wrapper/include/types.h | 12 +- lib/cunumeric_jl_wrapper/src/wrapper.cpp | 6 +- src/cuNumeric.jl | 17 +-- src/ndarray/binary.jl | 30 +++- src/ndarray/detail/ndarray.jl | 51 +++---- src/ndarray/ndarray.jl | 42 +++--- src/ndarray/unary.jl | 49 ++++++- test/runtests.jl | 166 +++++++++++++---------- test/tests/binary_tests.jl | 8 ++ test/tests/gemm.jl | 34 ++--- test/tests/stability.jl | 80 +++++------ test/tests/unary_tests.jl | 34 +++-- test/tests/util.jl | 14 +- 13 files changed, 323 insertions(+), 220 deletions(-) diff --git a/lib/cunumeric_jl_wrapper/include/types.h b/lib/cunumeric_jl_wrapper/include/types.h index 600b6ec4..c75754db 100644 --- a/lib/cunumeric_jl_wrapper/include/types.h +++ b/lib/cunumeric_jl_wrapper/include/types.h @@ -42,19 +42,19 @@ DEFINE_CODE_TO_CXX(UINT8, uint8_t) DEFINE_CODE_TO_CXX(UINT16, uint16_t) DEFINE_CODE_TO_CXX(UINT32, uint32_t) DEFINE_CODE_TO_CXX(UINT64, uint64_t) -#ifdef HAVE_CUDA +#if LEGATE_DEFINED(LEGATE_USE_CUDA) DEFINE_CODE_TO_CXX(FLOAT16, __half) #else -// Dummy type for FLOAT16 when CUDA is not available -// This allows compilation but we throw an error if actually used -struct __half_dummy {}; -DEFINE_CODE_TO_CXX(FLOAT16, __half_dummy) +struct __dummy {}; +DEFINE_CODE_TO_CXX(FLOAT16, __dummy) #endif DEFINE_CODE_TO_CXX(FLOAT32, float) DEFINE_CODE_TO_CXX(FLOAT64, double) DEFINE_CODE_TO_CXX(COMPLEX64, std::complex) DEFINE_CODE_TO_CXX(COMPLEX128, std::complex) -#undef DEFINE_CODE_TO_CXX + +using HalfType = typename code_to_cxx::type; + } // namespace legate_util // Unary op codes diff --git a/lib/cunumeric_jl_wrapper/src/wrapper.cpp b/lib/cunumeric_jl_wrapper/src/wrapper.cpp index 21a5ca63..29334333 100644 --- a/lib/cunumeric_jl_wrapper/src/wrapper.cpp +++ b/lib/cunumeric_jl_wrapper/src/wrapper.cpp @@ -72,21 +72,23 @@ JLCXX_MODULE define_julia_module(jlcxx::Module& mod) { using jlcxx::ParameterList; using jlcxx::Parametric; using jlcxx::TypeVar; + using legate_util::HalfType; // Map C++ complex types to Julia complex types mod.map_type>("ComplexF64"); mod.map_type>("ComplexF32"); + mod.map_type("Float16"); // These are the types/dims used to generate templated functions // i.e. only these types/dims can be used from Julia side - using fp_types = ParameterList; + using fp_types = ParameterList; using int_types = ParameterList; using uint_types = ParameterList; using all_types = ParameterList, - std::complex>; + std::complex, HalfType>; using allowed_dims = ParameterList< std::integral_constant, std::integral_constant, std::integral_constant, std::integral_constant>; diff --git a/src/cuNumeric.jl b/src/cuNumeric.jl index 0e5c15d1..9e8fee73 100644 --- a/src/cuNumeric.jl +++ b/src/cuNumeric.jl @@ -54,22 +54,15 @@ end const DEFAULT_FLOAT = Float32 const DEFAULT_INT = Int32 -const SUPPORTED_INT_TYPES = Union{Int32,Int64} -const SUPPORTED_FLOAT_TYPES = Union{Float32,Float64} +const SUPPORTED_INT_TYPES = Union{Int8,Int16,Int32,Int64,UInt8,UInt16,UInt32,UInt64} +const SUPPORTED_FLOAT_TYPES = Union{Float32,Float64} # Float16 not supported yet const SUPPORTED_COMPLEX_TYPES = Union{ComplexF32,ComplexF64} + const SUPPORTED_NUMERIC_TYPES = Union{ SUPPORTED_INT_TYPES,SUPPORTED_FLOAT_TYPES,SUPPORTED_COMPLEX_TYPES } -# const SUPPORTED_TYPES = Union{SUPPORTED_INT_TYPES,SUPPORTED_FLOAT_TYPES,Bool} #* TODO Test UInt, Complex - -const SUPPORTED_TYPES = Union{ - Bool, - Int8,Int16,Int32,Int64, - UInt8,UInt16,UInt32,UInt64, - Float16,Float32,Float64, - ComplexF32,ComplexF64, - String, -} +const SUPPORTED_ARRAY_TYPES = Union{Bool,SUPPORTED_NUMERIC_TYPES} +const SUPPORTED_TYPES = Union{SUPPORTED_ARRAY_TYPES,String} # const MAX_DIM = 6 # idk what we compiled? diff --git a/src/ndarray/binary.jl b/src/ndarray/binary.jl index a8715240..d5590e23 100644 --- a/src/ndarray/binary.jl +++ b/src/ndarray/binary.jl @@ -154,7 +154,6 @@ LinearAlgebra.mul!(out, a, b) function LinearAlgebra.mul!( out::NDArray{T,2}, rhs1::NDArray{A,2}, rhs2::NDArray{B,2} ) where {T<:SUPPORTED_NUMERIC_TYPES,A,B} - #! This will probably need more checks once we support Complex number size(rhs1, 2) == size(rhs2, 1) || throw(DimensionMismatch("Matrix dimensions incompatible: $(size(rhs1)) × $(size(rhs2))")) (size(out, 1) == size(rhs1, 1) && size(out, 2) == size(rhs2, 2)) || throw( @@ -162,12 +161,29 @@ function LinearAlgebra.mul!( "mul! output is $(size(out)), but inputs would produce $(size(rhs1,1))×$(size(rhs2,2))" ), ) - T_OUT = __my_promote_type(A, B) - ((T_OUT <: AbstractFloat) && (T <: Integer)) && throw( - ArgumentError( - "mul! output has integer type $(T), but inputs promote to floating point type: $(T_OUT)" - ), - ) + + T_REQUIRED = __my_promote_type(A, B) + if promote_type(T_REQUIRED, T) != T + if (T_REQUIRED <: Complex && !(T <: Complex)) + throw( + ArgumentError( + "Implicit promotion: mul! output has real type $(T), but inputs promote to complex type: $(T_REQUIRED)" + ), + ) + elseif (T_REQUIRED <: AbstractFloat && T <: Integer) + throw( + ArgumentError( + "Implicit promotion: mul! output has integer type $(T), but inputs promote to floating point type: $(T_REQUIRED)" + ), + ) + end + # General case (e.g. Float64 result into Float32) + throw( + ArgumentError( + "mul! output type $(T) cannot hold the promoted input type $(T_REQUIRED). Implicit promotion to wider type or complex result is disallowed." + ), + ) + end return nda_three_dot_arg(checked_promote_arr(rhs1, T), checked_promote_arr(rhs2, T), out) end diff --git a/src/ndarray/detail/ndarray.jl b/src/ndarray/detail/ndarray.jl index 7186baa6..62ec585a 100644 --- a/src/ndarray/detail/ndarray.jl +++ b/src/ndarray/detail/ndarray.jl @@ -33,40 +33,43 @@ The NDArray type represents a multi-dimensional array in cuNumeric. It is a wrapper around a Legate array and provides various methods for array manipulation and operations. Finalizer calls `nda_destroy_array` to clean up the underlying Legate array when the NDArray is garbage collected. """ -mutable struct NDArray{T,N,PADDED} <: AbstractNDArray{T,N} +mutable struct NDArray{T,N,PADDED,P} <: AbstractNDArray{T,N} ptr::NDArray_t nbytes::Int64 padding::Union{Nothing,NTuple{N,Int}} + parent::P function NDArray(ptr::NDArray_t, ::Type{T}, ::Val{N}) where {T,N} nbytes = cuNumeric.nda_nbytes(ptr) cuNumeric.register_alloc!(nbytes) - handle = new{T,N,false}(ptr, nbytes, nothing) + handle = new{T,N,false,Nothing}(ptr, nbytes, nothing, nothing) finalizer(handle) do h cuNumeric.nda_destroy_array(h.ptr) cuNumeric.register_free!(h.nbytes) end return handle end -end - -# Dynamic fallback, not great but required if we cannot infer things -NDArray(ptr::NDArray_t; T=get_julia_type(ptr), N::Integer=get_n_dim(ptr)) = NDArray(ptr, T, Val(N)) -# struct WrappedNDArray{T,N} <: AbstractNDArray{T,N} -# ndarr::NDArray{T,N} -# jlarr::Array{T,N} - -# function WrappedNDArray(ndarray::NDArray{T,N}, jlarray::Array{T,N}) where {T,N} -# ndarr = ndarray -# jlarr = jlarray -# end + # Explicit parent inner constructor + function NDArray(ptr::NDArray_t, ::Type{T}, ::Val{N}, parent::P) where {T,N,P} + nbytes = cuNumeric.nda_nbytes(ptr) + cuNumeric.register_alloc!(nbytes) + handle = new{T,N,false,P}(ptr, nbytes, nothing, parent) + finalizer(handle) do h + cuNumeric.nda_destroy_array(h.ptr) + cuNumeric.register_free!(h.nbytes) + end + return handle + end +end +# this here is to avoid if else patterns +@inline _NDArray(ptr, T, v, ::Nothing) = NDArray(ptr, T, v) +@inline _NDArray(ptr, T, v, parent) = NDArray(ptr, T, v, parent) -# function WrappedNDArray(ndarray::NDArray{T,N}) where {T,N} -# ndarr = ndarray -# jlarr = nothing -# end -# end +# Dynamic fallback +function NDArray(ptr::NDArray_t; T=get_julia_type(ptr), N::Integer=get_n_dim(ptr), parent=nothing) + return _NDArray(ptr, T, Val(N), parent) +end #! JUST USE FULL TO MAKE a 0D? # $ cuNumeric.nda_full_array(UInt64[], 2.0f0) @@ -314,7 +317,7 @@ function nda_attach_external(arr::AbstractArray{T,N}) where {T,N} # Use the CxxWrap method for type-safe interaction # This returns a raw pointer compatible with the NDArray constructor nda_ptr = cuNumeric.nda_store_to_ndarray(st.handle) - return NDArray(nda_ptr; T=T, N=N) + return NDArray(nda_ptr, T, Val(N), arr) end # return underlying logical store to the NDArray obj @@ -448,8 +451,8 @@ Emits warnings when array sizes or element types differ. - Iterates over elements using `CartesianIndices` to compare element-wise difference. """ function compare( - julia_array::AbstractArray{T,N}, arr::NDArray{T,N}, atol::Real, rtol::Real -) where {T,N} + julia_array::AbstractArray{T1,N}, arr::NDArray{T2,N}, atol::Real, rtol::Real +) where {T1,T2,N} if (shape(arr) != Base.size(julia_array)) @warn "NDArray has shape $(shape(arr)) and Julia array has shape $(Base.size(julia_array))!\n" return false @@ -468,8 +471,8 @@ function compare( end function compare( - arr::NDArray{T,N}, julia_array::AbstractArray{T,N}, atol::Real, rtol::Real -) where {T,N} + arr::NDArray{T2,N}, julia_array::AbstractArray{T1,N}, atol::Real, rtol::Real +) where {T1,T2,N} return compare(julia_array, arr, atol, rtol) end diff --git a/src/ndarray/ndarray.jl b/src/ndarray/ndarray.jl index 21bfbd31..adae3ebf 100644 --- a/src/ndarray/ndarray.jl +++ b/src/ndarray/ndarray.jl @@ -149,20 +149,17 @@ function (::Type{<:Array})(arr::NDArray{B}) where {B} end # conversion from Base Julia array to NDArray -function (::Type{<:NDArray{A}})(arr::Array{B}) where {A,B} - dims = Base.size(arr) - out = cuNumeric.zeros(A, dims) - attached = cuNumeric.nda_attach_external(arr) - copyto!(out, attached) # copy elems of attached to resulting out - return out +function (::Type{<:NDArray{T}})(arr::Array{T,N}) where {T,N} + return cuNumeric.nda_attach_external(arr) end -function (::Type{<:NDArray})(arr::Array{B}) where {B} - dims = Base.size(arr) - out = cuNumeric.zeros(B, dims) - attached = cuNumeric.nda_attach_external(arr) - copyto!(out, attached) - return out +function (::Type{<:NDArray{A}})(arr::Array{B,N}) where {A,B,N} + # If types differ, we cast in Julia first (creating a temp) then attach + return cuNumeric.nda_attach_external(A.(arr)) +end + +function (::Type{<:NDArray})(arr::Array{T,N}) where {T,N} + return cuNumeric.nda_attach_external(arr) end # Base.convert(::Type{<:NDArray{T}}, a::A) where {T, A} = NDArray(T(a))::NDArray{T} @@ -215,7 +212,7 @@ size(arr) size(arr, 2) ``` """ -Base.size(arr::NDArray{<:Any, N}) where N = cuNumeric.shape(arr) +Base.size(arr::NDArray{<:Any,N}) where {N} = cuNumeric.shape(arr) Base.size(arr::NDArray, dim::Int) = Base.size(arr)[dim] @doc""" @@ -339,6 +336,16 @@ function Base.setindex!(arr::NDArray{T,N}, value::T, idxs::Vararg{Int,N}) where _setindex!(Val{N}(), arr, value, idxs...) end +function Base.setindex!(arr::NDArray{Complex{T},N}, value::T, idxs::Vararg{Int,N}) where {T,N} + assertscalar("setindex!") + _setindex!(Val{N}(), arr, Complex{T}(value), idxs...) +end + +function Base.setindex!(arr::NDArray{T,N}, value, idxs::Vararg{Int,N}) where {T,N} + assertscalar("setindex!") + _setindex!(Val{N}(), arr, convert(T, value), idxs...) +end + function _setindex!(::Val{0}, arr::NDArray{T,0}, value::T) where {T<:SUPPORTED_NUMERIC_TYPES} acc = NDArrayAccessor{T,1}() write(acc, arr.ptr, StdVector(UInt64[0]), value) @@ -511,7 +518,6 @@ falses(dims::Dims) = cuNumeric.fill(false, dims) falses(dims::Int...) = cuNumeric.fill(false, dims) falses(dim::Int) = cuNumeric.fill(false, dim) - @doc""" cuNumeric.zeros([T=Float32,] dims::Int...) cuNumeric.zeros([T=Float32,] dims::Tuple) @@ -526,7 +532,7 @@ cuNumeric.zeros(Float64, 3) cuNumeric.zeros(Int32, (2,3)) ``` """ -function zeros(::Type{T}, dims::Dims{N}) where {T<:SUPPORTED_TYPES, N} +function zeros(::Type{T}, dims::Dims{N}) where {T<:SUPPORTED_TYPES,N} return nda_zeros_array(dims, T) end @@ -614,7 +620,9 @@ A = cuNumeric.zeros(2, 2); cuNumeric.rand!(A) ``` """ Random.rand!(arr::NDArray{Float64}) = cuNumeric.nda_random(arr, 0) -Random.rand!(arr::NDArray{T}) where T = error("rand! only supports NDArray{Float64} for now. Cast with cuNumeric.as_type.") +function Random.rand!(arr::NDArray{T}) where {T} + error("rand! only supports NDArray{Float64} for now. Cast with cuNumeric.as_type.") +end function rand(::Type{T}, dims::Dims) where {T<:AbstractFloat} arrfp64 = cuNumeric.nda_random_array(dims) @@ -648,7 +656,7 @@ end #*USNTABLE USE Val{false} IF WE REALLY WANT THIS FLAG function reshape(arr::NDArray, i::Int...; copy::Bool=false) - return reshape(arr, i; copy = copy) + return reshape(arr, i; copy=copy) end # Ignore the scalar indexing here... diff --git a/src/ndarray/unary.jl b/src/ndarray/unary.jl index 3b12e73d..8be56e77 100644 --- a/src/ndarray/unary.jl +++ b/src/ndarray/unary.jl @@ -110,6 +110,51 @@ function Base.:(-)(input::NDArray{T}) where {T} return nda_unary_op!(out, cuNumeric.NEGATIVE, input) end +function Base.real(input::NDArray{T}) where {T<:Complex} + T_OUT = Base.promote_op(real, T) + out = cuNumeric.zeros(T_OUT, size(input)) + return nda_unary_op!(out, cuNumeric.REAL, input) +end +Base.real(input::NDArray{<:Real}) = input + +function Base.imag(input::NDArray{T}) where {T<:Complex} + T_OUT = Base.promote_op(imag, T) + out = cuNumeric.zeros(T_OUT, size(input)) + return nda_unary_op!(out, cuNumeric.IMAG, input) +end +Base.imag(input::NDArray{T}) where {T<:Real} = cuNumeric.zeros(T, size(input)) + +function Base.conj(input::NDArray{T}) where {T<:Complex} + out = cuNumeric.zeros(T, size(input)) + return nda_unary_op!(out, cuNumeric.CONJ, input) +end +Base.conj(input::NDArray{<:Real}) = input + +# Broadcoast support for complex ops +@inline function __broadcast(f::typeof(Base.real), out::NDArray, input::NDArray{<:Complex}) + return nda_unary_op!(out, cuNumeric.REAL, input) +end +@inline function __broadcast(f::typeof(Base.imag), out::NDArray, input::NDArray{<:Complex}) + return nda_unary_op!(out, cuNumeric.IMAG, input) +end +@inline function __broadcast(f::typeof(Base.conj), out::NDArray, input::NDArray{<:Complex}) + return nda_unary_op!(out, cuNumeric.CONJ, input) +end + +# Fallbacks for Real types +@inline function __broadcast(f::typeof(Base.real), out::NDArray, input::NDArray{<:Real}) + # real(real_array) is just the array + return nda_unary_op!(out, cuNumeric.IDENTITY, input) +end +@inline function __broadcast(f::typeof(Base.imag), out::NDArray, input::NDArray{<:Real}) + # imag(real_array) is all zeros + return nda_binary_op!(out, cuNumeric.SUBTRACT, input, input) +end +@inline function __broadcast(f::typeof(Base.conj), out::NDArray, input::NDArray{<:Real}) + # conj(real_array) is just the array + return nda_unary_op!(out, cuNumeric.IDENTITY, input) +end + function Base.:(-)(input::NDArray{Bool}) return -(checked_promote_arr(input, DEFAULT_INT)) end @@ -157,8 +202,8 @@ end for (julia_fn, op_code) in unary_op_map_no_args @eval begin @inline function __broadcast( - f::typeof($julia_fn), out::NDArray{T}, input::NDArray{T} - ) where {T} + f::typeof($julia_fn), out::NDArray{A}, input::NDArray{B} + ) where {A,B} return nda_unary_op!(out, $(op_code), input) end end diff --git a/test/runtests.jl b/test/runtests.jl index bad87b28..9a958582 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -1,4 +1,4 @@ -#= Copyright 2025 Northwestern University, +#= Copyright 2026 Northwestern University, * Carnegie Mellon University University * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -77,8 +77,7 @@ end @testset verbose = true "GEMM" begin N = 50 M = 25 - @testset verbose = true for T in Base.uniontypes(cuNumeric.SUPPORTED_TYPES) - # @warn "SGEMM has some precision issues, using tol $(rtol(T)) 🥲" + @testset verbose = true for T in Base.uniontypes(cuNumeric.SUPPORTED_NUMERIC_TYPES) gemm(N, M, T, rtol(T)) end end @@ -87,8 +86,8 @@ end @testset verbose = true "Unary Ops w/o Args" begin N = 100 # keep as perfect square - @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_TYPES) - allowpromotion(T == Bool || T == Int32) do + @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_ARRAY_TYPES) + allowpromotion(true) do test_unary_function_set(cuNumeric.floaty_unary_ops_no_args, T, N) end @@ -108,31 +107,46 @@ end end end - #!SPECIAL CASES (!, -) + # Special cases for complex-related unary ops + @testset "Complex Unary Ops (real, imag, conj)" begin + if T <: Complex + arr = my_rand(T, N) + arr_cn = NDArray(arr) + + allowscalar() do + allowpromotion(true) do + @test cuNumeric.compare(real(arr), real(arr_cn), atol(T), rtol(T)) + @test cuNumeric.compare(imag(arr), imag(arr_cn), atol(T), rtol(T)) + @test cuNumeric.compare(conj(arr), conj(arr_cn), atol(T), rtol(T)) + + @test cuNumeric.compare(real.(arr), real.(arr_cn), atol(T), rtol(T)) + @test cuNumeric.compare(imag.(arr), imag.(arr_cn), atol(T), rtol(T)) + @test cuNumeric.compare(conj.(arr), conj.(arr_cn), atol(T), rtol(T)) + end + end + end + end end end @testset verbose = true "Unary Reductions" begin N = 100 - @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_TYPES) + @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_ARRAY_TYPES) julia_arr = my_rand(T, N) cunumeric_arr = @allowscalar NDArray(julia_arr) @testset "$(reduction)" for reduction in keys(cuNumeric.unary_reduction_map) - enable_sum_promotion = (T == Int32 || T == Bool) && (reduction == Base.sum) - enable_prod_promotion = (T == Int32) && (reduction == Base.prod) - - # Test promotion errors cause we can: - if enable_sum_promotion - @test_throws "Implicit promotion" reduction(cunumeric_arr) - end - - if enable_prod_promotion - @test_throws "Implicit promotion" reduction(cunumeric_arr) + # Skip reductions not supported by the cuNumeric backend for complex types + if T <: Complex && ( + reduction == Base.maximum || + reduction == Base.minimum || + reduction == Base.prod + ) + continue end - allowpromotion(enable_sum_promotion || enable_prod_promotion) do + allowpromotion(true) do cunumeric_res = reduction(cunumeric_arr) julia_res = reduction(julia_arr) @@ -156,72 +170,76 @@ end @testset verbose = true "Binary Ops" begin N = 100 - @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_TYPES) - allowpromotion(T == Bool || T == Int32) do + @testset for T in Base.uniontypes(cuNumeric.SUPPORTED_ARRAY_TYPES) + allowpromotion(true) do test_binary_function_set(cuNumeric.floaty_binary_op_map, T, N) - end - - allowpromotion(T == Bool) do test_binary_function_set(cuNumeric.binary_op_map, T, N) end - # Special cases - @testset "lcm, gcd, ==, !=" begin - arr_jl = my_rand(T, N) - arr_jl2 = my_rand(T, N) - arr_cn = @allowscalar NDArray(arr_jl) - arr_cn2 = @allowscalar NDArray(arr_jl2) + arr_jl = my_rand(T, N) + arr_jl2 = my_rand(T, N) + arr_cn = @allowscalar NDArray(arr_jl) + arr_cn2 = @allowscalar NDArray(arr_jl2) - if T <: cuNumeric.SUPPORTED_INT_TYPES - allowscalar() do - @test cuNumeric.compare( - lcm.(arr_jl, arr_jl2), lcm.(arr_cn, arr_cn2), atol(T), rtol(T) - ) - @test cuNumeric.compare( - gcd.(arr_jl, arr_jl2), gcd.(arr_cn, arr_cn2), atol(T), rtol(T) - ) - end - end + # lcm/gcd require specific handling for integers and avoid overflow + if T <: cuNumeric.SUPPORTED_INT_TYPES && T != Bool + range_limit = (T == Int8 || T == UInt8) ? 10 : 100 + arr_jl_small = my_rand(T, N; L=1, R=range_limit) + arr_jl2_small = my_rand(T, N; L=1, R=range_limit) + arr_cn_small = @allowscalar NDArray(arr_jl_small) + arr_cn2_small = @allowscalar NDArray(arr_jl2_small) allowscalar() do - @test unwrap(arr_cn == arr_cn) - @test !unwrap(arr_cn == arr_cn2) - @test unwrap(arr_cn != arr_cn2) - @test !unwrap(arr_cn != arr_cn) - @test unwrap(all(arr_cn .== arr_cn)) + @test cuNumeric.compare( + lcm.(arr_jl_small, arr_jl2_small), lcm.(arr_cn_small, arr_cn2_small), atol(T), + rtol(T), + ) + @test cuNumeric.compare( + gcd.(arr_jl_small, arr_jl2_small), gcd.(arr_cn_small, arr_cn2_small), atol(T), + rtol(T), + ) end end - end - - @testset "Type and Shape Promotion" begin - cunumeric_arr1 = cuNumeric.zeros(Float64, N) - cunumeric_arr3 = cuNumeric.zeros(Float32, N) - cunumeric_int64 = cuNumeric.zeros(Int64, N) - cunumeric_int32 = cuNumeric.zeros(Int32, N) - cunumeric_arr5 = cuNumeric.zeros(Float64, N - 1, N - 1) - @test_throws "Implicit promotion" cunumeric_arr3 .+ cunumeric_arr1 - @test_throws "Implicit promotion" map(+, cunumeric_arr3, cunumeric_arr1) - @test_throws DimensionMismatch cunumeric_arr1 .+ cunumeric_arr5 - @test_throws DimensionMismatch cunumeric_arr1 ./ cunumeric_arr5 allowscalar() do - @test cuNumeric.compare( - cunumeric_arr1, cunumeric_int64 .+ cunumeric_arr1, atol(Float64), rtol(Float64) - ) - r1 = @allowpromotion cunumeric_arr3 .+ cunumeric_arr1 - r2 = @allowpromotion map(+, cunumeric_arr3, cunumeric_arr1) - @test cuNumeric.compare(r1, r2, atol(Float64), rtol(Float64)) + @test unwrap(arr_cn == arr_cn) + @test !unwrap(arr_cn == arr_cn2) + @test unwrap(arr_cn != arr_cn2) + @test !unwrap(arr_cn != arr_cn) + @test unwrap(all(arr_cn .== arr_cn)) end end +end - @testset "Copy-To" begin - a = cuNumeric.zeros(2, 2) - b = cuNumeric.ones(2, 2) - copyto!(a, b); - @test is_same(a, b) +@testset "Type and Shape Promotion" begin + N = 100 + cunumeric_arr1 = cuNumeric.zeros(Float64, N) + cunumeric_arr3 = cuNumeric.zeros(Float32, N) + cunumeric_int64 = cuNumeric.zeros(Int64, N) + cunumeric_int32 = cuNumeric.zeros(Int32, N) + cunumeric_arr5 = cuNumeric.zeros(Float64, N - 1, N - 1) + @test_throws "Implicit promotion" cunumeric_arr3 .+ cunumeric_arr1 + @test_throws "Implicit promotion" map(+, cunumeric_arr3, cunumeric_arr1) + @test_throws DimensionMismatch cunumeric_arr1 .+ cunumeric_arr5 + @test_throws DimensionMismatch cunumeric_arr1 ./ cunumeric_arr5 + + allowscalar() do + @test cuNumeric.compare( + cunumeric_arr1, cunumeric_int64 .+ cunumeric_arr1, atol(Float64), rtol(Float64) + ) + r1 = @allowpromotion cunumeric_arr3 .+ cunumeric_arr1 + r2 = @allowpromotion map(+, cunumeric_arr3, cunumeric_arr1) + @test cuNumeric.compare(r1, r2, atol(Float64), rtol(Float64)) end end +@testset "Copy-To" begin + a = cuNumeric.zeros(2, 2) + b = cuNumeric.ones(2, 2) + copyto!(a, b); + @test is_same(a, b) +end + #TODO LOOP BINARY OPS WITH SCALARS @testset verbose = true "Scalars" begin N = 10 @@ -235,7 +253,7 @@ end allowscalar() do cunumeric_arr = NDArray(julia_arr) cunumeric_arr_2D = NDArray(julia_arr_2D) - allowpromotion(T == Int32) do + allowpromotion(true) do for cn_arr in (cunumeric_arr, cunumeric_arr_2D) @test cuNumeric.compare(s * julia_arr, s * cunumeric_arr, atol(T), rtol(T)) @test cuNumeric.compare(julia_arr * s, cunumeric_arr * s, atol(T), rtol(T)) @@ -289,7 +307,7 @@ end get_pwrs(::Type{F}) where {F<:AbstractFloat} = F.([-3.141, -2, -1, 0, 1, 2, 3.2, 4.41, 6.233]) get_pwrs(::Type{Bool}) = [true, false, true, false, false, true, false, true, true] - # TYPES = Base.uniontypes(cuNumeric.SUPPORTED_TYPES) + # TYPES = Base.uniontypes(cuNumeric.SUPPORTED_ARRAY_TYPES) TYPES = Base.uniontypes(cuNumeric.SUPPORTED_FLOAT_TYPES) @testset "$(BT) ^ $(PT)" for (BT, PT) in Iterators.product(TYPES, TYPES) @@ -321,7 +339,7 @@ end TEST_BROKEN = (BT <: Union{Int32,Int64} && PT == Bool) - allowpromotion(sizeof(BT) != sizeof(PT)) do + allowpromotion(true) do allowscalar() do # Power is array @test cuNumeric.compare( @@ -346,7 +364,7 @@ end # Cast julia result to whatever we do res_jl = T_OUT.(arr_jl .^ -1) - allowpromotion(T == Bool || T == Int32) do + allowpromotion(true) do res_cn = arr_cn .^ -1 res_cn2 = inv.(arr_cn) allowscalar() do @@ -366,8 +384,10 @@ end res_jl = arr_jl .^ 2 res_cn = arr_cn .^ 2 - allowscalar() do - @test cuNumeric.compare(res_jl, res_cn, atol(T_OUT), rtol(T_OUT)) + allowpromotion(true) do + allowscalar() do + @test cuNumeric.compare(res_jl, res_cn, atol(T_OUT), rtol(T_OUT)) + end end end end diff --git a/test/tests/binary_tests.jl b/test/tests/binary_tests.jl index ea69da43..917c6c5c 100644 --- a/test/tests/binary_tests.jl +++ b/test/tests/binary_tests.jl @@ -19,12 +19,20 @@ end function test_binary_function_set(func_dict, T, N) skip = (Base.lcm, Base.gcd) + # not defined for complex. + skip_on_complex = ( + Base.:(<), Base.:(<=), Base.:(>), Base.:(>=), Base.max, Base.min, Base.atan, Base.hypot + ) @testset "$func" for func in keys(func_dict) # This is tested separately func == Base.:(^) && continue + if T <: Complex && (func in skip_on_complex) + continue + end + (func in skip) && continue arrs_jl = make_julia_arrays(T, N, :uniform; count=2) diff --git a/test/tests/gemm.jl b/test/tests/gemm.jl index 57fdc12b..1b713eed 100644 --- a/test/tests/gemm.jl +++ b/test/tests/gemm.jl @@ -17,20 +17,17 @@ * Ethan Meitz =# - - function gemm(N, M, T, max_diff) - if T == Bool - a = cuNumeric.trues(5,5) - b = cuNumeric.as_type(cuNumeric.trues(5,5), Float32) - c = cuNumeric.as_type(cuNumeric.trues(5,5), Float64) + a = cuNumeric.trues(5, 5) + b = cuNumeric.as_type(cuNumeric.trues(5, 5), Float32) + c = cuNumeric.as_type(cuNumeric.trues(5, 5), Float64) @test_throws ArgumentError a * a # Bool * Bool not supported @allowpromotion d = a * b @allowpromotion e = a * c @test @allowscalar cuNumeric.compare(5 * ones(Float32, 5, 5), d, 0.0, max_diff) @test @allowscalar cuNumeric.compare(5 * ones(Float64, 5, 5), e, 0.0, max_diff) - return + return nothing end if T <: Integer @@ -40,10 +37,10 @@ function gemm(N, M, T, max_diff) b_jl = ones(Float32, 5, 5) @test_throws ArgumentError a * a @test @allowscalar cuNumeric.compare(a_jl * b_jl, a * b, 0.0, max_diff) - return + return nothing end - dims_to_test = [(N,N), (N, M), (M, N)] + dims_to_test = [(N, N), (N, M), (M, N)] @testset for dims in dims_to_test # Base julia arrays @@ -52,18 +49,9 @@ function gemm(N, M, T, max_diff) C_out_cpu = zeros(T, dims[1], dims[1]) # cunumeric arrays - A = cuNumeric.zeros(T, dims[1], dims[2]) - B = cuNumeric.zeros(T, dims[2], dims[1]) - C_out = cuNumeric.zeros(T, dims[1], dims[1]) - - # Initialize NDArrays with random values - # used in Julia arrays - @allowscalar for i in 1:dims[1] - for j in 1:dims[2] - A[i, j] = A_cpu[i, j] - B[j, i] = B_cpu[j, i] - end - end + A = cuNumeric.NDArray(A_cpu) + B = cuNumeric.NDArray(B_cpu) + C_out = cuNumeric.zeros(T, dims[1], dims[1]) # Julia result C_cpu = A_cpu * B_cpu @@ -79,8 +67,8 @@ function gemm(N, M, T, max_diff) LinearAlgebra.mul!(C_out, A, B) allowscalar() do - @test isapprox(C, C_cpu, rtol = max_diff) - @test isapprox(C, C_out, rtol = max_diff) + @test isapprox(C, C_cpu, rtol=max_diff) + @test isapprox(C, C_out, rtol=max_diff) if T != Float64 C_wider = cuNumeric.zeros(Float64, dims[1], dims[1]) diff --git a/test/tests/stability.jl b/test/tests/stability.jl index d097454a..628a7be4 100644 --- a/test/tests/stability.jl +++ b/test/tests/stability.jl @@ -1,50 +1,54 @@ @testset verbose = true "core" begin a = cuNumeric.zeros(5) b = cuNumeric.zeros(Float64, 3, 4) - @inferred size(a) - @inferred size(b) - @inferred cuNumeric.shape(a) - @inferred cuNumeric.shape(b) + @test @inferred(size(a)) !== nothing + @test @inferred(size(b)) !== nothing + @test @inferred(cuNumeric.shape(a)) !== nothing + @test @inferred(cuNumeric.shape(b)) !== nothing end @testset verbose = true "construction" begin # zeros, zeros_like, ones, rand, fill, trues, falses\ for constructor in (:zeros, :ones) @eval begin - @inferred cuNumeric.$(constructor)(Float64, 3, 2) - @inferred cuNumeric.$(constructor)(Float64, (3, 4)) - @inferred cuNumeric.$(constructor)(3, 5, 6) - @inferred cuNumeric.$(constructor)((3,)) - @inferred cuNumeric.$(constructor)() - @inferred cuNumeric.$(constructor)(Int64) + @test @inferred(cuNumeric.$(constructor)(Float64, 3, 2)) !== nothing + @test @inferred(cuNumeric.$(constructor)(Float64, (3, 4))) !== nothing + @test @inferred(cuNumeric.$(constructor)(3, 5, 6)) !== nothing + @test @inferred(cuNumeric.$(constructor)((3,))) !== nothing + @test @inferred(cuNumeric.$(constructor)()) !== nothing + @test @inferred(cuNumeric.$(constructor)(Int64)) !== nothing end end a = cuNumeric.zeros(Float64, 5, 3) - @inferred cuNumeric.zeros_like(a) + @test @inferred(cuNumeric.zeros_like(a)) !== nothing for constructor in (:trues, :falses) @eval begin - @inferred cuNumeric.$(constructor)(5) - @inferred cuNumeric.$(constructor)((5,4)) - @inferred cuNumeric.$(constructor)(3, 4, 5) + @test @inferred(cuNumeric.$(constructor)(5)) !== nothing + @test @inferred(cuNumeric.$(constructor)((5, 4))) !== nothing + @test @inferred(cuNumeric.$(constructor)(3, 4, 5)) !== nothing end end - @inferred cuNumeric.fill(2.0, 3, 4) - @inferred cuNumeric.fill(2, (3, 4)) - @inferred cuNumeric.fill(2.0, 3) + @test @inferred(cuNumeric.fill(2.0, 3, 4)) !== nothing + @test @inferred(cuNumeric.fill(2, (3, 4))) !== nothing + @test @inferred(cuNumeric.fill(2.0, 3)) !== nothing - @inferred cuNumeric.rand(4, 3) - @inferred cuNumeric.rand(Float32, 5) + @test @inferred(cuNumeric.rand(4, 3)) !== nothing + @test @inferred(cuNumeric.rand(Float32, 5)) !== nothing + + # NDArray from Julia Array (Parent-stable attachment) + @test @inferred(cuNumeric.NDArray(rand(10))) !== nothing + @test @inferred(cuNumeric.NDArray(rand(Float32, 3, 3))) !== nothing end @testset verbose = true "conversion" begin # cast to array, as_type a = cuNumeric.zeros(Float64, 5, 5) - @inferred Array(a) - @inferred Array{Float32}(a) - @inferred cuNumeric.as_type(a, Float32) - @inferred cuNumeric.as_type(a, Int64) + @test @inferred(Array(a)) !== nothing + @test @inferred(Array{Float32}(a)) !== nothing + @test @inferred(cuNumeric.as_type(a, Float32)) !== nothing + @test @inferred(cuNumeric.as_type(a, Int64)) !== nothing end @testset verbose = true "indexing" begin @@ -52,26 +56,26 @@ end a = cuNumeric.zeros(Float32, 5, 5) b = cuNumeric.zeros(Int32, 11) - @inferred a[1:3, 1:3] - @inferred a[2, 1:3] - @inferred a[1, 1:3] .+ b[1:3] - @inferred b[1:5] - # @inferred a[1:3, 1:end] + @test @inferred(a[1:3, 1:3]) !== nothing + @test @inferred(a[2, 1:3]) !== nothing + @test @inferred(a[1, 1:3] .+ b[1:3]) !== nothing + @test @inferred(b[1:5]) !== nothing + # @test @inferred(a[1:3, 1:end]) !== nothing allowscalar() do - @inferred a[1, 2] + @test @inferred(a[1, 2]) !== nothing end end @testset verbose = true "broadcasting" begin a = cuNumeric.ones(Float32, 3, 3) b = cuNumeric.ones(Int32, 3, 3) - @inferred 5 .* a - @inferred 5.0f0 .* a - @inferred 5 * a - @inferred 5.0f0 * a + @test @inferred(5 .* a) !== nothing + @test @inferred(5.0f0 .* a) !== nothing + @test @inferred(5 * a) !== nothing + @test @inferred(5.0f0 * a) !== nothing - @inferred a .* b - @inferred a .+ b - @inferred a ./ b - @inferred ((a .* b) .+ a) .* 2.0f0 -end \ No newline at end of file + @test @inferred(a .* b) !== nothing + @test @inferred(a .+ b) !== nothing + @test @inferred(a ./ b) !== nothing + @test @inferred(((a .* b) .+ a) .* 2.0f0) !== nothing +end diff --git a/test/tests/unary_tests.jl b/test/tests/unary_tests.jl index 4e349c7b..5a172b90 100644 --- a/test/tests/unary_tests.jl +++ b/test/tests/unary_tests.jl @@ -9,21 +9,19 @@ const SPECIAL_DOMAINS = Dict( Base.sqrt => :positive, ) - function test_unary_operation(func, julia_arr, cunumeric_arr, T) - T_OUT = Base.promote_op(func, T) - + # Pre-allocate output arrays cunumeric_in_place = cuNumeric.zeros(T_OUT, size(julia_arr)...) - + # Compute results using different methods julia_res = func.(julia_arr) - + cunumeric_res = func.(cunumeric_arr) cunumeric_in_place .= func.(cunumeric_arr) cunumeric_res2 = map(func, cunumeric_arr) - + allowscalar() do @test cuNumeric.compare(julia_res, cunumeric_in_place, atol(T_OUT), rtol(T_OUT)) @test cuNumeric.compare(julia_res, cunumeric_res, atol(T_OUT), rtol(T_OUT)) @@ -31,14 +29,22 @@ function test_unary_operation(func, julia_arr, cunumeric_arr, T) end end -function test_unary_function_set(func_dict, T, N) +skip_on_integer = (Base.acosh, Base.atanh, Base.atan, Base.acos, Base.asin) +skip_on_bool = (Base.:(-), skip_on_integer...) +skip_on_complex = ( + Base.tanh, + Base.deg2rad, Base.rad2deg, Base.sign, Base.cbrt, + Base.exp2, Base.expm1, Base.log10, Base.log1p, Base.log2, + Base.acos, Base.asin, Base.atan, Base.acosh, Base.asinh, Base.atanh, +) +function test_unary_function_set(func_dict, T, N) default_generator = (T == Bool) ? :uniform : :unit_interval - skip_on_integer = (Base.atanh, Base.atan, Base.acos, Base.asin) - skip_on_bool = (Base.:(-), skip_on_integer...) - @testset "$func" for func in keys(func_dict) + if func in skip_on_complex && (T <: Complex) + continue + end # The are only defined for like 3 integers (-1, 0, 1) so just skip them if func in skip_on_integer && (T <: Integer) @@ -56,9 +62,11 @@ function test_unary_function_set(func_dict, T, N) skip && continue julia_arr_1D, julia_arr_2D = make_julia_arrays(T, N, domain_type) - cunumeric_arr_1D, cunumeric_arr_2D = make_cunumeric_arrays([julia_arr_1D], [julia_arr_2D], T, N) - + cunumeric_arr_1D, cunumeric_arr_2D = make_cunumeric_arrays( + [julia_arr_1D], [julia_arr_2D], T, N + ) + test_unary_operation(func, julia_arr_1D, cunumeric_arr_1D, T) test_unary_operation(func, julia_arr_2D, cunumeric_arr_2D, T) end -end \ No newline at end of file +end diff --git a/test/tests/util.jl b/test/tests/util.jl index 856f1a88..c8de3c9a 100644 --- a/test/tests/util.jl +++ b/test/tests/util.jl @@ -5,7 +5,7 @@ const DOMAIN_GENERATORS = Dict{Symbol,Function}( :unit_interval => (T, N) -> T(2) .* rand(T, N) .- one(T), :normal => (T, N) -> randn(T, N), :uniform => (T, N) -> rand(T, N), - :positive => (T, N) -> abs.(rand(T, N)), + :positive => (T, N) -> (x=rand(T, N); T <: Signed ? abs.(max.(x, -typemax(T))) : x), ) rtol(::Type{Float16}) = 1e-2 @@ -27,8 +27,16 @@ is_same(arr1::Array, arr2::Array) = (arr1 == arr2) function my_rand(::Type{F}, dims...; L=F(-1000), R=F(1000)) where {F<:AbstractFloat} L .+ (R-L) .* rand(F, dims...) end -function my_rand(::Type{I}, dims...; L=I(-255), R=I(255)) where {I<:Signed} - L .+ floor.(I, (R - L + 1) .* rand(dims...)) +function my_rand(::Type{I}, dims...; L=nothing, R=nothing) where {I<:Integer} + L_default = I <: Unsigned ? 0 : max(-255, Int64(typemin(I))) + R_default = I <: Unsigned ? 255 : min(255, Int64(typemax(I))) + L_val = isnothing(L) ? I(L_default) : I(L) + R_val = isnothing(R) ? I(R_default) : I(R) + res = Float64(L_val) .+ floor.((Float64(R_val) - Float64(L_val) + 1.0) .* rand(dims...)) + return floor.(I, res) +end +function my_rand(::Type{CT}, dims...; L=T(-100), R=T(100)) where {T,CT<:Complex{T}} + Complex.(my_rand(T, dims...; L=L, R=R), my_rand(T, dims...; L=L, R=R)) end my_rand(::Type{Bool}, dims...) = rand(Bool, dims...) From 1cf07f2c5b6482d71ccd842e071a68c9b3ff9104 Mon Sep 17 00:00:00 2001 From: krasow Date: Tue, 7 Apr 2026 23:20:53 -0500 Subject: [PATCH 2/2] add lifetime.jl for the Julia Array to NDArray modifications --- test/runtests.jl | 4 ++ test/tests/axpy.jl | 2 +- test/tests/axpy_advanced.jl | 2 +- test/tests/binary_tests.jl | 18 +++++++ test/tests/elementwise.jl | 2 +- test/tests/gemm.jl | 2 +- test/tests/lifetime.jl | 92 ++++++++++++++++++++++++++++++++++ test/tests/scoping-advanced.jl | 19 ++++++- test/tests/scoping.jl | 19 ++++++- test/tests/stability.jl | 19 +++++++ test/tests/unary_tests.jl | 18 +++++++ test/tests/util.jl | 19 +++++++ 12 files changed, 210 insertions(+), 6 deletions(-) create mode 100644 test/tests/lifetime.jl diff --git a/test/runtests.jl b/test/runtests.jl index 9a958582..d2396b3f 100644 --- a/test/runtests.jl +++ b/test/runtests.jl @@ -404,6 +404,10 @@ end include("tests/stability.jl") end +@testset verbose = true "NDArray Lifetime and Zero-Copy" begin + include("tests/lifetime.jl") +end + @testset verbose = true "Scoping" begin N = 100 diff --git a/test/tests/axpy.jl b/test/tests/axpy.jl index 4ab7a754..a25a1060 100644 --- a/test/tests/axpy.jl +++ b/test/tests/axpy.jl @@ -1,4 +1,4 @@ -#= Copyright 2025 Northwestern University, +#= Copyright 2026 Northwestern University, * Carnegie Mellon University University * * Licensed under the Apache License, Version 2.0 (the "License"); diff --git a/test/tests/axpy_advanced.jl b/test/tests/axpy_advanced.jl index fe87d513..58c4f02b 100644 --- a/test/tests/axpy_advanced.jl +++ b/test/tests/axpy_advanced.jl @@ -1,4 +1,4 @@ -#= Copyright 2025 Northwestern University, +#= Copyright 2026 Northwestern University, * Carnegie Mellon University University * * Licensed under the Apache License, Version 2.0 (the "License"); diff --git a/test/tests/binary_tests.jl b/test/tests/binary_tests.jl index 917c6c5c..0e8031b3 100644 --- a/test/tests/binary_tests.jl +++ b/test/tests/binary_tests.jl @@ -1,3 +1,21 @@ +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# function test_binary_operation(func, julia_arr1, julia_arr2, cunumeric_arr1, cunumeric_arr2, T) T_OUT = Base.promote_op(func, T, T) diff --git a/test/tests/elementwise.jl b/test/tests/elementwise.jl index 0a378a5d..49f22bb2 100644 --- a/test/tests/elementwise.jl +++ b/test/tests/elementwise.jl @@ -1,4 +1,4 @@ -#= Copyright 2025 Northwestern University, +#= Copyright 2026 Northwestern University, * Carnegie Mellon University University * * Licensed under the Apache License, Version 2.0 (the "License"); diff --git a/test/tests/gemm.jl b/test/tests/gemm.jl index 1b713eed..646eb0e0 100644 --- a/test/tests/gemm.jl +++ b/test/tests/gemm.jl @@ -1,4 +1,4 @@ -#= Copyright 2025 Northwestern University, +#= Copyright 2026 Northwestern University, * Carnegie Mellon University University * * Licensed under the Apache License, Version 2.0 (the "License"); diff --git a/test/tests/lifetime.jl b/test/tests/lifetime.jl new file mode 100644 index 00000000..9e5f037e --- /dev/null +++ b/test/tests/lifetime.jl @@ -0,0 +1,92 @@ +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# + +@testset "Zero-Copy Verification" begin + A = rand(Float64, 4, 4) + NA = NDArray(A) + + @allowscalar begin + @test all(A .== Array(NA)) + end + + @test pointer(A) == cuNumeric.get_ptr(NA) + + # modify julia array, verify ndarray sees it + A[1, 1] = 99.0 + @allowscalar begin + @test NA[1, 1] == 99.0 + end + + # modify ndarray, verify julia array sees it + @allowscalar begin + NA[2, 2] = 88.0 + end + @test A[2, 2] == 88.0 +end + +@testset "Lifetime Protection" begin + function create_attached_ndarray() + local_A = rand(Float32, 100) + local_A[1] = 1.23f0 + return NDArray(local_A), local_A[1] + end + + NA, expected_val = create_attached_ndarray() + + # force gc to try and collect the local array + GC.gc(true) + GC.gc(true) # do it again + GC.gc(true) # and again lol + + # data should still be intact via parent reference + @allowscalar begin + @test NA[1] == expected_val + end + + # operations should still work + NA2 = NA .* 2.0f0 + @allowscalar begin + @test NA2[1] == expected_val * 2.0f0 + end +end + +@testset "Type Conversion Lifetime" begin + function create_typed_ndarray() + I = collect(1:10) + return NDArray{Float64}(I) + end + + NA = create_typed_ndarray() + + # the temporary Float64 array should be kept alive by NA.parent + GC.gc(true) + GC.gc(true) + GC.gc(true) + + @allowscalar begin + @test NA[1] == 1.0 + @test NA[10] == 10.0 + end + + # modification should work on the attached temporary + @allowscalar begin + NA[1] = 42.0 + @test NA[1] == 42.0 + end +end diff --git a/test/tests/scoping-advanced.jl b/test/tests/scoping-advanced.jl index 781c28e7..9977a2e6 100644 --- a/test/tests/scoping-advanced.jl +++ b/test/tests/scoping-advanced.jl @@ -1,4 +1,21 @@ -using cuNumeric +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# struct ParamsGS{T<:AbstractFloat} dx::T diff --git a/test/tests/scoping.jl b/test/tests/scoping.jl index 1df27515..aff2f28c 100644 --- a/test/tests/scoping.jl +++ b/test/tests/scoping.jl @@ -1,4 +1,21 @@ -using cuNumeric +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# function run_test(op, op_scope, FT, N) a = cuNumeric.rand(FT, (N, N)) diff --git a/test/tests/stability.jl b/test/tests/stability.jl index 628a7be4..648f4bc4 100644 --- a/test/tests/stability.jl +++ b/test/tests/stability.jl @@ -1,3 +1,22 @@ +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# + @testset verbose = true "core" begin a = cuNumeric.zeros(5) b = cuNumeric.zeros(Float64, 3, 4) diff --git a/test/tests/unary_tests.jl b/test/tests/unary_tests.jl index 5a172b90..88f284ff 100644 --- a/test/tests/unary_tests.jl +++ b/test/tests/unary_tests.jl @@ -1,3 +1,21 @@ +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# # Map functions to their required domains const SPECIAL_DOMAINS = Dict( diff --git a/test/tests/util.jl b/test/tests/util.jl index c8de3c9a..0035b776 100644 --- a/test/tests/util.jl +++ b/test/tests/util.jl @@ -1,3 +1,22 @@ +#= Copyright 2026 Northwestern University, + * Carnegie Mellon University University + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Author(s): David Krasowska + * Ethan Meitz +=# + # Domain-aware test data generation const DOMAIN_GENERATORS = Dict{Symbol,Function}(