Initial import
This commit is contained in:
commit
56a50eead3
820 changed files with 192077 additions and 0 deletions
52
experiments/CIC/cic_filter.py
Normal file
52
experiments/CIC/cic_filter.py
Normal file
|
|
@ -0,0 +1,52 @@
|
|||
import numpy as np
|
||||
import cosmotool as ct
|
||||
from tqdm import tqdm
|
||||
|
||||
Nmc=1
|
||||
L=1.0
|
||||
Ng=256
|
||||
Np=128**3
|
||||
Nbins=50
|
||||
alpha = (L/Ng)**3
|
||||
mass = (1.0*Ng**3) / Np
|
||||
|
||||
ik = np.fft.fftfreq(Ng, d=L/Ng)*2*np.pi
|
||||
kx,ky,kz=np.meshgrid(ik,ik,ik[:(Ng//2+1)],indexing='ij')
|
||||
k = np.sqrt(kx**2+ky**2+kz**2)
|
||||
var_in_k =Np * mass**2 * alpha**2 / L**3
|
||||
Hw,_ = np.histogram(k, range=(0,k.max()), bins=Nbins)
|
||||
|
||||
W= lambda q: (np.sinc(q/(2.*np.pi)*L/Ng)**2)
|
||||
|
||||
Href,_ = np.histogram(k, weights=(W(kx)*W(ky)*W(kz))**2, range=(0,k.max()), bins=Nbins)
|
||||
Href = (Href/Hw)[1:]
|
||||
|
||||
rho = np.zeros(Nbins-1)
|
||||
rho_stag = np.zeros(Nbins-1)
|
||||
|
||||
for i in tqdm( range(Nmc)):
|
||||
x = np.random.uniform(size=Np*3).reshape((Np,3))
|
||||
d = ct.project_cic(x, None, Ng, L, True, False)
|
||||
d_stag = 0.5*(d+ ct.project_cic(x+L/(2*Ng), None, Ng, L, True, False))
|
||||
d *= mass
|
||||
d_stag *= mass
|
||||
d -= 1
|
||||
d_stag -= 1
|
||||
|
||||
dhat = np.fft.rfftn(d) * alpha
|
||||
|
||||
H,b = np.histogram(k, weights=(dhat.real**2+dhat.imag**2)/L**3, range=(0,k.max()), bins=Nbins)
|
||||
dhat = np.fft.rfftn(d_stag) * alpha
|
||||
Hstag,b = np.histogram(k, weights=(dhat.real**2+dhat.imag**2)/L**3, range=(0,k.max()), bins=Nbins)
|
||||
|
||||
H = H / Hw
|
||||
Hstag = Hstag / Hw
|
||||
bc = 0.5*(b[1:]+b[:-1])
|
||||
bc = bc[1:]
|
||||
|
||||
rho += H[1:]
|
||||
rho_stag += Hstag[1:]
|
||||
|
||||
rho /= Nmc
|
||||
rho_stag /= Nmc
|
||||
|
||||
236
experiments/CIC/docic.py
Normal file
236
experiments/CIC/docic.py
Normal file
|
|
@ -0,0 +1,236 @@
|
|||
#+
|
||||
# ARES/HADES/BORG Package -- ./experiments/CIC/docic.py
|
||||
# Copyright (C) 2014-2020 Guilhem Lavaux <guilhem.lavaux@iap.fr>
|
||||
# Copyright (C) 2009-2020 Jens Jasche <jens.jasche@fysik.su.se>
|
||||
#
|
||||
# Additional contributions from:
|
||||
# Guilhem Lavaux <guilhem.lavaux@iap.fr> (2023)
|
||||
#
|
||||
#+
|
||||
import cosmotool as ct
|
||||
import numpy as np
|
||||
import pyopencl as cl
|
||||
import pyopencl.array as cl_array
|
||||
|
||||
CIC_PREKERNEL='''
|
||||
#define NDIM {ndim}
|
||||
typedef {cicType} BASIC_TYPE;
|
||||
|
||||
'''
|
||||
|
||||
CIC_KERNEL='''///CL///
|
||||
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
|
||||
|
||||
__kernel void init_pcell(__global int *p_cell, const int value)
|
||||
{
|
||||
int i = get_global_id(0);
|
||||
p_cell[i] = value;
|
||||
}
|
||||
|
||||
__kernel void build_indices(__global const BASIC_TYPE *pos,
|
||||
__global int *part_mesh, __global int *part_list, const int N, const BASIC_TYPE delta)
|
||||
{
|
||||
int i_part = get_global_id(0);
|
||||
long shifter = 1;
|
||||
long idx = 0;
|
||||
int d;
|
||||
|
||||
for (d = 0; d < NDIM; d++) {
|
||||
BASIC_TYPE x = pos[i_part*NDIM + d];
|
||||
int m = (int)floor(x*delta) %% N;
|
||||
|
||||
idx += shifter * m;
|
||||
shifter *= N;
|
||||
}
|
||||
|
||||
// Head of the list
|
||||
int initial_elt = atom_xchg(&part_mesh[idx], i_part);
|
||||
if (initial_elt == -1) {
|
||||
return;
|
||||
}
|
||||
// Point the next pointer of old_end to i_part
|
||||
part_list[i_part] = initial_elt;
|
||||
}
|
||||
|
||||
__kernel void reverse_list(__global int *part_mesh, __global int *part_list)
|
||||
{
|
||||
int mid = get_global_id(0);
|
||||
|
||||
int current_part = part_mesh[mid];
|
||||
if (current_part >= 0) {
|
||||
int next_part = part_list[current_part];
|
||||
part_list[current_part] = -1;
|
||||
while (next_part != -1) {
|
||||
int p = part_list[next_part];
|
||||
part_list[next_part] = current_part;
|
||||
current_part = next_part;
|
||||
next_part = p;
|
||||
}
|
||||
part_mesh[mid] = current_part;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void dance(__global const BASIC_TYPE *pos,
|
||||
__global BASIC_TYPE *density,
|
||||
__global int *part_mesh, __global int *part_list, const int N, const BASIC_TYPE delta)
|
||||
{
|
||||
int m[NDIM];
|
||||
int shifter = 1;
|
||||
int i;
|
||||
int first, i_part;
|
||||
int idx = 0;
|
||||
|
||||
for (i = 0; i < NDIM; i++) {
|
||||
m[i] = get_global_id(i);
|
||||
idx += shifter * m[i];
|
||||
shifter *= N;
|
||||
}
|
||||
|
||||
first = 1;
|
||||
|
||||
//BEGIN LOOPER
|
||||
%(looperFor)s
|
||||
//END LOOPER
|
||||
|
||||
int idx_dance = 0;
|
||||
BASIC_TYPE w = 0;
|
||||
//LOOPER INDEX
|
||||
int r[NDIM] = { %(looperVariables)s };
|
||||
//END LOOPER
|
||||
|
||||
i_part = part_mesh[idx];
|
||||
while (i_part != -1) {
|
||||
BASIC_TYPE w0 = 1;
|
||||
|
||||
for (int d = 0; d < NDIM; d++) {
|
||||
BASIC_TYPE x = pos[i_part*NDIM + d]*delta;
|
||||
BASIC_TYPE q = floor(x);
|
||||
BASIC_TYPE dx = x - q;
|
||||
|
||||
w0 *= (r[d] == 1) ? dx : ((BASIC_TYPE)1-dx);
|
||||
}
|
||||
|
||||
i_part = part_list[i_part];
|
||||
w += w0;
|
||||
}
|
||||
|
||||
shifter = 1;
|
||||
for (i = 0; i < NDIM; i++) {
|
||||
idx_dance += shifter * ((m[i]+r[i])%%N);
|
||||
shifter *= N;
|
||||
}
|
||||
|
||||
density[idx_dance] += w;
|
||||
|
||||
// One dance done. Wait for everybody for the next iteration
|
||||
barrier(CLK_GLOBAL_MEM_FENCE);
|
||||
%(looperForEnd)s
|
||||
}
|
||||
'''
|
||||
|
||||
class CIC_CL(object):
|
||||
|
||||
def __init__(self, context, ndim=2, ktype=np.float32):
|
||||
global CIC_PREKERNEL, CIC_KERNEL
|
||||
|
||||
translator = {}
|
||||
if ktype == np.float32:
|
||||
translator['cicType'] = 'float'
|
||||
pragmas = ''
|
||||
elif ktype == np.float64:
|
||||
translator['cicType'] = 'double'
|
||||
pragmas = '#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n'
|
||||
else:
|
||||
raise ValueError("Invalid ktype")
|
||||
|
||||
# 2 dimensions
|
||||
translator['ndim'] = ndim
|
||||
looperVariables = ','.join(['id%d' % d for d in xrange(ndim)])
|
||||
looperFor = '\n'.join(['for (int id{dim}=0; id{dim} < 2; id{dim}++) {{'.format(dim=d) for d in xrange(ndim)])
|
||||
looperForEnd = '}' * ndim
|
||||
|
||||
kern = pragmas + CIC_PREKERNEL.format(**translator) + (CIC_KERNEL % {'looperVariables': looperVariables, 'looperFor': looperFor, 'looperForEnd':looperForEnd})
|
||||
self.kern_code = kern
|
||||
self.ctx = context
|
||||
self.queue = cl.CommandQueue(context)#, properties=cl.OUT_OF_ORDER_EXEC_MODE_ENABLE)
|
||||
self.ktype = ktype
|
||||
self.ndim = ndim
|
||||
self.prog = cl.Program(self.ctx, kern).build()
|
||||
|
||||
def run(self, particles, Ng, L):
|
||||
assert particles.strides[1] == self.ktype().itemsize # This is C-ordering
|
||||
assert particles.shape[1] == self.ndim
|
||||
|
||||
print("Start again")
|
||||
|
||||
ndim = self.ndim
|
||||
part_pos = cl_array.to_device(self.queue, particles)
|
||||
part_mesh = cl_array.empty(self.queue, (Ng,)*ndim, np.int32, order='C')
|
||||
density = cl_array.zeros(self.queue, (Ng,)*ndim, self.ktype, order='C')
|
||||
part_list = cl_array.empty(self.queue, (particles.shape[0],), np.int32, order='C')
|
||||
|
||||
if True:
|
||||
delta = Ng/L
|
||||
|
||||
with ct.time_block("Init pcell array"):
|
||||
e = self.prog.init_pcell(self.queue, (Ng**ndim,), None, part_mesh.data, np.int32(-1))
|
||||
e.wait()
|
||||
|
||||
with ct.time_block("Init idx array"):
|
||||
e=self.prog.init_pcell(self.queue, (particles.shape[0],), None, part_list.data, np.int32(-1))
|
||||
e.wait()
|
||||
|
||||
with ct.time_block("Build indices"):
|
||||
self.prog.build_indices(self.queue, (particles.shape[0],), None,
|
||||
part_pos.data, part_mesh.data, part_list.data, np.int32(Ng), self.ktype(delta))
|
||||
|
||||
if True:
|
||||
with ct.time_block("Reverse list"):
|
||||
lastevt = self.prog.reverse_list(self.queue, (Ng**ndim,), None, part_mesh.data, part_list.data)
|
||||
# We require pmax pass, particles are ordered according to part_idx
|
||||
|
||||
with ct.time_block("dance"):
|
||||
self.prog.dance(self.queue, (Ng,)*ndim, None, part_pos.data, density.data, part_mesh.data, part_list.data, np.int32(Ng), self.ktype(delta))
|
||||
|
||||
print("Grab result")
|
||||
self.queue.finish()
|
||||
print("Del ppos")
|
||||
del part_pos
|
||||
print("Del pm")
|
||||
del part_mesh
|
||||
print("Del pl")
|
||||
del part_list
|
||||
print("Del density")
|
||||
with ct.time_block("download"):
|
||||
return density.get()
|
||||
|
||||
def go_test():
|
||||
ctx = cl.create_some_context()
|
||||
queue = cl.CommandQueue(ctx)
|
||||
|
||||
Nparts = 512**3
|
||||
Ng = 512
|
||||
cic = CIC_CL(ctx,ndim=3,ktype=np.float32)
|
||||
|
||||
print ("Generate")
|
||||
parts = np.random.rand(Nparts, 3).astype(np.float32)
|
||||
# ix=(np.arange(512)*1.0/512).astype(np.float32)
|
||||
# parts = [ix[:,None,None].repeat(512,axis=1).repeat(512,axis=2),
|
||||
# ix[None,:,None].repeat(512,axis=0).repeat(512,axis=2),
|
||||
# ix[None,None,:].repeat(512,axis=0).repeat(512,axis=1)]
|
||||
# parts = np.array(parts).transpose((3,0,1,2)).reshape((Nparts,3),order='C')
|
||||
|
||||
with ct.time_block("run cic"):
|
||||
cic.run(parts, Ng, 1.0)
|
||||
print("Done")
|
||||
|
||||
with ct.time_block("run cic 2"):
|
||||
density1 = cic.run(parts, Ng, 1.0)
|
||||
|
||||
with ct.time_block("classic cic"):
|
||||
density2 = ct.project_cic((parts-0.5).astype(ct.DTYPE), None, Ng, 1.0, True)
|
||||
|
||||
return density1, density2
|
||||
|
||||
|
||||
vv = go_test()
|
||||
47
experiments/CIC/timing.py
Normal file
47
experiments/CIC/timing.py
Normal file
|
|
@ -0,0 +1,47 @@
|
|||
#+
|
||||
# ARES/HADES/BORG Package -- ./experiments/CIC/timing.py
|
||||
# Copyright (C) 2014-2020 Guilhem Lavaux <guilhem.lavaux@iap.fr>
|
||||
# Copyright (C) 2009-2020 Jens Jasche <jens.jasche@fysik.su.se>
|
||||
#
|
||||
# Additional contributions from:
|
||||
# Guilhem Lavaux <guilhem.lavaux@iap.fr> (2023)
|
||||
#
|
||||
#+
|
||||
import time
|
||||
from contextlib import contextmanager
|
||||
|
||||
@contextmanager
|
||||
def time_block(name):
|
||||
ts = time.time()
|
||||
yield
|
||||
te = time.time()
|
||||
|
||||
print '%s %2.2f sec' % \
|
||||
(name, te-ts)
|
||||
|
||||
def timeit(method):
|
||||
|
||||
def timed(*args, **kw):
|
||||
ts = time.time()
|
||||
result = method(*args, **kw)
|
||||
te = time.time()
|
||||
|
||||
print '%r (%r, %r) %2.2f sec' % \
|
||||
(method.__name__, args, kw, te-ts)
|
||||
return result
|
||||
|
||||
return timed
|
||||
|
||||
def timeit_quiet(method):
|
||||
|
||||
def timed(*args, **kw):
|
||||
ts = time.time()
|
||||
result = method(*args, **kw)
|
||||
te = time.time()
|
||||
|
||||
print '%r %2.2f sec' % \
|
||||
(method.__name__, te-ts)
|
||||
return result
|
||||
|
||||
return timed
|
||||
|
||||
Loading…
Add table
Add a link
Reference in a new issue