Skip to content

Commit 0294869

Browse files
gbenrosciencegbenroscience
committed
llama model load an drun works on Intel(R) HD Graphics 520; crashes on AMD Radeon R5 M335, both very old GPUs. Working on tuning responses for Intel GPU, though.
1 parent 9c5581b commit 0294869

15 files changed

Lines changed: 1492 additions & 468 deletions

File tree

‎parser-ng-gpu-simd/src/main/java/com/github/gbenroscience/gpu/llm/cuda/CoreKernelBenchmark.java‎

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -591,25 +591,29 @@ private static float[] cpuQ8_0GemmTiled(byte[] xQ8Batch, byte[] wQ8, int T, int
591591
return out;
592592
}
593593

594+
/** B is [N, K] row-major -- matches the corrected f32_gemv kernel and GGUF's native Linear-weight layout. */
594595
private static float[] cpuF32Gemv(float[] a, float[] B, int K, int N) {
595596
float[] out = new float[N];
596597
for (int n = 0; n < N; n++) {
598+
int rowOff = n * K;
597599
float acc = 0f;
598600
for (int k = 0; k < K; k++) {
599-
acc += a[k] * B[k * N + n];
601+
acc += a[k] * B[rowOff + k];
600602
}
601603
out[n] = acc;
602604
}
603605
return out;
604606
}
605607

