Merge pull request #1 from mumax/master

updating this branch to latest mumax/3

README.md

@@ -60,3 +60,8 @@ To do all at once on Ubuntu:
 sudo apt-get install git golang-go gcc nvidia-cuda-toolkit nvidia-cuda-dev nvidia-340 gnuplot
 export GOPATH=$HOME go get -u -v github.com/mumax/3/cmd/mumax3
 ```
+
+Contributing
+------------
+
+Contributions are gratefully accepted. To contribute code, fork our repo on github and send a pull request.

bench/gpus.txt

@@ -8,3 +8,5 @@
 4.194304e+06 1.3222269984079185e+08 1.6945213680693907e-14  "GTX 980"  
 4.194304e+06 1.3064462210481596e+08 1.6945210251645377e-14  "K20XM"
 4.194304e+06 1.7967269114941204e+08 1.6945210251645377e-14  "GTX TITAN"
+4.194304e+06 1.8968187897582138e+08 1.6945207552204512e-14  "GTX 1080"
+4.194304e+06 1.9961744689897743e+08 1.6945207552204512e-14	"GTX 980 Ti"

cmd/mumax3/main.go

@@ -6,77 +6,44 @@ import (
 	"fmt"
 	"github.com/mumax/3/cuda"
 	"github.com/mumax/3/engine"
-	"github.com/mumax/3/prof"
 	"github.com/mumax/3/script"
-	"github.com/mumax/3/timer"
 	"github.com/mumax/3/util"
 	"log"
 	"os"
-	"runtime"
+	"os/exec"
+	"path"
 	"time"
 )
 
 var (
-	flag_cachedir      = flag.String("cache", "", "Kernel cache directory")
-	flag_cpuprof       = flag.Bool("cpuprof", false, "Record gopprof CPU profile")
-	flag_failfast      = flag.Bool("failfast", false, "If one simulation fails, stop entire batch immediately")
-	flag_forceclean    = flag.Bool("f", true, "Force start, clean existing output directory")
-	flag_gpu           = flag.Int("gpu", 0, "Specify GPU")
-	flag_interactive   = flag.Bool("i", false, "Open interactive browser session")
-	flag_launchtimeout = flag.Duration("launchtimeout", 0, "Launch timeout for CUDA calls")
-	flag_memprof       = flag.Bool("memprof", false, "Recored gopprof memory profile")
-	flag_od            = flag.String("o", "", "Override output directory")
-	flag_port          = flag.String("http", ":35367", "Port to serve web gui")
-	flag_selftest      = flag.Bool("paranoid", false, "Enable convolution self-test for cuFFT sanity.")
-	flag_silent        = flag.Bool("s", false, "Silent") // provided for backwards compatibility
-	flag_sync          = flag.Bool("sync", false, "Synchronize all CUDA calls (debug)")
-	flag_test          = flag.Bool("test", false, "Cuda test (internal)")
-	flag_version       = flag.Bool("v", true, "Print version")
-	flag_vet           = flag.Bool("vet", false, "Check input files for errors, but don't run them")
+	flag_failfast = flag.Bool("failfast", false, "If one simulation fails, stop entire batch immediately")
+	flag_test     = flag.Bool("test", false, "Cuda test (internal)")
+	flag_version  = flag.Bool("v", true, "Print version")
+	flag_vet      = flag.Bool("vet", false, "Check input files for errors, but don't run them")
+	// more flags in engine/gofiles.go
 )
 
 func main() {
 	flag.Parse()
 	log.SetPrefix("")
 	log.SetFlags(0)
 
-	cuda.Init(*flag_gpu)
-	runtime.GOMAXPROCS(runtime.NumCPU())
-	cuda.Synchronous = *flag_sync
+	cuda.Init(*engine.Flag_gpu)
+
+	cuda.Synchronous = *engine.Flag_sync
 	if *flag_version {
 		printVersion()
 	}
 
-	timer.Timeout = *flag_launchtimeout
-	if *flag_launchtimeout != 0 {
-		cuda.Synchronous = true
-	}
-
-	engine.TestDemag = *flag_selftest
-
 	// used by bootstrap launcher to test cuda
 	// successful exit means cuda was initialized fine
 	if *flag_test {
 		fmt.Println(cuda.GPUInfo)
 		os.Exit(0)
 	}
 
-	engine.CacheDir = *flag_cachedir
-	if *flag_cpuprof {
-		prof.InitCPU(".")
-	}
-	if *flag_memprof {
-		prof.InitMem(".")
-	}
-	defer prof.Cleanup()
 	defer engine.Close() // flushes pending output, if any
 
-	defer func() {
-		if *flag_sync {
-			timer.Print(os.Stdout)
-		}
-	}()
-
 	if *flag_vet {
 		vet()
 		return
@@ -99,7 +66,7 @@ func runInteractive() {
 	// setup outut dir
 	now := time.Now()
 	outdir := fmt.Sprintf("mumax-%v-%02d-%02d_%02dh%02d.out", now.Year(), int(now.Month()), now.Day(), now.Hour(), now.Minute())
-	engine.InitIO(outdir, outdir, *flag_forceclean)
+	engine.InitIO(outdir, outdir, *engine.Flag_forceclean)
 
 	engine.Timeout = 365 * 24 * time.Hour // basically forever
 
@@ -115,14 +82,21 @@ func runInteractive() {
 	engine.RunInteractive()
 }
 
-// Runs a script file.
 func runFileAndServe(fname string) {
+	if path.Ext(fname) == ".go" {
+		runGoFile(fname)
+	} else {
+		runScript(fname)
+	}
+}
 
+func runScript(fname string) {
 	outDir := util.NoExt(fname) + ".out"
-	if *flag_od != "" {
-		outDir = *flag_od
+	if *engine.Flag_od != "" {
+		outDir = *engine.Flag_od
 	}
-	engine.InitIO(fname, outDir, *flag_forceclean)
+	engine.InitIO(fname, outDir, *engine.Flag_forceclean)
+
 	fname = engine.InputFile
 
 	var code *script.BlockStmt
@@ -136,35 +110,50 @@ func runFileAndServe(fname string) {
 	// now the parser is not used anymore so it can handle web requests
 	goServeGUI()
 
-	if *flag_interactive {
-		openbrowser("http://127.0.0.1" + *flag_port)
+	if *engine.Flag_interactive {
+		openbrowser("http://127.0.0.1" + *engine.Flag_port)
 	}
 
 	// start executing the tree, possibly injecting commands from web gui
 	engine.EvalFile(code)
 
-	if *flag_interactive {
+	if *engine.Flag_interactive {
 		engine.RunInteractive()
 	}
 }
 
-//func runRemote(fname string) {
-//	URL, err := url.Parse(fname)
-//	util.FatalErr(err)
-//	host := URL.Host
-//	engine.MountHTTPFS("http://" + host)
-//	od := util.NoExt(URL.Path) + ".out"
-//	engine.InitIO(od, *flag_force)
-//	runFileAndServe(URL.Path) // TODO proxyserve?
-//}
+func runGoFile(fname string) {
+
+	// pass through flags
+	flags := []string{"run", fname}
+	flag.Visit(func(f *flag.Flag) {
+		if f.Name != "o" {
+			flags = append(flags, fmt.Sprintf("-%v=%v", f.Name, f.Value))
+		}
+	})
+
+	if *engine.Flag_od != "" {
+		flags = append(flags, fmt.Sprintf("-o=%v", *engine.Flag_od))
+	}
+
+	cmd := exec.Command("go", flags...)
+	log.Println("go", flags)
+	cmd.Stdout = os.Stdout
+	cmd.Stdin = os.Stdin
+	cmd.Stderr = os.Stderr
+	err := cmd.Run()
+	if err != nil {
+		os.Exit(1)
+	}
+}
 
 // start Gui server and return server address
 func goServeGUI() string {
-	if *flag_port == "" {
+	if *engine.Flag_port == "" {
 		log.Println(`//not starting GUI (-http="")`)
 		return ""
 	}
