From 6ae8c7e2b9687ff91f831af2f5f3689354992649 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sat, 1 Aug 2020 21:53:52 +0800 Subject: [PATCH 01/15] app --- taichi/backends/opengl/shaders/random.glsl.h | 12 ++++++------ taichi/program/kernel.cpp | 18 ++---------------- 2 files changed, 8 insertions(+), 22 deletions(-) diff --git a/taichi/backends/opengl/shaders/random.glsl.h b/taichi/backends/opengl/shaders/random.glsl.h index ec9cc3872de3c..98319bf424c81 100644 --- a/taichi/backends/opengl/shaders/random.glsl.h +++ b/taichi/backends/opengl/shaders/random.glsl.h @@ -5,11 +5,11 @@ STR( uvec4 _rand_; void _init_rand() { - uint i = (54321 + gl_GlobalInvocationID.x) * (12345 + _rand_state_); - _rand_.x = 123456789 * i * 1000000007; - _rand_.y = 362436069; - _rand_.z = 521288629; - _rand_.w = 88675123; + uint i = (54321u + gl_GlobalInvocationID.x) * (12345u + uint(_rand_state_)); + _rand_.x = 123456789u * i * 1000000007u; + _rand_.y = 362436069u; + _rand_.z = 521288629u; + _rand_.w = 88675123u; // Yes, this is not an atomic operation, but just fine since no matter // how `_rand_state_` changes, `gl_GlobalInvocationID.x` can still help @@ -22,7 +22,7 @@ uint _rand_u32() { uint t = _rand_.x ^ (_rand_.x << 11); _rand_.xyz = _rand_.yzw; _rand_.w = (_rand_.w ^ (_rand_.w >> 19)) ^ (t ^ (t >> 8)); - return _rand_.w * 1000000007; + return _rand_.w * 1000000007u; } float _rand_f32() { diff --git a/taichi/program/kernel.cpp b/taichi/program/kernel.cpp index e1bca3e410b1c..808fcfbf32181 100644 --- a/taichi/program/kernel.cpp +++ b/taichi/program/kernel.cpp @@ -119,12 +119,8 @@ void Kernel::set_arg_float(int i, float64 d) { TI_ASSERT_INFO( !args[i].is_nparray, "Assigning a scalar value to a numpy array argument is not allowed"); - - ActionRecorder::get_instance().record( - "set_kernel_arg_float64", {ActionArg("kernel_name", name), - ActionArg("arg_id", i), ActionArg("val", d)}); - auto dt = args[i].dt; + if (dt == DataType::f32) { program.context.set_arg(i, (float32)d); } else if (dt == DataType::f64) { @@ -154,12 +150,8 @@ void Kernel::set_arg_int(int i, int64 d) { TI_ASSERT_INFO( !args[i].is_nparray, "Assigning scalar value to numpy array argument is not allowed"); - - ActionRecorder::get_instance().record( - "set_kernel_arg_int64", {ActionArg("kernel_name", name), - ActionArg("arg_id", i), ActionArg("val", d)}); - auto dt = args[i].dt; + if (dt == DataType::i32) { program.context.set_arg(i, (int32)d); } else if (dt == DataType::i64) { @@ -247,12 +239,6 @@ void Kernel::set_arg_nparray(int i, uint64 ptr, uint64 size) { TI_ASSERT_INFO(args[i].is_nparray, "Assigning numpy array to scalar argument is not allowed"); - ActionRecorder::get_instance().record( - "set_kernel_arg_ext_ptr", - {ActionArg("kernel_name", name), ActionArg("arg_id", i), - ActionArg("address", fmt::format("0x{:x}", ptr)), - ActionArg("array_size_in_bytes", (int64)size)}); - args[i].size = size; program.context.set_arg(i, ptr); } From a4eba6cafff038ad924d62045c23a777f22d9f0d Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sat, 1 Aug 2020 21:59:50 +0800 Subject: [PATCH 02/15] misc --- misc/action_mpm88.py | 89 ++++++++++++++++++++++++++++++++++++++++ misc/action_simple_uv.py | 15 +++++++ misc/cc_action_record.py | 36 ---------------- 3 files changed, 104 insertions(+), 36 deletions(-) create mode 100644 misc/action_mpm88.py create mode 100644 misc/action_simple_uv.py delete mode 100644 misc/cc_action_record.py diff --git a/misc/action_mpm88.py b/misc/action_mpm88.py new file mode 100644 index 0000000000000..b448683a8e238 --- /dev/null +++ b/misc/action_mpm88.py @@ -0,0 +1,89 @@ +# MPM-MLS in 88 lines of Taichi code, originally created by @yuanming-hu +import taichi as ti +ti.core.start_recording('record.yml') +ti.init(arch=ti.cc) + +n_particles = 8192 +n_grid = 128 +dx = 1 / n_grid +dt = 2e-4 + +p_rho = 1 +p_vol = (dx * 0.5)**2 +p_mass = p_vol * p_rho +gravity = 9.8 +bound = 3 +E = 400 + +x = ti.Vector.field(2, ti.f32, n_particles) +v = ti.Vector.field(2, ti.f32, n_particles) +C = ti.Matrix.field(2, 2, ti.f32, n_particles) +J = ti.field(ti.f32, n_particles) + +grid_v = ti.Vector.field(2, ti.f32, (n_grid, n_grid)) +grid_m = ti.field(ti.f32, (n_grid, n_grid)) + +@ti.kernel +def substep(): + for i, j in grid_m: + grid_v[i, j] = [0, 0] + grid_m[i, j] = 0 + for p in x: + Xp = x[p] / dx + base = int(Xp - 0.5) + fx = Xp - base + w = [0.5 * (1.5 - fx)**2, 0.75 - (fx - 1)**2, 0.5 * (fx - 0.5)**2] + stress = -dt * 4 * E * p_vol * (J[p] - 1) / dx**2 + affine = ti.Matrix([[stress, 0], [0, stress]]) + p_mass * C[p] + for i, j in ti.static(ti.ndrange(3, 3)): + offset = ti.Vector([i, j]) + dpos = (offset - fx) * dx + weight = w[i].x * w[j].y + grid_v[base + offset] += weight * (p_mass * v[p] + affine @ dpos) + grid_m[base + offset] += weight * p_mass + for i, j in grid_m: + if grid_m[i, j] > 0: + grid_v[i, j] /= grid_m[i, j] + grid_v[i, j][1] -= dt * gravity + if i < bound and grid_v[i, j].x < 0: + grid_v[i, j].x = 0 + if i > n_grid - bound and grid_v[i, j].x > 0: + grid_v[i, j].x = 0 + if j < bound and grid_v[i, j].y < 0: + grid_v[i, j].y = 0 + if j > n_grid - bound and grid_v[i, j].y > 0: + grid_v[i, j].y = 0 + for p in x: + Xp = x[p] / dx + base = int(Xp - 0.5) + fx = Xp - base + w = [0.5 * (1.5 - fx)**2, 0.75 - (fx - 1)**2, 0.5 * (fx - 0.5)**2] + new_v = ti.Vector.zero(ti.f32, 2) + new_C = ti.Matrix.zero(ti.f32, 2, 2) + for i, j in ti.static(ti.ndrange(3, 3)): + offset = ti.Vector([i, j]) + dpos = (offset - fx) * dx + weight = w[i].x * w[j].y + g_v = grid_v[base + offset] + new_v += weight * g_v + new_C += 4 * weight * g_v.outer_product(dpos) / dx**2 + v[p] = new_v + x[p] += dt * v[p] + J[p] *= 1 + dt * new_C.trace() + C[p] = new_C + +@ti.kernel +def init(): + for i in range(n_particles): + x[i] = [ti.random() * 0.4 + 0.2, ti.random() * 0.4 + 0.2] + v[i] = [0, -1] + J[i] = 1 + +init() +gui = ti.GUI('MPM88') +while gui.running and not gui.get_event(gui.ESCAPE): + for s in range(50): + substep() + gui.clear(0x112F41) + gui.circles(x.to_numpy(), radius=1.5, color=0x068587) + gui.show() diff --git a/misc/action_simple_uv.py b/misc/action_simple_uv.py new file mode 100644 index 0000000000000..7379447e79d5e --- /dev/null +++ b/misc/action_simple_uv.py @@ -0,0 +1,15 @@ +import taichi as ti +ti.core.start_recording('record.yml') +ti.init(arch=ti.cc) + +n = 512 +x = ti.Vector(3, ti.f32, (n, n)) + + +@ti.kernel +def render(): + for i, j in x: + x[i, j] = [i / x.shape[0], j / x.shape[1], 0] + + +render() diff --git a/misc/cc_action_record.py b/misc/cc_action_record.py deleted file mode 100644 index e437a90f505af..0000000000000 --- a/misc/cc_action_record.py +++ /dev/null @@ -1,36 +0,0 @@ -import taichi as ti -ti.core.start_recording('record.yml') -ti.init(arch=ti.cc) - -n = 512 -x = ti.Vector(3, ti.f32, (n, n)) - - -@ti.kernel -def render(): - for i, j in x: - x[i, j] = [i / x.shape[0], j / x.shape[1], 0] - - -@ti.kernel -def dump_ppm(tensor: ti.template()): - if ti.static(isinstance(tensor, ti.Matrix)): - print('P3') - else: - print('P2') - print(tensor.shape[0], tensor.shape[1], 255) - for _ in range(1): - for i in range(x.shape[0]): - for j in range(x.shape[1]): - c = min(255, - max(0, int(tensor[j, x.shape[1] - 1 - i] * 255 + 0.5))) - if ti.static(isinstance(tensor, ti.Matrix)): - r, g, b = c - print(r, g, b) - else: - print(c) - - -render() -dump_ppm(x) -ti.imshow(x) From b7d1c662e2534d6dabe960d4d6e04e1258935311 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sat, 1 Aug 2020 22:53:36 +0800 Subject: [PATCH 03/15] tmp save --- misc/cc_action.py | 64 +++++++++++++++++++++++++++++++++++++++++++++ taichi/util/str.cpp | 1 - 2 files changed, 64 insertions(+), 1 deletion(-) create mode 100644 misc/cc_action.py diff --git a/misc/cc_action.py b/misc/cc_action.py new file mode 100644 index 0000000000000..ed3cb73d9d065 --- /dev/null +++ b/misc/cc_action.py @@ -0,0 +1,64 @@ +import yaml +import warnings + + +class Composer: + def __init__(self, entries): + self.entries = entries + self.launches = [] + + def run(self): + for e in self.entries: + action = e['action'] + func = getattr(self, 'do_' + action, self.do_unknown) + func(e) + + def do_unknown(self, e): + print(f"Unknown action type: {e['action']}") + + def do_compile_runtime(self, e): + header = e['runtime_header'] + source = e['runtime_source'] + + print(header) + print(source) + + def do_compile_layout(self, e): + source = e['layout_source'] + + print(source) + print('struct Ti_S0root Ti_root;') + print('') + + def do_allocate_buffer(self, e): + root_size = e['root_size'] + gtmp_size = e['gtmp_size'] + + print(f'Ti_i8 Ti_gtmp[{gtmp_size}];') + print('union Ti_BitCast Ti_args[8];') + print('Ti_i32 Ti_earg[8 * 8];') + print('') + print('struct Ti_Context Ti_ctx = {') + print(' &Ti_root, Ti_gtmp, Ti_args, Ti_earg,') + print('};') + print('') + + def do_compile_kernel(self, e): + name = e['kernel_name'] + source = e['kernel_source'] + + print(source) + print('') + + def do_launch_kernel(self, e): + self.launches.append(e) + + + +if __name__ == '__main__': + with open('record.yml') as fin: + warnings.filterwarnings('ignore') + obj = yaml.load(fin) + + comp = Composer(obj) + comp.run() diff --git a/taichi/util/str.cpp b/taichi/util/str.cpp index 1caa7963f6c6a..d2d9366415d8a 100644 --- a/taichi/util/str.cpp +++ b/taichi/util/str.cpp @@ -17,7 +17,6 @@ std::string c_quoted(std::string const &str) { REG_ESC('\n', "n"); REG_ESC('\a', "a"); REG_ESC('\b', "b"); - REG_ESC('\?', "?"); REG_ESC('\v', "v"); REG_ESC('\t', "t"); REG_ESC('\f', "f"); From 68555156b37607833b4a22999c7c8cecedad1547 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 00:44:55 +0800 Subject: [PATCH 04/15] emscripten --- misc/action_simple_uv.py | 1 + misc/cc_action.py | 61 ++++++++++++++++++++++++++++++++++-- taichi/backends/cc/context.h | 1 + 3 files changed, 61 insertions(+), 2 deletions(-) diff --git a/misc/action_simple_uv.py b/misc/action_simple_uv.py index 7379447e79d5e..eff86cd79d6e1 100644 --- a/misc/action_simple_uv.py +++ b/misc/action_simple_uv.py @@ -13,3 +13,4 @@ def render(): render() +x.to_numpy() diff --git a/misc/cc_action.py b/misc/cc_action.py index ed3cb73d9d065..11c6a0ef315d3 100644 --- a/misc/cc_action.py +++ b/misc/cc_action.py @@ -3,16 +3,54 @@ class Composer: - def __init__(self, entries): + def __init__(self, entries, emscripten=False): self.entries = entries + self.emscripten = emscripten self.launches = [] def run(self): + if self.emscripten: + print('#include ') + for e in self.entries: action = e['action'] func = getattr(self, 'do_' + action, self.do_unknown) func(e) + if self.emscripten: + print('') + print('EMSCRIPTEN_KEEPALIVE') + print('void Ti_set_arg_i32(struct Ti_Context *ti_ctx, ') + print(' Ti_i32 index, Ti_i32 value) {') + print(' ti_ctx->args[index].val_i32 = value;') + print('}') + print('') + print('EMSCRIPTEN_KEEPALIVE') + print('void Ti_set_arg_i64(struct Ti_Context *ti_ctx, ') + print(' Ti_i32 index, Ti_i64 value) {') + print(' ti_ctx->args[index].val_i64 = value;') + print('}') + print('') + print('EMSCRIPTEN_KEEPALIVE') + print('void Ti_set_arg_f32(struct Ti_Context *ti_ctx, ') + print(' Ti_i32 index, Ti_f32 value) {') + print(' ti_ctx->args[index].val_f32 = value;') + print('}') + print('') + print('EMSCRIPTEN_KEEPALIVE') + print('void Ti_set_arg_f64(struct Ti_Context *ti_ctx, ') + print(' Ti_i32 index, Ti_f64 value) {') + print(' ti_ctx->args[index].val_f64 = value;') + print('}') + print('') + print('EMSCRIPTEN_KEEPALIVE') + print('void Ti_set_arg_ptr(struct Ti_Context *ti_ctx, ') + print(' Ti_i32 index, void *ptr) {') + print(' ti_ctx->args[index].ptr_void = ptr;') + print('}') + print('') + print('') + def do_unknown(self, e): print(f"Unknown action type: {e['action']}") @@ -22,22 +60,39 @@ def do_compile_runtime(self, e): print(header) print(source) + print('') def do_compile_layout(self, e): source = e['layout_source'] print(source) + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') print('struct Ti_S0root Ti_root;') print('') def do_allocate_buffer(self, e): root_size = e['root_size'] gtmp_size = e['gtmp_size'] + extr_size = 4 * 1024 * 1024 # pinpoint: 4 MB + + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') + print(f'Ti_i8 Ti_extr[{extr_size}];') print(f'Ti_i8 Ti_gtmp[{gtmp_size}];') + + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') + print('union Ti_BitCast Ti_args[8];') + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') print('Ti_i32 Ti_earg[8 * 8];') print('') + + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') print('struct Ti_Context Ti_ctx = {') print(' &Ti_root, Ti_gtmp, Ti_args, Ti_earg,') print('};') @@ -47,6 +102,8 @@ def do_compile_kernel(self, e): name = e['kernel_name'] source = e['kernel_source'] + if self.emscripten: + print('EMSCRIPTEN_KEEPALIVE') print(source) print('') @@ -60,5 +117,5 @@ def do_launch_kernel(self, e): warnings.filterwarnings('ignore') obj = yaml.load(fin) - comp = Composer(obj) + comp = Composer(obj, emscripten=True) comp.run() diff --git a/taichi/backends/cc/context.h b/taichi/backends/cc/context.h index c6a68bf32acc5..464802d0e700a 100644 --- a/taichi/backends/cc/context.h +++ b/taichi/backends/cc/context.h @@ -23,6 +23,7 @@ union Ti_BitCast { Ti_u8 *ptr_u8; Ti_f32 *ptr_f32; Ti_f64 *ptr_f64; + void *ptr_void; }; struct Ti_Context { From b49d00ef744129f093fc51968285fa8052a6c4a8 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 06:36:44 +0800 Subject: [PATCH 05/15] TI_ACTION_RECORD=record.yml env var --- misc/{cc_action.py => cc_compose.py} | 0 python/taichi/core/util.py | 4 ++++ 2 files changed, 4 insertions(+) rename misc/{cc_action.py => cc_compose.py} (100%) diff --git a/misc/cc_action.py b/misc/cc_compose.py similarity index 100% rename from misc/cc_action.py rename to misc/cc_compose.py diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 7ac07b9b190c1..0849d0bf18ec2 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -423,6 +423,10 @@ def at_startup(): ti_core.set_core_state_python_imported(True) + record_file = os.environ.get('TI_ACTION_RECORD') + if record_file: + ti_core.start_recording(record_file) + def start_memory_monitoring(output_fn, pid=-1, interval=1): # removing dependency on psutil From 0f55f59690b0587eb1f405f6578d8e03adc65c8d Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 06:38:42 +0800 Subject: [PATCH 06/15] clean --- misc/action_mpm88.py | 89 ---------------------------------------- misc/action_simple_uv.py | 16 -------- 2 files changed, 105 deletions(-) delete mode 100644 misc/action_mpm88.py delete mode 100644 misc/action_simple_uv.py diff --git a/misc/action_mpm88.py b/misc/action_mpm88.py deleted file mode 100644 index b448683a8e238..0000000000000 --- a/misc/action_mpm88.py +++ /dev/null @@ -1,89 +0,0 @@ -# MPM-MLS in 88 lines of Taichi code, originally created by @yuanming-hu -import taichi as ti -ti.core.start_recording('record.yml') -ti.init(arch=ti.cc) - -n_particles = 8192 -n_grid = 128 -dx = 1 / n_grid -dt = 2e-4 - -p_rho = 1 -p_vol = (dx * 0.5)**2 -p_mass = p_vol * p_rho -gravity = 9.8 -bound = 3 -E = 400 - -x = ti.Vector.field(2, ti.f32, n_particles) -v = ti.Vector.field(2, ti.f32, n_particles) -C = ti.Matrix.field(2, 2, ti.f32, n_particles) -J = ti.field(ti.f32, n_particles) - -grid_v = ti.Vector.field(2, ti.f32, (n_grid, n_grid)) -grid_m = ti.field(ti.f32, (n_grid, n_grid)) - -@ti.kernel -def substep(): - for i, j in grid_m: - grid_v[i, j] = [0, 0] - grid_m[i, j] = 0 - for p in x: - Xp = x[p] / dx - base = int(Xp - 0.5) - fx = Xp - base - w = [0.5 * (1.5 - fx)**2, 0.75 - (fx - 1)**2, 0.5 * (fx - 0.5)**2] - stress = -dt * 4 * E * p_vol * (J[p] - 1) / dx**2 - affine = ti.Matrix([[stress, 0], [0, stress]]) + p_mass * C[p] - for i, j in ti.static(ti.ndrange(3, 3)): - offset = ti.Vector([i, j]) - dpos = (offset - fx) * dx - weight = w[i].x * w[j].y - grid_v[base + offset] += weight * (p_mass * v[p] + affine @ dpos) - grid_m[base + offset] += weight * p_mass - for i, j in grid_m: - if grid_m[i, j] > 0: - grid_v[i, j] /= grid_m[i, j] - grid_v[i, j][1] -= dt * gravity - if i < bound and grid_v[i, j].x < 0: - grid_v[i, j].x = 0 - if i > n_grid - bound and grid_v[i, j].x > 0: - grid_v[i, j].x = 0 - if j < bound and grid_v[i, j].y < 0: - grid_v[i, j].y = 0 - if j > n_grid - bound and grid_v[i, j].y > 0: - grid_v[i, j].y = 0 - for p in x: - Xp = x[p] / dx - base = int(Xp - 0.5) - fx = Xp - base - w = [0.5 * (1.5 - fx)**2, 0.75 - (fx - 1)**2, 0.5 * (fx - 0.5)**2] - new_v = ti.Vector.zero(ti.f32, 2) - new_C = ti.Matrix.zero(ti.f32, 2, 2) - for i, j in ti.static(ti.ndrange(3, 3)): - offset = ti.Vector([i, j]) - dpos = (offset - fx) * dx - weight = w[i].x * w[j].y - g_v = grid_v[base + offset] - new_v += weight * g_v - new_C += 4 * weight * g_v.outer_product(dpos) / dx**2 - v[p] = new_v - x[p] += dt * v[p] - J[p] *= 1 + dt * new_C.trace() - C[p] = new_C - -@ti.kernel -def init(): - for i in range(n_particles): - x[i] = [ti.random() * 0.4 + 0.2, ti.random() * 0.4 + 0.2] - v[i] = [0, -1] - J[i] = 1 - -init() -gui = ti.GUI('MPM88') -while gui.running and not gui.get_event(gui.ESCAPE): - for s in range(50): - substep() - gui.clear(0x112F41) - gui.circles(x.to_numpy(), radius=1.5, color=0x068587) - gui.show() diff --git a/misc/action_simple_uv.py b/misc/action_simple_uv.py deleted file mode 100644 index eff86cd79d6e1..0000000000000 --- a/misc/action_simple_uv.py +++ /dev/null @@ -1,16 +0,0 @@ -import taichi as ti -ti.core.start_recording('record.yml') -ti.init(arch=ti.cc) - -n = 512 -x = ti.Vector(3, ti.f32, (n, n)) - - -@ti.kernel -def render(): - for i, j in x: - x[i, j] = [i / x.shape[0], j / x.shape[1], 0] - - -render() -x.to_numpy() From 0ee7da0c799148b51906dd683036df145c6fd6b7 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 07:14:28 +0800 Subject: [PATCH 07/15] cli --- misc/cc_compose.py | 121 ------------------------- python/taichi/cc_compose.py | 130 +++++++++++++++++++++++++++ python/taichi/core/util.py | 11 ++- python/taichi/main.py | 26 ++++++ taichi/backends/cc/cc_configuation.h | 2 +- 5 files changed, 164 insertions(+), 126 deletions(-) delete mode 100644 misc/cc_compose.py create mode 100644 python/taichi/cc_compose.py diff --git a/misc/cc_compose.py b/misc/cc_compose.py deleted file mode 100644 index 11c6a0ef315d3..0000000000000 --- a/misc/cc_compose.py +++ /dev/null @@ -1,121 +0,0 @@ -import yaml -import warnings - - -class Composer: - def __init__(self, entries, emscripten=False): - self.entries = entries - self.emscripten = emscripten - self.launches = [] - - def run(self): - if self.emscripten: - print('#include ') - - for e in self.entries: - action = e['action'] - func = getattr(self, 'do_' + action, self.do_unknown) - func(e) - - if self.emscripten: - print('') - print('EMSCRIPTEN_KEEPALIVE') - print('void Ti_set_arg_i32(struct Ti_Context *ti_ctx, ') - print(' Ti_i32 index, Ti_i32 value) {') - print(' ti_ctx->args[index].val_i32 = value;') - print('}') - print('') - print('EMSCRIPTEN_KEEPALIVE') - print('void Ti_set_arg_i64(struct Ti_Context *ti_ctx, ') - print(' Ti_i32 index, Ti_i64 value) {') - print(' ti_ctx->args[index].val_i64 = value;') - print('}') - print('') - print('EMSCRIPTEN_KEEPALIVE') - print('void Ti_set_arg_f32(struct Ti_Context *ti_ctx, ') - print(' Ti_i32 index, Ti_f32 value) {') - print(' ti_ctx->args[index].val_f32 = value;') - print('}') - print('') - print('EMSCRIPTEN_KEEPALIVE') - print('void Ti_set_arg_f64(struct Ti_Context *ti_ctx, ') - print(' Ti_i32 index, Ti_f64 value) {') - print(' ti_ctx->args[index].val_f64 = value;') - print('}') - print('') - print('EMSCRIPTEN_KEEPALIVE') - print('void Ti_set_arg_ptr(struct Ti_Context *ti_ctx, ') - print(' Ti_i32 index, void *ptr) {') - print(' ti_ctx->args[index].ptr_void = ptr;') - print('}') - print('') - print('') - - def do_unknown(self, e): - print(f"Unknown action type: {e['action']}") - - def do_compile_runtime(self, e): - header = e['runtime_header'] - source = e['runtime_source'] - - print(header) - print(source) - print('') - - def do_compile_layout(self, e): - source = e['layout_source'] - - print(source) - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - print('struct Ti_S0root Ti_root;') - print('') - - def do_allocate_buffer(self, e): - root_size = e['root_size'] - gtmp_size = e['gtmp_size'] - extr_size = 4 * 1024 * 1024 # pinpoint: 4 MB - - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - print(f'Ti_i8 Ti_extr[{extr_size}];') - - print(f'Ti_i8 Ti_gtmp[{gtmp_size}];') - - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - - print('union Ti_BitCast Ti_args[8];') - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - print('Ti_i32 Ti_earg[8 * 8];') - print('') - - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - print('struct Ti_Context Ti_ctx = {') - print(' &Ti_root, Ti_gtmp, Ti_args, Ti_earg,') - print('};') - print('') - - def do_compile_kernel(self, e): - name = e['kernel_name'] - source = e['kernel_source'] - - if self.emscripten: - print('EMSCRIPTEN_KEEPALIVE') - print(source) - print('') - - def do_launch_kernel(self, e): - self.launches.append(e) - - - -if __name__ == '__main__': - with open('record.yml') as fin: - warnings.filterwarnings('ignore') - obj = yaml.load(fin) - - comp = Composer(obj, emscripten=True) - comp.run() diff --git a/python/taichi/cc_compose.py b/python/taichi/cc_compose.py new file mode 100644 index 0000000000000..b8e4b321ac5b6 --- /dev/null +++ b/python/taichi/cc_compose.py @@ -0,0 +1,130 @@ +import yaml +import warnings +import sys + + +class Composer: + def __init__(self, fout, entries, emscripten=False): + self.fout = fout + self.entries = entries + self.emscripten = emscripten + self.launches = [] + + def emit(self, line): + print(line, out=self.fout) + + def run(self): + if self.emscripten: + self.emit('#include ') + + for e in self.entries: + action = e['action'] + func = getattr(self, 'do_' + action, self.do_unknown) + func(e) + + if self.emscripten: + self.emit('') + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('void Ti_set_arg_i32(struct Ti_Context *ti_ctx, ') + self.emit(' Ti_i32 index, Ti_i32 value) {') + self.emit(' ti_ctx->args[index].val_i32 = value;') + self.emit('}') + self.emit('') + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('void Ti_set_arg_i64(struct Ti_Context *ti_ctx, ') + self.emit(' Ti_i32 index, Ti_i64 value) {') + self.emit(' ti_ctx->args[index].val_i64 = value;') + self.emit('}') + self.emit('') + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('void Ti_set_arg_f32(struct Ti_Context *ti_ctx, ') + self.emit(' Ti_i32 index, Ti_f32 value) {') + self.emit(' ti_ctx->args[index].val_f32 = value;') + self.emit('}') + self.emit('') + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('void Ti_set_arg_f64(struct Ti_Context *ti_ctx, ') + self.emit(' Ti_i32 index, Ti_f64 value) {') + self.emit(' ti_ctx->args[index].val_f64 = value;') + self.emit('}') + self.emit('') + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('void Ti_set_arg_ptr(struct Ti_Context *ti_ctx, ') + self.emit(' Ti_i32 index, void *ptr) {') + self.emit(' ti_ctx->args[index].ptr_void = ptr;') + self.emit('}') + self.emit('') + self.emit('') + + def do_unknown(self, e): + self.emit(f"Unknown action type: {e['action']}") + + def do_compile_runtime(self, e): + header = e['runtime_header'] + source = e['runtime_source'] + + self.emit(header) + self.emit(source) + self.emit('') + + def do_compile_layout(self, e): + source = e['layout_source'] + + self.emit(source) + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('struct Ti_S0root Ti_root;') + self.emit('') + + def do_allocate_buffer(self, e): + root_size = e['root_size'] + gtmp_size = e['gtmp_size'] + extr_size = 4 * 1024 * 1024 # pinpoint: 4 MB + + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit(f'Ti_i8 Ti_extr[{extr_size}];') + + self.emit(f'Ti_i8 Ti_gtmp[{gtmp_size}];') + + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + + self.emit('union Ti_BitCast Ti_args[8];') + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('Ti_i32 Ti_earg[8 * 8];') + self.emit('') + + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit('struct Ti_Context Ti_ctx = {') + self.emit(' &Ti_root, Ti_gtmp, Ti_args, Ti_earg,') + self.emit('};') + self.emit('') + + def do_compile_kernel(self, e): + name = e['kernel_name'] + source = e['kernel_source'] + + if self.emscripten: + self.emit('EMSCRIPTEN_KEEPALIVE') + self.emit(source) + self.emit('') + + def do_launch_kernel(self, e): + self.launches.append(e) + + +def main(fin_name, fout_name, emscripten=False): + with open(fin_name) as fin: + warnings.filterwarnings('ignore') + obj = yaml.load(fin) + + with open(fout_name) as fout: + comp = Composer(fout, obj, emscripten) + comp.run() + + +if __name__ == '__main__': + main(sys.argv[1], sys.argv[2], len(sys.argv) > 3) diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 0849d0bf18ec2..b811f231be21f 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -251,17 +251,19 @@ def check_exists(src): ) -def prepare_sandbox(src): +def prepare_sandbox(src=None): global g_tmp_dir - check_exists(src) + if src is not None: + check_exists(src) import atexit import shutil from tempfile import mkdtemp tmp_dir = mkdtemp(prefix='taichi-') atexit.register(shutil.rmtree, tmp_dir) print(f'[Taichi] preparing sandbox at {tmp_dir}') - dest = os.path.join(tmp_dir, 'taichi_core.so') - shutil.copy(src, dest) + if src is not None: + dest = os.path.join(tmp_dir, 'taichi_core.so') + shutil.copy(src, dest) os.mkdir(os.path.join(tmp_dir, 'runtime/')) return tmp_dir @@ -285,6 +287,7 @@ def get_unique_task_id(): import_ti_core() if get_os_name() != 'win': dll = ctypes.CDLL(get_core_shared_object(), mode=ctypes.RTLD_LOCAL) + ti_core.set_tmp_dir(prepare_sandbox()) ti_core.set_python_package_dir(package_root()) os.makedirs(ti_core.get_repo_dir(), exist_ok=True) diff --git a/python/taichi/main.py b/python/taichi/main.py index 6e633b0561157..44f810b44ae78 100644 --- a/python/taichi/main.py +++ b/python/taichi/main.py @@ -1003,6 +1003,32 @@ def dist(self, arguments: list = sys.argv[2:]): sys.argv.append(args.mode) runpy.run_path('build.py') + @register + def cc_compose(self, arguments: list = sys.argv[2:]): + """Compose C backend action record into a complete C file""" + parser = argparse.ArgumentParser(prog='ti cc_compose', + description=f"{self.cc_compose.__doc__}") + parser.add_argument( + 'fin_name', + required=True, + help='Action record YAML file name from C backend, e.g. program.yml') + parser.add_argument( + 'fout_name', + required=True, + help='The output C source file name, e.g. program.c') + parser.add_argument( + '-e', + '--emscripten', + required=False, + default=False, + dest='emscripten', + action='store_true', + help='Generate output C file for Emscripten instead of raw C') + args = parser.parse_args(arguments) + + from .cc_compose import main + main(args.fin_name, args.fout_name, args.emscripten) + def main(): cli = TaichiMain() diff --git a/taichi/backends/cc/cc_configuation.h b/taichi/backends/cc/cc_configuation.h index 1c63803e0172f..1fa5a4dc0c9cc 100644 --- a/taichi/backends/cc/cc_configuation.h +++ b/taichi/backends/cc/cc_configuation.h @@ -9,7 +9,7 @@ struct CCConfiguation { std::string compile_cmd, link_cmd; CCConfiguation() - : compile_cmd("gcc -Wc99-c11-compat -c -o '{}' '{}'"), + : compile_cmd("gcc -Wc99-c11-compat -c -o '{}' '{}' -O3"), link_cmd("gcc -shared -fPIC -o '{}' '{}'") { } }; From 0eeca935c39ad01ff81195e48e4933b1ecf00f22 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 07:55:04 +0800 Subject: [PATCH 08/15] fix cli --- python/taichi/cc_compose.py | 6 +++--- python/taichi/main.py | 2 -- 2 files changed, 3 insertions(+), 5 deletions(-) diff --git a/python/taichi/cc_compose.py b/python/taichi/cc_compose.py index b8e4b321ac5b6..b6bc7cf3b6589 100644 --- a/python/taichi/cc_compose.py +++ b/python/taichi/cc_compose.py @@ -11,7 +11,7 @@ def __init__(self, fout, entries, emscripten=False): self.launches = [] def emit(self, line): - print(line, out=self.fout) + print(line, file=self.fout) def run(self): if self.emscripten: @@ -117,11 +117,11 @@ def do_launch_kernel(self, e): def main(fin_name, fout_name, emscripten=False): - with open(fin_name) as fin: + with open(fin_name, 'r') as fin: warnings.filterwarnings('ignore') obj = yaml.load(fin) - with open(fout_name) as fout: + with open(fout_name, 'w') as fout: comp = Composer(fout, obj, emscripten) comp.run() diff --git a/python/taichi/main.py b/python/taichi/main.py index 44f810b44ae78..b8aa3fac7d798 100644 --- a/python/taichi/main.py +++ b/python/taichi/main.py @@ -1010,11 +1010,9 @@ def cc_compose(self, arguments: list = sys.argv[2:]): description=f"{self.cc_compose.__doc__}") parser.add_argument( 'fin_name', - required=True, help='Action record YAML file name from C backend, e.g. program.yml') parser.add_argument( 'fout_name', - required=True, help='The output C source file name, e.g. program.c') parser.add_argument( '-e', From 230b22e13bfaf8bbe7a8c306bf3e099dc18f4157 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 12:39:10 +0800 Subject: [PATCH 09/15] [skip ci] cleaner --- python/taichi/cc_compose.py | 34 ---------------------------------- 1 file changed, 34 deletions(-) diff --git a/python/taichi/cc_compose.py b/python/taichi/cc_compose.py index b6bc7cf3b6589..d147e54cc999e 100644 --- a/python/taichi/cc_compose.py +++ b/python/taichi/cc_compose.py @@ -22,40 +22,6 @@ def run(self): func = getattr(self, 'do_' + action, self.do_unknown) func(e) - if self.emscripten: - self.emit('') - self.emit('EMSCRIPTEN_KEEPALIVE') - self.emit('void Ti_set_arg_i32(struct Ti_Context *ti_ctx, ') - self.emit(' Ti_i32 index, Ti_i32 value) {') - self.emit(' ti_ctx->args[index].val_i32 = value;') - self.emit('}') - self.emit('') - self.emit('EMSCRIPTEN_KEEPALIVE') - self.emit('void Ti_set_arg_i64(struct Ti_Context *ti_ctx, ') - self.emit(' Ti_i32 index, Ti_i64 value) {') - self.emit(' ti_ctx->args[index].val_i64 = value;') - self.emit('}') - self.emit('') - self.emit('EMSCRIPTEN_KEEPALIVE') - self.emit('void Ti_set_arg_f32(struct Ti_Context *ti_ctx, ') - self.emit(' Ti_i32 index, Ti_f32 value) {') - self.emit(' ti_ctx->args[index].val_f32 = value;') - self.emit('}') - self.emit('') - self.emit('EMSCRIPTEN_KEEPALIVE') - self.emit('void Ti_set_arg_f64(struct Ti_Context *ti_ctx, ') - self.emit(' Ti_i32 index, Ti_f64 value) {') - self.emit(' ti_ctx->args[index].val_f64 = value;') - self.emit('}') - self.emit('') - self.emit('EMSCRIPTEN_KEEPALIVE') - self.emit('void Ti_set_arg_ptr(struct Ti_Context *ti_ctx, ') - self.emit(' Ti_i32 index, void *ptr) {') - self.emit(' ti_ctx->args[index].ptr_void = ptr;') - self.emit('}') - self.emit('') - self.emit('') - def do_unknown(self, e): self.emit(f"Unknown action type: {e['action']}") From 2a4fceb96e924e6791d90a9cdcf1deada7ec6eae Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Sun, 2 Aug 2020 13:59:18 +0800 Subject: [PATCH 10/15] prepare_sandbox now take no arg for simplicity --- python/taichi/core/util.py | 22 +++++++++++++--------- python/taichi/testing.py | 2 -- 2 files changed, 13 insertions(+), 11 deletions(-) diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index b811f231be21f..7c9904c714ca2 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -251,19 +251,17 @@ def check_exists(src): ) -def prepare_sandbox(src=None): - global g_tmp_dir - if src is not None: - check_exists(src) +def prepare_sandbox(): + ''' + return a runtime temporary directory name, which will be automatically + deleted on exit. + ''' import atexit import shutil from tempfile import mkdtemp tmp_dir = mkdtemp(prefix='taichi-') atexit.register(shutil.rmtree, tmp_dir) print(f'[Taichi] preparing sandbox at {tmp_dir}') - if src is not None: - dest = os.path.join(tmp_dir, 'taichi_core.so') - shutil.copy(src, dest) os.mkdir(os.path.join(tmp_dir, 'runtime/')) return tmp_dir @@ -287,6 +285,7 @@ def get_unique_task_id(): import_ti_core() if get_os_name() != 'win': dll = ctypes.CDLL(get_core_shared_object(), mode=ctypes.RTLD_LOCAL) + # TODO: add tmp dir for win too, so that C backend can be functional: ti_core.set_tmp_dir(prepare_sandbox()) ti_core.set_python_package_dir(package_root()) @@ -298,12 +297,15 @@ def get_unique_task_id(): os.environ['DYLD_FALLBACK_LIBRARY_PATH'] = get_runtime_directory() lib_path = os.path.join(bin_dir, 'libtaichi_core.dylib') tmp_cwd = os.getcwd() - tmp_dir = prepare_sandbox(lib_path) + tmp_dir = prepare_sandbox() + check_exists(lib_path) + shutil.copy(lib_path, os.path.join(tmp_dir, 'taichi_core.so')) os.chdir(tmp_dir) sys.path.append(tmp_dir) import taichi_core as ti_core os.chdir(tmp_cwd) + # TODO: unify importing infrastructure: elif get_os_name() == 'linux': bin_dir = get_bin_directory() if 'LD_LIBRARY_PATH' in os.environ: @@ -313,7 +315,9 @@ def get_unique_task_id(): lib_path = os.path.join(bin_dir, 'libtaichi_core.so') check_exists(lib_path) tmp_cwd = os.getcwd() - tmp_dir = prepare_sandbox(lib_path) + tmp_dir = prepare_sandbox() + check_exists(lib_path) + shutil.copy(lib_path, os.path.join(tmp_dir, 'taichi_core.so')) os.chdir(tmp_dir) sys.path.append(tmp_dir) try: diff --git a/python/taichi/testing.py b/python/taichi/testing.py index 1f961b7dbeeeb..8c322cf4a1a68 100644 --- a/python/taichi/testing.py +++ b/python/taichi/testing.py @@ -1,7 +1,5 @@ import taichi as ti -print('[Taichi] loading test module') - ## Helper functions def approx(expected, **kwargs): From 71bc7b530c5b5341d55cc74341c6f2a75cf21899 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=BD=AD=E4=BA=8E=E6=96=8C?= <1931127624@qq.com> Date: Mon, 3 Aug 2020 00:42:45 +0800 Subject: [PATCH 11/15] [skip ci] Update python/taichi/core/util.py Co-authored-by: Yuanming Hu --- python/taichi/core/util.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 7c9904c714ca2..2660c210c1f93 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -253,8 +253,8 @@ def check_exists(src): def prepare_sandbox(): ''' - return a runtime temporary directory name, which will be automatically - deleted on exit. + Returns a temporary directory, which will be automatically deleted on exit. + It may contain the taichi_core shared object or some misc. files. ''' import atexit import shutil From 1b325454ffcc4507ae7b21debc64d14cfd0d6f21 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Mon, 3 Aug 2020 00:44:06 +0800 Subject: [PATCH 12/15] revert kernel.cpp since yuanming's insist --- taichi/program/kernel.cpp | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) diff --git a/taichi/program/kernel.cpp b/taichi/program/kernel.cpp index 808fcfbf32181..e1bca3e410b1c 100644 --- a/taichi/program/kernel.cpp +++ b/taichi/program/kernel.cpp @@ -119,8 +119,12 @@ void Kernel::set_arg_float(int i, float64 d) { TI_ASSERT_INFO( !args[i].is_nparray, "Assigning a scalar value to a numpy array argument is not allowed"); - auto dt = args[i].dt; + ActionRecorder::get_instance().record( + "set_kernel_arg_float64", {ActionArg("kernel_name", name), + ActionArg("arg_id", i), ActionArg("val", d)}); + + auto dt = args[i].dt; if (dt == DataType::f32) { program.context.set_arg(i, (float32)d); } else if (dt == DataType::f64) { @@ -150,8 +154,12 @@ void Kernel::set_arg_int(int i, int64 d) { TI_ASSERT_INFO( !args[i].is_nparray, "Assigning scalar value to numpy array argument is not allowed"); - auto dt = args[i].dt; + ActionRecorder::get_instance().record( + "set_kernel_arg_int64", {ActionArg("kernel_name", name), + ActionArg("arg_id", i), ActionArg("val", d)}); + + auto dt = args[i].dt; if (dt == DataType::i32) { program.context.set_arg(i, (int32)d); } else if (dt == DataType::i64) { @@ -239,6 +247,12 @@ void Kernel::set_arg_nparray(int i, uint64 ptr, uint64 size) { TI_ASSERT_INFO(args[i].is_nparray, "Assigning numpy array to scalar argument is not allowed"); + ActionRecorder::get_instance().record( + "set_kernel_arg_ext_ptr", + {ActionArg("kernel_name", name), ActionArg("arg_id", i), + ActionArg("address", fmt::format("0x{:x}", ptr)), + ActionArg("array_size_in_bytes", (int64)size)}); + args[i].size = size; program.context.set_arg(i, ptr); } From 05f7dfa5c18e404de110bd18ca90f0a966a0f2c8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=BD=AD=E4=BA=8E=E6=96=8C?= <1931127624@qq.com> Date: Mon, 3 Aug 2020 12:32:00 +0800 Subject: [PATCH 13/15] Apply suggestions from code review --- python/taichi/core/util.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 2660c210c1f93..0d21dd9f97bc1 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -285,8 +285,8 @@ def get_unique_task_id(): import_ti_core() if get_os_name() != 'win': dll = ctypes.CDLL(get_core_shared_object(), mode=ctypes.RTLD_LOCAL) - # TODO: add tmp dir for win too, so that C backend can be functional: - ti_core.set_tmp_dir(prepare_sandbox()) + # C backend need a tmp dir for containing the generated .c and compiled .so files: + ti_core.set_tmp_dir(prepare_sandbox()) # TODO: always allocate a tmp_dir for all situations ti_core.set_python_package_dir(package_root()) os.makedirs(ti_core.get_repo_dir(), exist_ok=True) From 914b2fba16b0de810770d5130acda40e814ff552 Mon Sep 17 00:00:00 2001 From: archibate <1931127624@qq.com> Date: Tue, 4 Aug 2020 16:47:49 +0800 Subject: [PATCH 14/15] rev part --- python/taichi/testing.py | 1 + taichi/backends/opengl/shaders/random.glsl.h | 12 ++++++------ 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/python/taichi/testing.py b/python/taichi/testing.py index 8c322cf4a1a68..8e42f95830958 100644 --- a/python/taichi/testing.py +++ b/python/taichi/testing.py @@ -1,5 +1,6 @@ import taichi as ti +print('[Taichi] loading test module') ## Helper functions def approx(expected, **kwargs): diff --git a/taichi/backends/opengl/shaders/random.glsl.h b/taichi/backends/opengl/shaders/random.glsl.h index 98319bf424c81..ec9cc3872de3c 100644 --- a/taichi/backends/opengl/shaders/random.glsl.h +++ b/taichi/backends/opengl/shaders/random.glsl.h @@ -5,11 +5,11 @@ STR( uvec4 _rand_; void _init_rand() { - uint i = (54321u + gl_GlobalInvocationID.x) * (12345u + uint(_rand_state_)); - _rand_.x = 123456789u * i * 1000000007u; - _rand_.y = 362436069u; - _rand_.z = 521288629u; - _rand_.w = 88675123u; + uint i = (54321 + gl_GlobalInvocationID.x) * (12345 + _rand_state_); + _rand_.x = 123456789 * i * 1000000007; + _rand_.y = 362436069; + _rand_.z = 521288629; + _rand_.w = 88675123; // Yes, this is not an atomic operation, but just fine since no matter // how `_rand_state_` changes, `gl_GlobalInvocationID.x` can still help @@ -22,7 +22,7 @@ uint _rand_u32() { uint t = _rand_.x ^ (_rand_.x << 11); _rand_.xyz = _rand_.yzw; _rand_.w = (_rand_.w ^ (_rand_.w >> 19)) ^ (t ^ (t >> 8)); - return _rand_.w * 1000000007u; + return _rand_.w * 1000000007; } float _rand_f32() { From 5ec8dc6caf2ebb2bafff3deeed817d970a48330c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E5=BD=AD=E4=BA=8E=E6=96=8C?= <1931127624@qq.com> Date: Tue, 4 Aug 2020 23:15:04 +0800 Subject: [PATCH 15/15] [skip ci] Update python/taichi/core/util.py Co-authored-by: Yuanming Hu --- python/taichi/core/util.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/taichi/core/util.py b/python/taichi/core/util.py index 0d21dd9f97bc1..e3468eabc30a3 100644 --- a/python/taichi/core/util.py +++ b/python/taichi/core/util.py @@ -285,7 +285,7 @@ def get_unique_task_id(): import_ti_core() if get_os_name() != 'win': dll = ctypes.CDLL(get_core_shared_object(), mode=ctypes.RTLD_LOCAL) - # C backend need a tmp dir for containing the generated .c and compiled .so files: + # The C backend needs a temporary directory for the generated .c and compiled .so files: ti_core.set_tmp_dir(prepare_sandbox()) # TODO: always allocate a tmp_dir for all situations ti_core.set_python_package_dir(package_root())