From f5243571d1a0a9fc2fa415021778233a5ac14884 Mon Sep 17 00:00:00 2001
From: Zonghang Li <lizhuestc@gmail.com>
Date: Sun, 8 Dec 2024 10:14:52 +0400
Subject: [PATCH] use sequential vram read test

---
 ggml/src/ggml-cuda/read.cu | 25 ++++---------------------
 1 file changed, 4 insertions(+), 21 deletions(-)

diff --git a/ggml/src/ggml-cuda/read.cu b/ggml/src/ggml-cuda/read.cu
index 47ca250b..343d908b 100644
--- a/ggml/src/ggml-cuda/read.cu
+++ b/ggml/src/ggml-cuda/read.cu
@@ -2,20 +2,11 @@
 #include "read.cuh"
 
 __global__ void read_vram_f32(
-    const float * data, int64_t ne,
-    int64_t nb00, int64_t nb01, int64_t nb02, int64_t nb03,
-    int64_t ne00, int64_t ne01, int64_t ne02
-) {
+    const float * data, int64_t ne) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx >= ne) return;
 
-    int i = idx % ne00;
-    int j = (idx / ne00) % ne01;
-    int k = (idx / (ne00 * ne01)) % ne02;
-
-    int64_t offset = i * nb00 + j * nb01 + k * nb02;
-
-    volatile float value = data[offset / sizeof(float)];
+    volatile float value = data[idx];
     asm volatile("" : : "f"(value) : "memory");
 }
 
@@ -25,23 +16,15 @@ void ggml_cuda_read(ggml_tensor * dst) {
     GGML_ASSERT(ggml_is_contiguous(dst));
     GGML_ASSERT(dst->type == GGML_TYPE_F32);
 
-    const int64_t ne00 = dst->ne[0];
-    const int64_t ne01 = dst->ne[1];
-    const int64_t ne02 = dst->ne[2];
-
-    const int64_t nb00 = dst->nb[0];
-    const int64_t nb01 = dst->nb[1];
-    const int64_t nb02 = dst->nb[2];
-    const int64_t nb03 = dst->nb[3];
-
     const char * dst_ddc = (const char *)dst->data;
 
     cudaStream_t stream;
     cudaStreamCreate(&stream);
 
     const int num_blocks = (ne + CUDA_READ_BLOCK_SIZE - 1) / CUDA_READ_BLOCK_SIZE;
+
     read_vram_f32<<<num_blocks, CUDA_READ_BLOCK_SIZE, 0, stream>>>(
-        (const float *)dst_ddc, ne, nb00, nb01, nb02, nb03, ne00, ne01, ne02
+        (const float *)dst_ddc, ne
     );
 
     cudaStreamSynchronize(stream);