608+
/** B is [N, K] row-major -- matches the corrected f32_gemm_tiled kernel. */
606609
private static float[] cpuF32GemmTiled(float[] A, float[] B, int T, int K, int N) {
607610
float[] out = new float[T * N];
608611
for (int t = 0; t < T; t++) {
609612
for (int n = 0; n < N; n++) {
613+
int bRowOff = n * K;
610614
float acc = 0f;
611615
for (int k = 0; k < K; k++) {
612-
acc += A[t * K + k] * B[k * N + n];
616+
acc += A[t * K + k] * B[bRowOff + k];
613617
}
614618
out[t * N + n] = acc;
615619
}

‎parser-ng-gpu-simd/src/main/java/com/github/gbenroscience/gpu/llm/cuda/KernelSource.java‎

Lines changed: 12 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -448,6 +448,11 @@ __device__ __forceinline__ float gpu_gelu_f(float x) {
448448
x[i] += y[i];
449449
}
450450
451+
// NOTE (fixed): B is [N, K] row-major -- N=out_features rows,
452+
// K=in_features cols, matching GGUF's native Linear-weight layout.
453+
// The earlier version read B as [K, N] (B[k*N+n]) -- silently
454+
// transposed for wo (K==N==dim, no shape check could catch it) and
455+
// for the LM head (K=dim, N=vocab, still in-bounds, still wrong).
451456
extern "C" __global__ void f32_gemv(
452457
const float* a,
453458
const float* B,
@@ -458,9 +463,10 @@ __device__ __forceinline__ float gpu_gelu_f(float x) {
458463
const int n = blockIdx.x * blockDim.x + threadIdx.x;
459464
if (n >= N) return;
460465
466+
const int rowOff = n * K;
461467
float acc = 0.0f;
462468
for (int k = 0; k < K; k++) {
463-
acc += a[k] * B[k * N + n];
469+
acc += a[k] * B[rowOff + k];
464470
}
465471
out[n] = acc;
466472
}
@@ -541,9 +547,11 @@ __device__ __forceinline__ float gpu_gelu_f(float x) {
541547
out[t * N + n] = acc;
542548
}
543549
550+
// NOTE (fixed): B is [N, K] row-major, matching GGUF's native
551+
// Linear-weight layout -- same fix and rationale as f32_gemv above.
544552
extern "C" __global__ void f32_gemm_tiled(
545553
const float* A, // [T, K]
546-
const float* B, // [K, N]
554+
const float* B, // [N, K]
547555
float* out, // [T, N]
548556
const int T,
549557
const int K,
@@ -553,9 +561,10 @@ __device__ __forceinline__ float gpu_gelu_f(float x) {
553561
const int n = blockIdx.x * blockDim.x + threadIdx.x;
554562
if (t >= T || n >= N) return;
555563
564+
const int bRowOff = n * K;
556565
float acc = 0.0f;
557566
for (int k = 0; k < K; k++) {
558-
acc += A[t * K + k] * B[k * N + n];
567+
acc += A[t * K + k] * B[bRowOff + k];
559568
}
560569
out[t * N + n] = acc;
561570
}

‎parser-ng-gpu-simd/src/main/java/com/github/gbenroscience/gpu/llm/opencl/CoreKernelBenchmark.java‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -61,7 +61,6 @@ private CoreKernelBenchmark() {
6161
private static final float TOLERANCE_ABS = 1e-2f;
6262

6363
public static void main(String[] args) throws Throwable {
64-
OpenCLDeviceSelector.selectDevice(OpenCLDeviceSelector.GpuVendor.INTEL);
6564
GpuContext ctx = new GpuContext();
6665
try {
6766
// Prime every kernel this file touches before anything is
@@ -623,25 +622,29 @@ private static float[] cpuQ8_0GemmTiled(byte[] xQ8Batch, byte[] wQ8, int T, int
623622
return out;
624623
}
625624

625+
/** B is [N, K] row-major -- matches the corrected f32_gemv kernel and GGUF's native Linear-weight layout. */
626626
private static float[] cpuF32Gemv(float[] a, float[] B, int K, int N) {
627627
float[] out = new float[N];
628628
for (int n = 0; n < N; n++) {
629+
int rowOff = n * K;
629630
float acc = 0f;
630631
for (int k = 0; k < K; k++) {
631-
acc += a[k] * B[k * N + n];
632+
acc += a[k] * B[rowOff + k];
632633
}
633634
out[n] = acc;
634635
}
635636
return out;
636637
}
637638

639+
/** B is [N, K] row-major -- matches the corrected f32_gemm_tiled kernel. */
638640
private static float[] cpuF32GemmTiled(float[] A, float[] B, int T, int K, int N) {
639641
float[] out = new float[T * N];
640642
for (int t = 0; t < T; t++) {
641643
for (int n = 0; n < N; n++) {
644+
int bRowOff = n * K;
642645
float acc = 0f;
643646
for (int k = 0; k < K; k++) {
644-
acc += A[t * K + k] * B[k * N + n];
647+
acc += A[t * K + k] * B[bRowOff + k];
645648
}
646649
out[t * N + n] = acc;
647650
}

‎parser-ng-gpu-simd/src/main/java/com/github/gbenroscience/gpu/llm/opencl/KernelSource.java‎

Lines changed: 16 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -514,6 +514,15 @@ __kernel void residual_add(
514514
x[i] += y[i];
515515
}
516516
517+
// NOTE (fixed): B is [N, K] row-major -- N=out_features rows, K=in_features
518+
// cols, matching GGUF's native Linear-weight layout (PyTorch dumps
519+
// [out_features, in_features] as-is, no transpose on conversion).
520+
// This now matches q8_0_gemv_plain's convention. The earlier version
521+
// read B as [K, N] (B[k*N+n]) -- silently transposed for wo (where
522+
// K==N==dim, so no shape check could ever catch it) and for the LM
523+
// head (K=dim, N=vocab, still in-bounds, still wrong). See
524+
// LlamaLayer's/this kernel's git history -- or ask the person who
525+
// introduced this fix -- for the full derivation.
517526
__kernel void f32_gemv(
518527
__global const float* a,
519528
__global const float* B,
@@ -524,9 +533,10 @@ __kernel void f32_gemv(
524533
const int n = get_group_id(0) * get_local_size(0) + get_local_id(0);
525534
if (n >= N) return;
526535
536+
const int rowOff = n * K;
527537
float acc = 0.0f;
528538
for (int k = 0; k < K; k++) {
529-
acc += a[k] * B[k * N + n];
539+
acc += a[k] * B[rowOff + k];
530540
}
531541
out[n] = acc;
532542
}
@@ -597,9 +607,11 @@ __kernel void q8_0_gemm_tiled(
597607
out[t * N + n] = acc;
598608
}
599609
610+
// NOTE (fixed): B is [N, K] row-major, matching GGUF's native
611+
// Linear-weight layout -- same fix and rationale as f32_gemv above.
600612
__kernel void f32_gemm_tiled(
601613
__global const float* A, // [T, K]
602-
__global const float* B, // [K, N]
614+
__global const float* B, // [N, K]
603615
__global float* out, // [T, N]
604616
const int T,
605617
const int K,
@@ -609,9 +621,10 @@ __kernel void f32_gemm_tiled(
609621
const int n = get_group_id(0) * get_local_size(0) + get_local_id(0);
610622
if (t >= T || n >= N) return;
611623
624+
const int bRowOff = n * K;
612625
float acc = 0.0f;
613626
for (int k = 0; k < K; k++) {
614-
acc += A[t * K + k] * B[k * N + n];
627+
acc += A[t * K + k] * B[bRowOff + k];
615628
}
616629
out[t * N + n] = acc;
617630
}

‎parser-ng-gpu-simd/src/main/java/com/github/gbenroscience/gpu/llm/opencl/OpenCLDemo.java‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ private OpenCLDemo() {
4343
public static void main(String[] args) throws Throwable {
4444
if (args.length == 0) {
4545
args = new String[]{"C:\\Users\\GBEMIRO\\Documents\\NetBeansProjects\\ai-models\\llama-3.2-1b-instruct-q8_0.gguf",
46-
"What is 23*34?"};
46+
"The capital of France is"};
4747
}
4848
if (args.length < 2) {
4949
System.err.println("Usage: OpenCLDemo <path-to-model.gguf> <prompt>");

‎parser-ng/src/main/java/com/github/gbenroscience/math/differentialcalculus/equations/coeffextractor/clext/ArgumentIsolator.java‎

Lines changed: 1 addition & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1,26 +1,9 @@
1-
/*
2-
* Copyright 2026 GBEMIRO.
3-
*
4-
* Licensed under the Apache License, Version 2.0 (the "License");
5-
* you may not use this file except in compliance with the License.
6-
* You may obtain a copy of the License at
7-
*
8-
* http://www.apache.org/licenses/LICENSE-2.0
9-
*
10-
* Unless required by applicable law or agreed to in writing, software
11-
* distributed under the License is distributed on an "AS IS" BASIS,
12-
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13-
* See the License for the specific language governing permissions and
14-
* limitations under the License.
15-
*/
161
package com.github.gbenroscience.math.differentialcalculus.equations.coeffextractor.clext;
17-
18-
2+
193
import java.util.ArrayList;
204
import java.util.List;
215

226
/**
23-
* @author GBEMIRO
247
* Isolates the first top-level argument of a function call from its token
258
* stream, e.g. given the full scan of
269
* {@code diffeqn(<equation>, 1, 0, anon1)}, returns just the {@code

0 commit comments

Comments
 (0)