-	addr := engine.GoServe(*flag_port)
+	addr := engine.GoServe(*engine.Flag_port)
 	fmt.Print("//starting GUI at http://127.0.0.1", addr, "\n")
 	return addr
 }

cmd/mumax3/queue.go

@@ -24,7 +24,7 @@ var (
 func RunQueue(files []string) {
 	s := NewStateTab(files)
 	s.PrintTo(os.Stdout)
-	go s.ListenAndServe(*flag_port)
+	go s.ListenAndServe(*engine.Flag_port)
 	s.Run()
 	fmt.Println(numOK.get(), "OK, ", numFailed.get(), "failed")
 	os.Exit(int(exitStatus))
@@ -115,7 +115,7 @@ func run(inFile string, gpu int, webAddr string) {
 	// pass through flags
 	flags := []string{gpuFlag, httpFlag}
 	flag.Visit(func(f *flag.Flag) {
-		if f.Name != "gpu" && f.Name != "http" {
+		if f.Name != "gpu" && f.Name != "http" && f.Name != "failfast" {
 			flags = append(flags, fmt.Sprintf("-%v=%v", f.Name, f.Value))
 		}
 	})

cuda/amul.h

@@ -0,0 +1,29 @@
+#ifndef _AMUL_H_
+#define _AMUL_H_
+
+#include "float3.h"
+
+// Returns mul * arr[i], or mul when arr == NULL;
+inline __device__ float amul(float *arr, float mul, int i) {
+    return (arr == NULL)? (mul): (mul * arr[i]);
+}
+
+// Returns m * a[i], or m when a == NULL;
+inline __device__ float3 vmul(float *ax, float *ay, float *az,
+                              float mx, float my, float mz, int i) {
+    return make_float3(amul(ax, mx, i),
+                       amul(ay, my, i),
+                       amul(az, mz, i));
+}
+
+// Returns 1/Msat, or 0 when Msat == 0.
+inline __device__ float inv_Msat(float *Ms_, float Ms_mul, int i) {
+    float ms = amul(Ms_, Ms_mul, i);
+    if (ms == 0.0f) {
+        return 0.0f;
+    } else {
+        return 1.0f / ms;
+    }
+}
+
+#endif

cuda/anisotropy.go

@@ -1,40 +1,51 @@
 package cuda
 
 import (
-	"unsafe"
-
 	"github.com/mumax/3/data"
 	"github.com/mumax/3/util"
 )
 
-// Adds cubic anisotropy field to Beff.
-// see cubicanisotropy.cu
-func AddCubicAnisotropy(Beff, m *data.Slice, k1_red, k2_red, k3_red LUTPtr, c1, c2 LUTPtrs, regions *Bytes) {
+// Add uniaxial magnetocrystalline anisotropy field to Beff.
+// see uniaxialanisotropy.cu
+func AddCubicAnisotropy2(Beff, m *data.Slice, Msat, k1, k2, k3, c1, c2 MSlice) {
 	util.Argument(Beff.Size() == m.Size())
 
 	N := Beff.Len()
 	cfg := make1DConf(N)
-
-	k_addcubicanisotropy_async(
+	k_addcubicanisotropy2_async(
 		Beff.DevPtr(X), Beff.DevPtr(Y), Beff.DevPtr(Z),
 		m.DevPtr(X), m.DevPtr(Y), m.DevPtr(Z),
-		unsafe.Pointer(k1_red), unsafe.Pointer(k2_red), unsafe.Pointer(k3_red),
-		c1[X], c1[Y], c1[Z],
-		c2[X], c2[Y], c2[Z],
-		regions.Ptr, N, cfg)
+		Msat.DevPtr(0), Msat.Mul(0),
+		k1.DevPtr(0), k1.Mul(0),
+		k2.DevPtr(0), k2.Mul(0),
+		k3.DevPtr(0), k3.Mul(0),
+		c1.DevPtr(X), c1.Mul(X),
+		c1.DevPtr(Y), c1.Mul(Y),
+		c1.DevPtr(Z), c1.Mul(Z),
+		c2.DevPtr(X), c2.Mul(X),
+		c2.DevPtr(Y), c2.Mul(Y),
+		c2.DevPtr(Z), c2.Mul(Z),
+		N, cfg)
 }
 
 // Add uniaxial magnetocrystalline anisotropy field to Beff.
 // see uniaxialanisotropy.cu
-func AddUniaxialAnisotropy(Beff, m *data.Slice, k1_red, k2_red LUTPtr, u LUTPtrs, regions *Bytes) {
+func AddUniaxialAnisotropy2(Beff, m *data.Slice, Msat, k1, k2, u MSlice) {
 	util.Argument(Beff.Size() == m.Size())
 
+	checkSize(Beff, m, k1, k2, u, Msat)
+
 	N := Beff.Len()
 	cfg := make1DConf(N)
 
-	k_adduniaxialanisotropy_async(Beff.DevPtr(X), Beff.DevPtr(Y), Beff.DevPtr(Z),
+	k_adduniaxialanisotropy2_async(
+		Beff.DevPtr(X), Beff.DevPtr(Y), Beff.DevPtr(Z),
 		m.DevPtr(X), m.DevPtr(Y), m.DevPtr(Z),
-		unsafe.Pointer(k1_red), unsafe.Pointer(k2_red),
-		u[X], u[Y], u[Z],
-		regions.Ptr, N, cfg)
+		Msat.DevPtr(0), Msat.Mul(0),
+		k1.DevPtr(0), k1.Mul(0),
+		k2.DevPtr(0), k2.Mul(0),
+		u.DevPtr(X), u.Mul(X),
+		u.DevPtr(Y), u.Mul(Y),
+		u.DevPtr(Z), u.Mul(Z),
+		N, cfg)
 }

cuda/buffer.go

@@ -61,6 +61,9 @@ func Recycle(s *data.Slice) {
 	// put each component buffer back on the stack
 	for i := 0; i < s.NComp(); i++ {
 		ptr := s.DevPtr(i)
+		if ptr == unsafe.Pointer(uintptr(0)) {
+			continue
+		}
 		if _, ok := buf_check[ptr]; !ok {
 			log.Panic("recyle: was not obtained with getbuffer")
 		}

cuda/conv_copypad.go

@@ -1,8 +1,6 @@
 package cuda
 
 import (
-	"unsafe"
-
 	"github.com/mumax/3/data"
 	"github.com/mumax/3/util"
 )
@@ -22,13 +20,13 @@ func copyUnPad(dst, src *data.Slice, dstsize, srcsize [3]int) {
 // Copies src into dst, which is larger, and multiplies by vol*Bsat.
 // The remainder of dst is not filled with zeros.
 // Used to zero-pad magnetization before convolution and in the meanwhile multiply m by its length.
-func copyPadMul(dst, src, vol *data.Slice, dstsize, srcsize [3]int, Bsat LUTPtr, regions *Bytes) {
+func copyPadMul(dst, src, vol *data.Slice, dstsize, srcsize [3]int, Msat MSlice) {
 	util.Argument(dst.NComp() == 1 && src.NComp() == 1)
 	util.Assert(dst.Len() == prod(dstsize) && src.Len() == prod(srcsize))
 
 	cfg := make3DConf(srcsize)
 
-	k_copypadmul_async(dst.DevPtr(0), dstsize[X], dstsize[Y], dstsize[Z],
-		src.DevPtr(0), vol.DevPtr(0), srcsize[X], srcsize[Y], srcsize[Z],
-		unsafe.Pointer(Bsat), regions.Ptr, cfg)
+	k_copypadmul2_async(dst.DevPtr(0), dstsize[X], dstsize[Y], dstsize[Z],
+		src.DevPtr(0), srcsize[X], srcsize[Y], srcsize[Z],
+		Msat.DevPtr(0), Msat.Mul(0), vol.DevPtr(0), cfg)
 }