Skip to content

Commit 2748e7b

Browse files
committed
add some tests
1 parent 3a8f1ec commit 2748e7b

File tree

3 files changed

+139
-3
lines changed

3 files changed

+139
-3
lines changed

Project.toml

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -20,10 +20,9 @@ SPIRVIntrinsics = "71d1d633-e7e8-4a92-83a1-de8814b09ba8"
2020
SPIRV_LLVM_Backend_jll = "4376b9bf-cff8-51b6-bb48-39421dff0d0c"
2121
SPIRV_Tools_jll = "6ac6d60f-d740-5983-97d7-a4482c0689f4"
2222
StaticArrays = "90137ffa-7385-5640-81b9-e52037218182"
23-
pocl_jll = "627d6b7a-bbe6-5189-83e7-98cc0a5aeadd"
2423

2524
[sources]
26-
GPUCompiler = {path = "/home/simeon/.julia/dev/GPUCompiler"}
25+
GPUCompiler = {rev = "sds/additional_args", url = "https://github.com/JuliaGPU/GPUCompiler.jl.git"}
2726
SPIRVIntrinsics = {path = "lib/intrinsics"}
2827

2928
[compat]
@@ -45,4 +44,3 @@ SPIRV_LLVM_Backend_jll = "20"
4544
SPIRV_Tools_jll = "2025.1"
4645
StaticArrays = "1"
4746
julia = "1.10"
48-
pocl_jll = "7.0.0"

test/Project.toml

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3,6 +3,7 @@ Adapt = "79e6a3ab-5dfb-504d-930d-738a2a938a0e"
33
Dates = "ade2ca70-3891-5945-98fb-dc099432e06a"
44
Distributed = "8ba89e20-285c-5b6f-9357-94700520ee1b"
55
GPUArrays = "0c68f7d7-f131-5f86-a1c3-88cf8149b2d7"
6+
GPUCompiler = "61eb1bfa-7361-4325-ad38-22787b887f55"
67
IOCapture = "b5f81e59-6552-4d32-b1f0-c071b021bf89"
78
InteractiveUtils = "b77e0a4c-d291-57a0-90e8-8db25a27a240"
89
JLD2 = "033835bb-8acc-5ee8-8aae-3f567f8a3819"
@@ -22,5 +23,8 @@ Statistics = "10745b16-79ce-11e8-11f9-7d13ad32a3b2"
2223
Test = "8dfed614-e22c-5e08-85e1-65c5234f0b40"
2324
pocl_jll = "627d6b7a-bbe6-5189-83e7-98cc0a5aeadd"
2425

26+
[sources]
27+
GPUCompiler = {rev = "sds/additional_args", url = "https://github.com/JuliaGPU/GPUCompiler.jl.git"}
28+
2529
[compat]
2630
pocl_jll = "7.0"

test/device/random.jl

