Skip to content

Commit

Permalink
rename back to gpuarrays
Browse files Browse the repository at this point in the history
  • Loading branch information
SimonDanisch committed Mar 27, 2017
1 parent bae703b commit e405a8b
Show file tree
Hide file tree
Showing 23 changed files with 131 additions and 131 deletions.
2 changes: 1 addition & 1 deletion .travis.yml
Expand Up @@ -21,4 +21,4 @@ notifications:
# uncomment the following lines to override the default test script
script:
- if [[ -a .git/shallow ]]; then git fetch --unshallow; fi
- julia -e 'Pkg.clone(pwd()); Pkg.build("JTensors"); Pkg.test("JTensors"; coverage=true)'
- julia -e 'Pkg.clone(pwd()); Pkg.build("GPUArrays"); Pkg.test("GPUArrays"; coverage=true)'
2 changes: 1 addition & 1 deletion LICENSE.md
@@ -1,4 +1,4 @@
The JTensors.jl package is licensed under the MIT "Expat" License:
The GPUArrays.jl package is licensed under the MIT "Expat" License:

> Copyright (c) 2016: Simon Danisch.
>
Expand Down
22 changes: 11 additions & 11 deletions README.md
@@ -1,6 +1,6 @@
# JTensors (formerly GPUArrays)
# GPUArrays (formerly GPUArrays)

