Skip to content

Instantly share code, notes, and snippets.

View dadeba's full-sized avatar

N.Nakasato dadeba

7 followers · 0 following
  • University of Aizu
  • Japan
View GitHub Profile
@dadeba
dadeba / bnn_pynq.py
Created August 28, 2018 14:23
A python implementation for the PYNQ-BNN MNIST example
import numpy as np
from udmabuf import Udmabuf
from uio import Uio
import os
import struct
import math
# set path to the BNN-PYQN repo.
BNN_PYNQ_ROOT = '.'
PARAM = 'bnn/params/mnist'
@dadeba
dadeba / gist:2f8b4b49a72c1a2762f975811cbba449
Created August 28, 2018 14:28
diff for udmabuf.py (https://github.com/ikwzm/FPGA-SoC-Linux-Example-1-ZYBO-Z7)
+ def config_sync(self, size, direc):
+ with open(self.class_path + '/sync_size', mode='w') as f:
+ f.write(str(size))
+ with open(self.class_path + '/sync_direction', mode='w') as f:
+ f.write(str(direc))
+
+ def sync_for_cpu(self):
+ with open(self.class_path + '/sync_for_cpu', mode='w') as f:
+ f.write(str(1))
+
@dadeba
dadeba / pipeline_top.vhd
Created March 28, 2018 20:31
VHDL FP32
-------------------------------------------------------------------------------
-------------------------------------------------------------------------------
library ieee;
use ieee.std_logic_1164.all;
use ieee.std_logic_unsigned.all;
entity pipeline_top is
port (
xi : in std_logic_vector(31 downto 0);
yi : in std_logic_vector(31 downto 0);
@dadeba
dadeba / g5.q
Last active March 22, 2018 21:13
pipeline description
/VARI xi, yi, zi, e2;
/VARJ xj, yj, zj, mj;
/VARF ax, ay, az, ap;
dx = xj - xi;
dy = yj - yi;
dz = zj - zi;
r2 = dx*dx + dy*dy + dz*dz + e2;
r1i = powm12(r2);
r2i = r1i*r1i;
r1im = mj*r1i;
@dadeba
dadeba / sph1.q
Created March 20, 2018 19:49
input
/VARI xi, yi, zi, vxi, vyi, vzi, hi;
/VARJ xj, yj, zj, vxj, vyj, vzj, hj, mj;
/VARF rho, rox, roy, roz, dd;
dx = xi - xj;
dy = yi - yj;
dz = zi - zj;
r2 = dx*dx + dy*dy + dz*dz;
r1 = sqrt(r2);
@dadeba
dadeba / HKeys.hpp
Created March 30, 2012 08:50
Convert the Morton key to the Hilbert key in 3-D
// A B C D E F G H I J K M
// 0 1 2 3 4 5 6 7 8 9 10 11
struct HKeys {
static const uint64 T0[12][8];
static const uint64 T1[12][8];
uint64 poi(uint64 key, uint64 level)
{
return (key >> 3*level)&0x7ULL;
}
@dadeba
dadeba / gravity_v4v2.asm
Created March 23, 2012 13:01
ASM file for gravity_v4v2.cl (https://gist.github.com/2167470)
.file "/tmp/5d41e25b-e85e-4f10-836a-5b23eab3f6a7.TMP"
.text
.globl _Z12native_rsqrtDv8_f
.align 16, 0x90
.type _Z12native_rsqrtDv8_f,@function
_Z12native_rsqrtDv8_f: # @_Z12native_rsqrtDv8_f
# BB#0:
vrsqrtps YMM0, YMM0
ret
.Ltmp0:
@dadeba
dadeba / html.rb
Created March 23, 2012 06:21
get
def self.get(url)
return self.new(Open3.popen3("curl", "--location", "--compressed", url) {|i,o,e| o.read })
end
@dadeba
dadeba / gravity_v4v2.cl
Created March 23, 2012 05:59
OpenCL: A vectorized kernel for gravity interaction inspired by Tanikawa etal. 2012
#define READONLY_P const * restrict
float4 sum(float8 x)
{
float4 tmp;
tmp.x = x.s0 + x.s4;
tmp.y = x.s1 + x.s5;
tmp.z = x.s2 + x.s6;
tmp.w = x.s3 + x.s7;
return tmp;
@dadeba
dadeba / gravity_v4.cl
Created March 23, 2012 02:43
OpenCL: A vectorized kernel for gravity interaction
#define READONLY_P const * restrict
__kernel
void
grav1(
__global float4 READONLY_P x,
__global float4 READONLY_P y,
__global float4 READONLY_P z,
__global float4 READONLY_P m,
__global float4 *ax,