Lines changed: 134 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,134 @@
1+
using Random
2+
3+
const n = 256
4+
5+
function apply_seed(seed)
6+
if seed === missing
7+
# should result in different numbers across launches
8+
Random.seed!()
9+
# XXX: this currently doesn't work, because of the definition in Base,
10+
# `seed!(r::MersenneTwister=default_rng())`, which breaks overriding
11+
# `default_rng` with a non-MersenneTwister RNG.
12+
elseif seed !== nothing
13+
# should result in the same numbers
14+
Random.seed!(seed)
15+
elseif seed === nothing
16+
# should result in different numbers across launches,
17+
# as determined by the seed set during module loading.
18+
end
19+
end
20+
21+
@testset "rand($T), seed $seed" for T in (
22+
Int32, UInt32, Int64, UInt64, Float32, Float64,
23+
), seed in (nothing, #=missing,=# 1234)
24+
# different kernel invocations should get different numbers
25+
@testset "across launches" begin
26+
function kernel(A::AbstractArray{T}, seed) where {T}
27+
apply_seed(seed)
28+
tid = get_global_id(1)
29+
A[tid] = rand(T)
30+
return nothing
31+
end
32+
33+
a = OpenCL.zeros(T, n)
34+
b = OpenCL.zeros(T, n)
35+
36+
@opencl global_size=n local_size=n kernel(a, seed)
37+
@opencl global_size=n local_size=n kernel(b, seed)
38+
39+
if seed === nothing || seed === missing
40+
@test Array(a) != Array(b)
41+
else
42+
@test Array(a) == Array(b)
43+
end
44+
end
45+
46+
# multiple calls to rand should get different numbers
47+
@testset "across calls" begin
48+
function kernel(A::AbstractArray{T}, B::AbstractArray{T}, seed) where {T}
49+
apply_seed(seed)
50+
tid = get_global_id(1)
51+
A[tid] = rand(T)
52+
B[tid] = rand(T)
53+
return nothing
54+
end
55+
56+
a = OpenCL.zeros(T, n)
57+
b = OpenCL.zeros(T, n)
58+
59+
@opencl global_size=n local_size=n kernel(a, b, seed)
60+
61+
@test Array(a) != Array(b)
62+
end
63+
64+
# different threads should get different numbers
65+
@testset "across threads, dim $active_dim" for active_dim in 1:6
66+
function kernel(A::AbstractArray{T}, seed) where {T}
67+
apply_seed(seed)
68+
id = get_local_id(1) * get_local_id(2) * get_local_id(3) *
69+
get_group_id(1) * get_group_id(2) * get_group_id(3)
70+
if 1 <= id <= length(A)
71+
A[id] = rand(T)
72+
end
73+
return nothing
74+
end
75+
76+
tx, ty, tz, bx, by, bz = [dim == active_dim ? 3 : 1 for dim in 1:6]
77+
gx, gy, gz = tx*bx, ty*by, tz*bz
78+
a = OpenCL.zeros(T, 3)
79+
80+
@opencl local_size=(tx, ty, tz) global_size=(gx, gy, gz) kernel(a, seed)
81+
82+
# NOTE: we don't just generate two numbers and compare them, instead generating a
83+
# couple more and checking they're not all the same, in order to avoid
84+
# occasional collisions with lower-precision types (i.e., Float16).
85+
# TODO: why is the third dimension broken?
86+
@test length(unique(Array(a))) > 1 broken = active_dim == 3 || active_dim == 6
87+
end
88+
end
89+
90+
@testset "basic randn($T), seed $seed" for T in (
91+
Float32, Float64,
92+
), seed in (nothing, #=missing,=# 1234)
93+
function kernel(A::AbstractArray{T}, seed) where {T}
94+
apply_seed(seed)
95+
tid = get_global_id(1)
96+
A[tid] = randn(T)
97+
return
98+
end
99+
100+
a = OpenCL.zeros(T, n)
101+
b = OpenCL.zeros(T, n)
102+
103+
@opencl global_size=n local_size=n kernel(a, seed)
104+
@opencl global_size=n local_size=n kernel(b, seed)
105+
106+
if seed === nothing || seed === missing
107+
@test Array(a) != Array(b)
108+
else
109+
@test Array(a) == Array(b)
110+
end
111+
end
112+
113+
@testset "basic randexp($T), seed $seed" for T in (
114+
Float32, Float64,
115+
), seed in (nothing, #=missing,=# 1234)
116+
function kernel(A::AbstractArray{T}, seed) where {T}
117+
apply_seed(seed)
118+
tid = get_global_id(1)
119+
A[tid] = randexp(T)
120+
return
121+
end
122+
123+
a = OpenCL.zeros(T, n)
124+
b = OpenCL.zeros(T, n)
125+
126+
@opencl global_size=n local_size=n kernel(a, seed)
127+
@opencl global_size=n local_size=n kernel(b, seed)
128+
129+
if seed === nothing || seed === missing
130+
@test Array(a) != Array(b)
131+
else
132+
@test Array(a) == Array(b)
133+
end
134+
end

0 commit comments

Comments
 (0)