[![Build Status](https://travis-ci.org/SimonDanisch/JTensors.jl.svg?branch=master)](https://travis-ci.org/SimonDanisch/JTensors.jl)
[![Build Status](https://travis-ci.org/SimonDanisch/GPUArrays.jl.svg?branch=master)](https://travis-ci.org/SimonDanisch/GPUArrays.jl)


Prototype for a GPU Array library.
Expand All @@ -14,20 +14,20 @@ CUDAnative.jl is using (via LLVM + SPIR-V).

This allows to get more involved functionality, like complex arithmetic, for free, since we can compile what's already in Julia Base.

JTensors relies heavily on dot broadcasting. The great thing about dot broadcasting in Julia is, that it [actually fuses operations syntactically](http://julialang.org/blog/2017/01/moredots), which is vital for performance on the GPU.
GPUArrays relies heavily on dot broadcasting. The great thing about dot broadcasting in Julia is, that it [actually fuses operations syntactically](http://julialang.org/blog/2017/01/moredots), which is vital for performance on the GPU.
E.g.:

```Julia
out .= a .+ b ./ c .+ 1
```

Will result in one GPU kernel call to a function that combines the operations without any extra allocations.
This allows JTensors to offer a lot of functionality with minimal code.
This allows GPUArrays to offer a lot of functionality with minimal code.

#### Main type:

```Julia
type JTensor{T, N, B, C} <: DenseArray{T, N}
type GPUArray{T, N, B, C} <: DenseArray{T, N}
buffer::B # GPU buffer, allocated by context
size::NTuple{N, Int} # size of the array
context::C # GPU context
Expand All @@ -43,26 +43,26 @@ Planned backends: OpenGL, Vulkan
Implemented for all backends:

```Julia
map(f, ::JTensor...)
map!(f, dest::JTensor, ::JTensor...)
map(f, ::GPUArray...)
map!(f, dest::GPUArray, ::GPUArray...)

# maps
mapidx(f, A::JTensor, args...) do idx, a, args...
mapidx(f, A::GPUArray, args...) do idx, a, args...
# e.g
if idx < length(A)
a[idx+1] = a[idx]
end
end


broadcast(f, ::JTensor...)
broadcast!(f, dest::JTensor, ::JTensor...)
broadcast(f, ::GPUArray...)
broadcast!(f, dest::GPUArray, ::GPUArray...)

```


CLFFT, CUFFT, CLBLAS and CUBLAS will soon be supported.
A prototype of generic support of these libraries can be found in [blas.jl](https://github.com/JuliaGPU/JTensors.jl/blob/sd/glsl/src/blas.jl).
A prototype of generic support of these libraries can be found in [blas.jl](https://github.com/JuliaGPU/GPUArrays.jl/blob/sd/glsl/src/blas.jl).
The OpenCL backend already supports mat mul via `CLBLAS.gemm!` and `fft!`/`ifft!`.
CUDAnative could support these easily as well, but we currently run into problems with the interactions of `CUDAdrv` and `CUDArt`.

Expand Down
4 changes: 2 additions & 2 deletions appveyor.yml
Expand Up @@ -28,7 +28,7 @@ build_script:
- C:\projects\julia\bin\julia -e "Pkg.clone(\"https://github.com/SimonDanisch/Matcha.jl.git\")"
- C:\projects\julia\bin\julia -e "Pkg.clone(\"https://github.com/SimonDanisch/Sugar.jl.git\")"
- C:\projects\julia\bin\julia -e "versioninfo();
Pkg.clone(pwd(), \"JTensors\"); Pkg.build(\"JTensors\")"
Pkg.clone(pwd(), \"GPUArrays\"); Pkg.build(\"GPUArrays\")"

test_script:
- C:\projects\julia\bin\julia --check-bounds=yes -e "Pkg.test(\"JTensors\")"
- C:\projects\julia\bin\julia --check-bounds=yes -e "Pkg.test(\"GPUArrays\")"
8 changes: 4 additions & 4 deletions deps/build.jl
@@ -1,6 +1,6 @@
info("""
This process will figure out which acceleration Packages you have installed
and therefore which backends JTensors can offer.
and therefore which backends GPUArrays can offer.
Theoretically available:
:cudanative, :julia, :opencl
Expand All @@ -21,10 +21,10 @@ install_cudanative = true
if !isdir(cudanative_dir)
info("""
Not installing CUDAnative backend. If you've installed CUDAnative.jl not in the
default location, consider building JTensors like this:
default location, consider building GPUArrays like this:
```
ENV[CUDANATIVE_PATH] = "path/to/CUDAnative/"
Pkg.build("JTensors")
Pkg.build("GPUArrays")
```
If not installed, you can get CUDAnative like this:
```
Expand All @@ -48,7 +48,7 @@ install_cudanative = try
true
catch e
info("CUDAnative doesn't seem to be usable and it won't be installed as a backend. Error: $e")
info("If error fixed, try Pkg.build(\"JTensors\") again!")
info("If error fixed, try Pkg.build(\"GPUArrays\") again!")
false
end
if install_cudanative
Expand Down
18 changes: 9 additions & 9 deletions examples/blackscholes.jl
@@ -1,6 +1,6 @@
using JTensors
import JTensors: JLBackend, CLBackend
import JTensors.JLBackend: JLArray
using GPUArrays
import GPUArrays: JLBackend, CLBackend
import GPUArrays.JLBackend: JLArray

function blackscholes(
sptprice,
Expand Down Expand Up @@ -45,12 +45,12 @@ result = similar(time)
result .= blackscholes.(sptprice, initStrike, rate, volatility, time)

ctx = CLBackend.init()
sptprice_gpu = JTensor(sptprice, flag = :r)
initStrike_gpu = JTensor(initStrike, flag = :r)
rate_gpu = JTensor(rate, flag = :r)
volatility_gpu = JTensor(volatility, flag = :r)
time_gpu = JTensor(time, flag = :r)
result_gpu = JTensor(result, flag = :w)
sptprice_gpu = GPUArray(sptprice, flag = :r)
initStrike_gpu = GPUArray(initStrike, flag = :r)
rate_gpu = GPUArray(rate, flag = :r)
volatility_gpu = GPUArray(volatility, flag = :r)
time_gpu = GPUArray(time, flag = :r)
result_gpu = GPUArray(result, flag = :w)

using BenchmarkTools
@time test(
Expand Down
18 changes: 9 additions & 9 deletions examples/custom_kernels.jl
@@ -1,5 +1,5 @@
using JTensors
using JTensors.CLBackend
using GPUArrays
using GPUArrays.CLBackend
CLBackend.init()


Expand All @@ -10,8 +10,8 @@ function clmap!(f, out, b)
end
# we need a `guiding array`, to get the context and indicate on what size we
# want to execute the kernel! This kind of scheme might change in the future
x = JTensor(rand(Float32, 100))
y = JTensor(rand(Float32, 100))
x = GPUArray(rand(Float32, 100))
y = GPUArray(rand(Float32, 100))
func = CLFunction(x, clmap!, sin, x, y)
# same here, x is just passed to supply a kernel size!
func(x, sin, x, y)
Expand Down Expand Up @@ -47,7 +47,7 @@ __kernel void julia(
w = 2048 * 2;
h = 2048 * 2;
q = [Complex64(r,i) for i=1:-(2.0/w):-1, r=-1.5:(3.0/h):1.5];
q_buff = JTensor(q)
q_buff = GPUArray(q)
o_buff = similar(q_buff, UInt16)

# you need to pass the name of the kernel you'd like to compile
Expand All @@ -64,8 +64,8 @@ x /= maximum(x)
save("test.jpg", Gray.(x))


using JTensors
using JTensors.CUBackend
using GPUArrays
using GPUArrays.CUBackend
using CUDAnative
CUBackend.init()
function clmap!(f, out, b)
Expand All @@ -74,8 +74,8 @@ function clmap!(f, out, b)
return
end
# CUDA works exactly the same!
x = JTensor(rand(Float32, 100))
y = JTensor(rand(Float32, 100))
x = GPUArray(rand(Float32, 100))
y = GPUArray(rand(Float32, 100))
# CUFunction instead of CLFunction
# Note, that we need to use the sin of CUDAnative.
# This necessity will hopefully be removed soon
Expand Down
2 changes: 1 addition & 1 deletion examples/reduction.jl
@@ -1,4 +1,4 @@
using JTensors: CLBackend
using GPUArrays: CLBackend
ctx = CLBackend.init()
using OpenCL: cl
source = """
Expand Down
4 changes: 2 additions & 2 deletions src/JTensors.jl → src/GPUArrays.jl
@@ -1,11 +1,11 @@
__precompile__(true)
module JTensors
module GPUArrays

abstract Context
using Sugar

include("abstractarray.jl")
export JTensor, mapidx, linear_index
export GPUArray, mapidx, linear_index

include(joinpath("backends", "backends.jl"))
export is_backend_supported, supported_backends
Expand Down
14 changes: 7 additions & 7 deletions src/abstractarray.jl
@@ -1,6 +1,6 @@
abstract AbstractAccArray{T, N} <: DenseArray{T, N}

type JTensor{T, N, B, C} <: AbstractAccArray{T, N}
type GPUArray{T, N, B, C} <: AbstractAccArray{T, N}
buffer::B
size::NTuple{N, Int}
context::C
Expand Down Expand Up @@ -30,7 +30,7 @@ Base.eltype{T}(::AbstractAccArray{T}) = T
Base.size(A::AbstractAccArray) = A.size

function Base.show(io::IO, mt::MIME"text/plain", A::AbstractAccArray)
println(io, "JTensor with ctx: $(context(A)): ")
println(io, "GPUArray with ctx: $(context(A)): ")
show(io, mt, Array(A))
end
function Base.showcompact(io::IO, mt::MIME"text/plain", A::AbstractAccArray)
Expand All @@ -46,12 +46,12 @@ end
function Base.similar{T <: AbstractAccArray, N}(x::T, dims::NTuple{N, Int})
similar(x, eltype(x), dims)
end
function Base.similar{T <: JTensor, ET, N}(
function Base.similar{T <: GPUArray, ET, N}(
::Type{T}, ::Type{ET}, sz::NTuple{N, Int};
context::Context = current_context(), kw_args...
)
b = create_buffer(context, ET, sz; kw_args...)
JTensor{ET, N, typeof(b), typeof(context)}(b, sz, context)
GPUArray{ET, N, typeof(b), typeof(context)}(b, sz, context)
end


Expand Down Expand Up @@ -227,7 +227,7 @@ end
#############################
# reduce

# horrible hack to get around of fetching the first element of the JTensor
# horrible hack to get around of fetching the first element of the GPUArray
# as a startvalue, which is a bit complicated with the current reduce implementation
function startvalue(f, T)
error("Please supply a starting value for mapreduce. E.g: mapreduce($f, $op, 1, A)")
Expand Down Expand Up @@ -293,11 +293,11 @@ else
error("No Serialization type found. Probably unsupported Julia version")
end

function Base.serialize{T<:JTensor}(s::BaseSerializer, t::T)
function Base.serialize{T<:GPUArray}(s::BaseSerializer, t::T)
Base.serialize_type(s, T)
serialize(s, Array(t))
end
function Base.deserialize{T<:JTensor}(s::BaseSerializer, ::Type{T})
function Base.deserialize{T<:GPUArray}(s::BaseSerializer, ::Type{T})
A = deserialize(s)
T(A)
end
16 changes: 8 additions & 8 deletions src/backends/cudanative/cudanative.jl
@@ -1,12 +1,12 @@

module CUBackend

using ..JTensors, CUDAnative, StaticArrays, Compat
using ..GPUArrays, CUDAnative, StaticArrays, Compat

import CUDAdrv, CUDArt #, CUFFT

import JTensors: buffer, create_buffer, acc_broadcast!, acc_mapreduce, mapidx
import JTensors: Context, JTensor, context, broadcast_index, linear_index
import GPUArrays: buffer, create_buffer, acc_broadcast!, acc_mapreduce, mapidx
import GPUArrays: Context, GPUArray, context, broadcast_index, linear_index
using CUDAdrv: CuDefaultStream

immutable GraphicsResource{T}
Expand All @@ -28,8 +28,8 @@ function any_context()
CUContext(ctx, dev)
end

#typealias GLArrayImg{T, N} JTensor{T, N, gl.Texture{T, N}, GLContext}
@compat const CUArray{T, N, B} = JTensor{T, N, B, CUContext} #, GLArrayImg{T, N}}
#typealias GLArrayImg{T, N} GPUArray{T, N, gl.Texture{T, N}, GLContext}
@compat const CUArray{T, N, B} = GPUArray{T, N, B, CUContext} #, GLArrayImg{T, N}}
@compat const CUArrayBuff{T, N} = CUArray{T, N, CUDAdrv.CuArray{T, N}}


Expand All @@ -38,7 +38,7 @@ let contexts = CUContext[]
all_contexts() = copy(contexts)::Vector{CUContext}
current_context() = last(contexts)::CUContext
function init(;ctx = any_context())
JTensors.make_current(ctx)
GPUArrays.make_current(ctx)
push!(contexts, ctx)
ctx
end
Expand All @@ -60,7 +60,7 @@ end
function Base.similar{T, N, ET}(x::CUArray{T, N}, ::Type{ET}, sz::NTuple{N, Int}; kw_args...)
ctx = context(x)
b = create_buffer(ctx, ET, sz; kw_args...)
JTensor{ET, N, typeof(b), typeof(ctx)}(b, sz, ctx)
GPUArray{ET, N, typeof(b), typeof(ctx)}(b, sz, ctx)
end

function thread_blocks_heuristic(A::AbstractArray)
Expand Down Expand Up @@ -282,7 +282,7 @@ end


# TODO figure out how interact with CUDArt and CUDAdr
#GFFT = JTensor(Complex64, div(size(G,1),2)+1, size(G,2))
#GFFT = GPUArray(Complex64, div(size(G,1),2)+1, size(G,2))
# function Base.fft!(A::CUArray)
# G, GFFT = CUFFT.RCpair(A)
# fft!(G, GFFT)
Expand Down
16 changes: 8 additions & 8 deletions src/backends/interop/gl_cu.jl
Expand Up @@ -4,7 +4,7 @@ import CUDAdrv, CUDAnative, GLAbstraction
const cu = CUDArt.rt


cu_register(buff::JTensor) = cu_register(buffer(buff))
cu_register(buff::GPUArray) = cu_register(buffer(buff))

let graphic_resource_map_buff = Dict{UInt32, Ref{cu.cudaGraphicsResource_t}}(), graphic_resource_map_tex = Dict{UInt32, Ref{cu.cudaGraphicsResource_t}}()
function cu_register(
Expand Down Expand Up @@ -34,7 +34,7 @@ let graphic_resource_map_buff = Dict{UInt32, Ref{cu.cudaGraphicsResource_t}}(),
end
end

function cu_map(f, buff::JTensor)
function cu_map(f, buff::GPUArray)
cu_array = cu_map(buff)
f(cu_array)
cu_unmap(buff)
Expand All @@ -53,7 +53,7 @@ function cu_map{T, N}(A::GLBackend.GLArrayBuff{T, N}, graphics_ref=cu_register(A
cudevptr = Base.unsafe_convert(CUDAdrv.DevicePtr{T}, Ptr{T}(cubuff[]))
sz = size(A)
cuarr = CUDAdrv.CuArray{T}(sz, cudevptr)
JTensor(cuarr, sz, context=ctx)
GPUArray(cuarr, sz, context=ctx)
end

function cu_map{T, N}(A::GLBackend.GLArrayTex{T, N}, graphics_ref=cu_register(A), ctx=CUBackend.current_context())
Expand All @@ -67,7 +67,7 @@ function cu_map{T, N}(A::GLBackend.GLArrayTex{T, N}, graphics_ref=cu_register(A)
cudevptr = Base.unsafe_convert(CUDAdrv.DevicePtr{T}, Ptr{T}(cubuff[]))
sz = size(A)
cuarr = CUDAdrv.CuArray{T}(sz, cudevptr)
JTensor(cuarr, sz, context=ctx)
GPUArray(cuarr, sz, context=ctx)
end

function cu_unmap{T, N}(A::GLBackend.GLArray{T, N}, graphics_ref=cu_register(A))
Expand Down Expand Up @@ -95,7 +95,7 @@ end



# type CUDAGLBuffer{T} <: JTensor{T, 1}
# type CUDAGLBuffer{T} <: GPUArray{T, 1}
# buffer::GLBuffer{T}
# graphics_resource::Ref{cu.cudaGraphicsResource_t}
# ismapped::Bool
Expand Down Expand Up @@ -165,7 +165,7 @@ end
#
#
#
# function register_with_cuda{T, ND}(tex::JTensor{Texture{T, ND}}, image::CUImage)
# function register_with_cuda{T, ND}(tex::GPUArray{Texture{T, ND}}, image::CUImage)
# graphics_ref = Ref{cu.cudaGraphicsResource_t}()
# cuGraphicsGLRegisterImage(
# graphics_ref,
Expand All @@ -185,12 +185,12 @@ end
# )
# cuGraphicsUnmapResources(1, graphics_ref, C_NULL)
# cubuff = CuArray(cubuff[], size(tex), length(tex))
# JTensor(cubuff, size(tex), cuctx)
# GPUArray(cubuff, size(tex), cuctx)
# end
#
#
# function register_with_cuda{T}(
# glbuffer::JTensor{GLBuffer{T}};
# glbuffer::GPUArray{GLBuffer{T}};
# cuctx = first(cuda_contexts()),
# flag=cu.cudaGraphicsMapFlagsNone
# )
Expand Down

0 comments on commit e405a8b

Please sign in to comment.