diff --git a/2d/total/blur.go b/2d/total/blur.go new file mode 100644 index 0000000..2d9510e --- /dev/null +++ b/2d/total/blur.go @@ -0,0 +1,34 @@ +package total + +import ( + "fmt" + "image" + "os" + + "github.com/disintegration/imaging" +) + +// Blur blurs the image using a gaussian function. Sigma parameter must be +// positive and indicates how much the image will be blurred. It runs on the +// Intel GPU when available and falls back to the CPU implementation +// (disintegration/imaging's Blur) otherwise. +// +// Blur 使用高斯函数模糊图像。Sigma 参数必须为正数，表示模糊程度。 +// 当 Intel GPU 可用时在 GPU 上运行，否则回退到 CPU 实现 +// （disintegration/imaging 的 Blur）。 +func Blur(img image.Image, sigma float64) *image.NRGBA { + if sigma <= 0 { + return imaging.Clone(img) + } + + if canUseBlurKernel { + dst, err := gpuBlur(img, sigma) + if err == nil { + return dst + } + fmt.Fprintln(os.Stderr, "[gg.blur_ocl] gpuBlur err:", err, "fallback to cpu") + canUseBlurKernel = false + } + + return imaging.Blur(img, sigma) +} diff --git a/2d/total/blur_ocl.go b/2d/total/blur_ocl.go new file mode 100644 index 0000000..5067217 --- /dev/null +++ b/2d/total/blur_ocl.go @@ -0,0 +1,308 @@ +package total + +import ( + _ "embed" + "errors" + "fmt" + "image" + "image/color" + "math" + "os" + "unsafe" + + "github.com/fumiama/gozel/gozel" + "github.com/fumiama/gozel/ze" + + "github.com/FloatTech/gg/gpu" +) + +//go:generate ocloc compile -file build/blur_ocl.cl -spv_only -options "-cl-mad-enable -cl-fast-relaxed-math -cl-finite-math-only -cl-single-precision-constant" -internal_options "-O3" -output build/blur_ocl +//go:generate llvm-spirv -to-text build/blur_ocl_.spv -o build/blur_ocl.spt + +//go:embed build/blur_ocl_.spv +var blurspv []byte + +var ( + canUseBlurKernel = false + blurModel ze.ModuleHandle +) + +// maxBlurRadius caps the gaussian window half-size to avoid excessive device loops. +// +// maxBlurRadius 限制高斯窗口半径上限，避免设备端循环过长。 +const maxBlurRadius = 64 + +func init() { + if !gpu.IsAvailable() { + return + } + + var err error + blurModel, err = gpu.ModuleCreateAndCheckKernels(blurspv, "blurh", "blurv") + if err != nil { + fmt.Fprintln(os.Stderr, "[gg.blur_ocl] gpu init err:", err) + return + } + + canUseBlurKernel = true +} + +// gaussianBlurWeights computes the normalized gaussian weights of window size +// 2*radius+1, where radius = ceil(3*sigma) capped at maxBlurRadius, mirroring +// disintegration/imaging's Blur. +func gaussianBlurWeights(sigma float64) (radius int, weights []float32) { + radius = int(math.Ceil(sigma * 3.0)) + if radius < 1 { + radius = 1 + } + if radius > maxBlurRadius { + radius = maxBlurRadius + } + weights = make([]float32, 2*radius+1) + sum := 0.0 + for i := range weights { + x := float64(i - radius) + w := math.Exp(-(x * x) / (2 * sigma * sigma)) + weights[i] = float32(w) + sum += w + } + inv := float32(1 / sum) + for i := range weights { + weights[i] *= inv + } + return radius, weights +} + +func gpuBlur(img image.Image, sigma float64) (dst *image.NRGBA, err error) { + if sigma <= 0 { + return nil, errors.New("sigma must be positive") + } + rgbaimg := ImageToRGBA(img) + w, h := rgbaimg.Bounds().Dx(), rgbaimg.Bounds().Dy() + if w <= 0 || h <= 0 { + return nil, errors.New("empty image") + } + npix := uintptr(w) * uintptr(h) + if npix*16 > 1<<31 { + return nil, errors.New("image too large to blur on gpu") + } + pixels := rgbaimg.Pix + + radius, weights := gaussianBlurWeights(sigma) + + krnH, err := blurModel.KernelCreate("blurh") + if err != nil { + return nil, err + } + defer krnH.Destroy() + krnV, err := blurModel.KernelCreate("blurv") + if err != nil { + return nil, err + } + defer krnV.Destroy() + + // Allocate input and output image buffers (host + device) + imgSize := npix * unsafe.Sizeof(color.RGBA{}) + inHost, inDev, err := gpu.MemAllocHostDevicePair(imgSize, unsafe.Sizeof(color.RGBA{})) + if err != nil { + return nil, err + } + defer gpu.MemFree(inHost) + defer gpu.MemFree(inDev) + + outHost, outDev, err := gpu.MemAllocHostDevicePair(imgSize, unsafe.Sizeof(color.RGBA{})) + if err != nil { + return nil, err + } + defer gpu.MemFree(outHost) + defer gpu.MemFree(outDev) + + // Allocate intermediate premultiplied float4 buffer (device only) + tmpDev, err := gpu.MemAllocDevice(npix*unsafe.Sizeof([4]float32{}), unsafe.Sizeof([4]float32{})) + if err != nil { + return nil, err + } + defer gpu.MemFree(tmpDev) + + // Allocate weights buffer (host + device) + wtsSize := uintptr(len(weights)) * unsafe.Sizeof(float32(0)) + wtsHost, wtsDev, err := gpu.MemAllocHostDevicePair(wtsSize, unsafe.Sizeof(float32(0))) + if err != nil { + return nil, err + } + defer gpu.MemFree(wtsHost) + defer gpu.MemFree(wtsDev) + + copy(unsafe.Slice((*uint8)(inHost), len(pixels)), pixels) + copy(unsafe.Slice((*float32)(wtsHost), len(weights)), weights) + + // Set horizontal pass arguments + err = krnH.SetArgumentValue(0, &inDev) + if err != nil { + return nil, err + } + err = krnH.SetArgumentValue(1, &tmpDev) + if err != nil { + return nil, err + } + err = krnH.SetArgumentValue(2, uint32(w)) + if err != nil { + return nil, err + } + err = krnH.SetArgumentValue(3, uint32(h)) + if err != nil { + return nil, err + } + err = krnH.SetArgumentValue(4, int32(radius)) + if err != nil { + return nil, err + } + err = krnH.SetArgumentValue(5, &wtsDev) + if err != nil { + return nil, err + } + + // Set vertical pass arguments + err = krnV.SetArgumentValue(0, &tmpDev) + if err != nil { + return nil, err + } + err = krnV.SetArgumentValue(1, &outDev) + if err != nil { + return nil, err + } + err = krnV.SetArgumentValue(2, uint32(w)) + if err != nil { + return nil, err + } + err = krnV.SetArgumentValue(3, uint32(h)) + if err != nil { + return nil, err + } + err = krnV.SetArgumentValue(4, int32(radius)) + if err != nil { + return nil, err + } + err = krnV.SetArgumentValue(5, &wtsDev) + if err != nil { + return nil, err + } + + // Determine group size + gX, gY, _, err := krnH.SuggestGroupSize(uint32(w), uint32(h), 1) + if err != nil { + return nil, err + } + err = krnH.SetGroupSize(gX, gY, 1) + if err != nil { + return nil, err + } + err = krnV.SetGroupSize(gX, gY, 1) + if err != nil { + return nil, err + } + gcx := uint32(math.Ceil(float64(w) / float64(gX))) + gcy := uint32(math.Ceil(float64(h) / float64(gY))) + + // Build command list + lst, err := gpu.CommandListCreate() + if err != nil { + return nil, err + } + defer lst.Destroy() + + // Event: input pixels copy done + inEv, cl1, err := gpu.EventCreate(gozel.ZE_EVENT_SCOPE_FLAG_HOST, 0) + if err != nil { + return nil, err + } + defer cl1() + defer inEv.Destroy() + + // Event: weights copy done + wtEv, cl2, err := gpu.EventCreate(gozel.ZE_EVENT_SCOPE_FLAG_HOST, 0) + if err != nil { + return nil, err + } + defer cl2() + defer wtEv.Destroy() + + // Copy input pixels: host -> device + err = lst.AppendMemoryCopy(inDev, inHost, imgSize, inEv) + if err != nil { + return nil, err + } + + // Copy weights: host -> device + err = lst.AppendMemoryCopy(wtsDev, wtsHost, wtsSize, wtEv) + if err != nil { + return nil, err + } + + // Event: horizontal pass done + hEv, cl3, err := gpu.EventCreate(gozel.ZE_EVENT_SCOPE_FLAG_HOST, 0) + if err != nil { + return nil, err + } + defer cl3() + defer hEv.Destroy() + + // Launch horizontal pass + err = lst.AppendLaunchKernel(krnH, &gozel.ZeGroupCount{ + Groupcountx: gcx, Groupcounty: gcy, Groupcountz: 1, + }, hEv, inEv, wtEv) + if err != nil { + return nil, err + } + + // Event: vertical pass done + vEv, cl4, err := gpu.EventCreate(gozel.ZE_EVENT_SCOPE_FLAG_HOST, 0) + if err != nil { + return nil, err + } + defer cl4() + defer vEv.Destroy() + + // Launch vertical pass + err = lst.AppendLaunchKernel(krnV, &gozel.ZeGroupCount{ + Groupcountx: gcx, Groupcounty: gcy, Groupcountz: 1, + }, vEv, hEv) + if err != nil { + return nil, err + } + + // Event: output copy done + outEv, cl5, err := gpu.EventCreate(gozel.ZE_EVENT_SCOPE_FLAG_HOST, 0) + if err != nil { + return nil, err + } + defer cl5() + defer outEv.Destroy() + + // Copy output: device -> host (wait for vertical pass) + err = lst.AppendMemoryCopy(outHost, outDev, imgSize, outEv, vEv) + if err != nil { + return nil, err + } + + // Close and execute + err = lst.Close() + if err != nil { + return nil, err + } + + err = gpu.ExecCommandLists(lst) + if err != nil { + return nil, err + } + + // Wait for output copy to complete + err = outEv.HostSynchronize(math.MaxUint64) + if err != nil { + return nil, err + } + + dst = image.NewNRGBA(image.Rect(0, 0, w, h)) + copy(dst.Pix, unsafe.Slice((*uint8)(outHost), len(dst.Pix))) + return dst, nil +} diff --git a/2d/total/blur_test.go b/2d/total/blur_test.go new file mode 100644 index 0000000..ae497bc --- /dev/null +++ b/2d/total/blur_test.go @@ -0,0 +1,206 @@ +package total + +import ( + "image" + "image/color" + "math" + "math/rand" + "testing" + + "github.com/disintegration/imaging" +) + +func TestGaussianBlurWeights(t *testing.T) { + for _, sigma := range []float64{0.3, 0.5, 1, 2.3, 9, 25, 100} { + radius, weights := gaussianBlurWeights(sigma) + if len(weights) != 2*radius+1 { + t.Fatalf("sigma %v: len(weights) = %d, want 2*%d+1", sigma, len(weights), radius) + } + sum := float32(0) + for _, w := range weights { + sum += w + } + if math.Abs(float64(sum-1)) > 1e-5 { + t.Fatalf("sigma %v: weights sum = %v, want 1", sigma, sum) + } + for i, w := range weights { + if math.Abs(float64(weights[len(weights)-1-i]-w)) > 1e-6 { + t.Fatalf("sigma %v: weights not symmetric at index %d", sigma, i) + } + } + } +} + +func randomTestImage(w, h int, seed int64) *image.RGBA { + img := image.NewRGBA(image.Rect(0, 0, w, h)) + rnd := rand.New(rand.NewSource(seed)) + for y := 0; y < h; y++ { + for x := 0; x < w; x++ { + // 生成带大片低 alpha 和透明区域的像素，考验预乘与边缘处理 + a := uint8(0) + switch rnd.Intn(4) { + case 0: + a = 0 + case 1: + a = uint8(rnd.Intn(64)) + case 2: + a = uint8(64 + rnd.Intn(128)) + default: + a = 255 + } + // 预乘 + r8, g8, b8 := uint8(rnd.Intn(256)), uint8(rnd.Intn(256)), uint8(rnd.Intn(256)) + img.SetRGBA(x, y, color.RGBA{ + R: uint8(uint16(r8) * uint16(a) / 255), + G: uint8(uint16(g8) * uint16(a) / 255), + B: uint8(uint16(b8) * uint16(a) / 255), + A: a, + }) + } + } + return img +} + +// simulateBlur 在 CPU 上模拟 blur_ocl.cl 两个 kernel 的数学流程， +// 用于验证算法与 imaging.Blur 语义一致。 +func simulateBlur(img image.Image, sigma float64) *image.NRGBA { + rgbaimg := ImageToRGBA(img) + w, h := rgbaimg.Bounds().Dx(), rgbaimg.Bounds().Dy() + radius, weights := gaussianBlurWeights(sigma) + + px := func(x, y int) [4]float32 { + o := (y*w + x) * 4 + return [4]float32{ + float32(rgbaimg.Pix[o]), + float32(rgbaimg.Pix[o+1]), + float32(rgbaimg.Pix[o+2]), + float32(rgbaimg.Pix[o+3]), + } + } + + // pass 1: blurh —— 预乘域加权平均，按入界权重和归一 + tmp := make([][4]float32, w*h) + for y := 0; y < h; y++ { + for x := 0; x < w; x++ { + var acc [4]float32 + wsum := float32(0) + lo, hi := max(-radius, -x), min(radius, w-1-x) + for i := lo; i <= hi; i++ { + wt := weights[i+radius] + p := px(x+i, y) + for c := 0; c < 4; c++ { + acc[c] += p[c] * wt + } + wsum += wt + } + if wsum < 1e-12 { + wsum = 1e-12 + } + for c := 0; c < 4; c++ { + acc[c] /= wsum + } + tmp[y*w+x] = acc + } + } + + // pass 2: blurv —— 预乘域加权平均，alpha 按权重和归一，rgb 反预乘 + dst := image.NewNRGBA(image.Rect(0, 0, w, h)) + for y := 0; y < h; y++ { + for x := 0; x < w; x++ { + var acc [4]float32 + wsum := float32(0) + lo, hi := max(-radius, -y), min(radius, h-1-y) + for i := lo; i <= hi; i++ { + wt := weights[i+radius] + p := tmp[(y+i)*w+x] + for c := 0; c < 4; c++ { + acc[c] += p[c] * wt + } + wsum += wt + } + if wsum < 1e-12 { + wsum = 1e-12 + } + for c := 0; c < 4; c++ { + acc[c] /= wsum + } + o := (y*w + x) * 4 + if acc[3] > 0 { + k := 255 / acc[3] + for c := 0; c < 3; c++ { + dst.Pix[o+c] = sat8(acc[c] * k) + } + dst.Pix[o+3] = sat8(acc[3]) + } + } + } + return dst +} + +func sat8(v float32) uint8 { + if v <= 0 { + return 0 + } + if v >= 255 { + return 255 + } + return uint8(math.Round(float64(v))) +} + +func compareWithImaging(t *testing.T, label string, got *image.NRGBA, want *image.NRGBA, sigma float64) { + t.Helper() + b := got.Bounds() + var sum, maxv float64 + n := b.Dx() * b.Dy() + for y := 0; y < b.Dy(); y++ { + for x := 0; x < b.Dx(); x++ { + i := (y*b.Dx() + x) * 4 + for c := 0; c < 4; c++ { + d := math.Abs(float64(got.Pix[i+c]) - float64(want.Pix[i+c])) + sum += d + if d > maxv { + maxv = d + } + } + } + } + mean := sum / float64(n*4) + t.Logf("[%s sigma=%.2f] mean abs diff = %.3f, max abs diff = %.0f", label, sigma, mean, maxv) + if mean > 2.0 { + t.Fatalf("[%s sigma=%.2f] mean abs diff %.3f > 2.0", label, sigma, mean) + } +} + +func TestBlurAlgorithmMatchesImaging(t *testing.T) { + for _, sigma := range []float64{0.8, 2.3, 9} { + img := randomTestImage(97, 61, int64(sigma*10)) + want := imaging.Blur(img, sigma) + got := simulateBlur(img, sigma) + compareWithImaging(t, "simulate", got, want, sigma) + } +} + +func TestBlurGPUMatchesImaging(t *testing.T) { + if !canUseBlurKernel { + t.Skip("no available Intel GPU on this machine") + } + for _, sigma := range []float64{0.8, 2.3, 9} { + img := randomTestImage(97, 61, int64(sigma*10)) + want := imaging.Blur(img, sigma) + got, err := gpuBlur(img, sigma) + if err != nil { + t.Fatalf("sigma %v: gpuBlur err: %v", sigma, err) + } + compareWithImaging(t, "gpu", got, want, sigma) + } +} + +func TestBlur(t *testing.T) { + img := randomTestImage(33, 21, 42) + for _, sigma := range []float64{-1, 0, 0.5, 5} { + dst := Blur(img, sigma) + if dst.Bounds().Dx() != 33 || dst.Bounds().Dy() != 21 { + t.Fatalf("sigma %v: bad bounds %v", sigma, dst.Bounds()) + } + } +} diff --git a/2d/total/build/blur_ocl.cl b/2d/total/build/blur_ocl.cl new file mode 100644 index 0000000..396cc26 --- /dev/null +++ b/2d/total/build/blur_ocl.cl @@ -0,0 +1,70 @@ +// Separable gaussian blur, mirroring disintegration/imaging's Blur math: +// - pass 1 (blurh) walks horizontally and stores premultiplied float4 (rgb premul, alpha) +// renormalized by the sum of in-bounds weights; +// - pass 2 (blurv) walks vertically over premultiplied values, renormalizes alpha by the +// sum of in-bounds weights and un-premultiplies rgb, writing straight RGBA8. +// Both passes are edge-aware: out-of-bounds taps are skipped, not clamped. + +kernel void blurh( + global const uchar* src, // premultiplied RGBA8, w * h * 4 bytes + global float* dst, // float4 per pixel, w * h * 4 floats + uint w, uint h, + int radius, + global const float* wts) // 2 * radius + 1 normalized gaussian weights +{ + uint x = get_global_id(0); + uint y = get_global_id(1); + if (x >= w || y >= h) { + return; + } + + int lo = max(-radius, -(int)x); + int hi = min(radius, (int)w - 1 - (int)x); + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + float wsum = 0.0f; + for (int i = lo; i <= hi; i++) { + float wt = wts[i + radius]; + uchar4 p = vload4(0, src + ((ulong)y * w + (uint)((int)x + i)) * 4); + acc += convert_float4(p) * wt; + wsum += wt; + } + acc /= max(wsum, 1e-12f); + vstore4(acc, 0, dst + ((ulong)y * w + x) * 4); +} + +kernel void blurv( + global const float* src, // premultiplied float4 per pixel, w * h * 4 floats + global uchar* dst, // straight RGBA8, w * h * 4 bytes + uint w, uint h, + int radius, + global const float* wts) // 2 * radius + 1 normalized gaussian weights +{ + uint x = get_global_id(0); + uint y = get_global_id(1); + if (x >= w || y >= h) { + return; + } + + int lo = max(-radius, -(int)y); + int hi = min(radius, (int)h - 1 - (int)y); + + float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + float wsum = 0.0f; + for (int i = lo; i <= hi; i++) { + float wt = wts[i + radius]; + acc += vload4(0, src + ((ulong)(uint)((int)y + i) * w + x) * 4) * wt; + wsum += wt; + } + float4 f = acc / max(wsum, 1e-12f); + + float a = f.w; + float4 o; + if (a > 0.0f) { + float k = 255.0f / a; + o = (float4)(f.x * k, f.y * k, f.z * k, a); + } else { + o = (float4)(0.0f, 0.0f, 0.0f, 0.0f); + } + vstore4(convert_uchar4_sat_rte(o), 0, dst + ((ulong)y * w + x) * 4); +} diff --git a/2d/total/build/blur_ocl_.spv b/2d/total/build/blur_ocl_.spv new file mode 100644 index 0000000..5938ffb Binary files /dev/null and b/2d/total/build/blur_ocl_.spv differ diff --git a/compat.go b/compat.go new file mode 100644 index 0000000..ee3d447 --- /dev/null +++ b/compat.go @@ -0,0 +1,40 @@ +package gg + +import ( + "image" + "image/color" + + "github.com/FloatTech/gg/2d/total" + "github.com/FloatTech/gg/2d/unit" +) + +// Radians converts degrees to radians. It is kept in the root package for +// backward compatibility with rendercard and other dependents. +// +// Radians 将角度转换为弧度。保留在根包中以兼容 rendercard 等依赖方。 +func Radians(degrees float64) float64 { + return unit.Radians(degrees) +} + +// ImageToNRGBA converts an image.Image to *image.NRGBA. It is kept in the +// root package for backward compatibility with rendercard and other dependents. +// +// ImageToNRGBA 将 image.Image 转换为 *image.NRGBA。保留在根包中以兼容 +// rendercard 等依赖方。 +func ImageToNRGBA(src image.Image) *image.NRGBA { + return total.ImageToNRGBA(src) +} + +// TakeThemeColorsKMeans extracts the k dominant colors from an image using +// k-means. It is kept in the root package with its legacy signature (no error +// return) for backward compatibility with rendercard and other dependents. +// +// TakeThemeColorsKMeans 使用 k-means 算法从图像中提取 k 个主色。保留在根包 +// 并沿用旧签名（无 error 返回值）以兼容 rendercard 等依赖方。 +func TakeThemeColorsKMeans(img image.Image, k uint16) []color.RGBA { + colors, err := total.TakeThemeColorsKMeans(img, k) + if err != nil { + return nil + } + return colors +} diff --git a/effects.go b/effects.go index a3bab93..eb5c19b 100644 --- a/effects.go +++ b/effects.go @@ -4,6 +4,7 @@ import ( "image" "math" + "github.com/FloatTech/gg/2d/total" "github.com/disintegration/imaging" ) @@ -44,5 +45,5 @@ func (dc *Context) Blur(s float64) { if math.Abs(s) < 0.001 { return } - dc.im = (*image.RGBA)(imaging.Blur(dc.im, s)) + dc.im = (*image.RGBA)(total.Blur(dc.im, s)) }