From a3b4d8d31eec48fe9977b4e5150f85947f9c2871 Mon Sep 17 00:00:00 2001 From: bssrdf Date: Wed, 29 Oct 2025 21:46:15 -0400 Subject: [PATCH] clean up --- ggml/src/ggml-cuda/cpy.cu | 121 +--------- ggml/src/ggml-cuda/cpy.cuh | 5 - ggml/src/ggml.c | 1 - tests/CMakeLists.txt | 2 - tests/test-conv2d-implicit.cpp | 413 --------------------------------- tests/test-transpose.cpp | 375 ------------------------------ 6 files changed, 5 insertions(+), 912 deletions(-) delete mode 100644 tests/test-conv2d-implicit.cpp delete mode 100644 tests/test-transpose.cpp diff --git a/ggml/src/ggml-cuda/cpy.cu b/ggml/src/ggml-cuda/cpy.cu index 514657537f..c0a568f4ab 100644 --- a/ggml/src/ggml-cuda/cpy.cu +++ b/ggml/src/ggml-cuda/cpy.cu @@ -37,90 +37,6 @@ static __global__ void cpy_flt(const char * cx, char * cdst_direct, const int ne cpy_1(cx + x_offset, cdst + dst_offset); } - -template -static __global__ void cpy_flt_transpose(const char * cx, char * cdst_direct, const int ne, - const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02, - const int nb03, const int ne10, const int ne11, const int ne12, const int nb10, const int nb11, - const int nb12, const int nb13, char ** cdst_indirect, int graph_cpynode_index) { - - char * cdst = (cdst_indirect != nullptr) ? cdst_indirect[graph_cpynode_index]: cdst_direct; - - const T* src = reinterpret_cast(cx); - T* dst = reinterpret_cast(cdst); - - const int64_t nmat = ne / (ne00 * ne01); - const int64_t n = ne00 * ne01; - int width = ne01; - int height = ne00; - int x = blockIdx.x * TILE_DIM + threadIdx.x; - int y = blockIdx.y * TILE_DIM + threadIdx.y; - int tx = blockIdx.y * TILE_DIM + threadIdx.x; // transpose block offset - int ty = blockIdx.x * TILE_DIM + threadIdx.y; - - __shared__ T tile[TILE_DIM][TILE_DIM]; - - for(int i = 0; i < BLOCK_NM; ++i){ - - const unsigned int imat = blockIdx.z * BLOCK_NM + i; - if(imat >= nmat) - break; - for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS){ - // if(imat < nmat && x < width && y + j < height){ - if(x < width && y + j < height){ - const unsigned int idx = (y+j)*width + x; - const int row = threadIdx.y+j; - const int col = threadIdx.x ^ row; - // tile[threadIdx.y+j][threadIdx.x] = src[imat*n + idx]; - tile[row][col] = src[imat*n + idx]; - } - } - __syncthreads(); - - - // if(threadIdx.x == 0 && threadIdx.y == 0 && blockIdx.x == 0 && blockIdx.y == 0 && blockIdx.z == 0){ - // printf("BEGIN %d\n", i); - // for(int jj = 0; jj < TILE_DIM; ++jj){ - // for(int ii = 0; ii < TILE_DIM; ++ii) - // printf("%.f, ", tile[jj][ii]); - // printf("]\n"); - // } - // } - - for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS){ - // if(imat < nmat && ty + j < width && tx < height){ - if(ty + j < width && tx < height){ - const unsigned int idx = (ty+j)*height + tx; - const int col = (threadIdx.y+j) ^ threadIdx.x; - // dst[imat*n + idx] = tile[threadIdx.x][threadIdx.y + j]; - dst[imat*n + idx] = tile[threadIdx.x][col]; - // if(imat*n + idx == 4*ne00){ - // printf("DEBUG: (%u, %u, %u, %u, %u), j=%d, tx=%d, ty=%d, imat=%u idx=%u dst[%u]=%.2f, %f\n", - // threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, blockIdx.z, j, tx, ty, - // imat, idx, imat*n + idx, dst[imat*n + idx], tile[threadIdx.x][threadIdx.y + j]); - // } - } - } - } - - // if(threadIdx.x == 0 && threadIdx.y == 0 && blockIdx.x == 0 && blockIdx.y == 0 && blockIdx.z == 0){ - // // for(int j = 0; j < 32; ++j){ - // // j = 0; - // for(int i = 0; i < 32; ++i) - // // printf("%.2f, ", src[j*48+i]); - // // printf("%.2f, ", src[j*48+i]); - // printf("%.2f, ", __half2float(src[i])); - // printf("]\n"); - // // } - // printf("==============================\n"); - // // for(int j = 0; j < 32; ++j){ - // for(int i = 0; i < 32; ++i) - // printf("%.2f, ", __half2float(dst[i])); - // printf("]\n"); - // // } - // } -} - static __device__ void cpy_blck_q8_0_f32(const char * cxi, char * cdsti) { float * cdstf = (float *)(cdsti); @@ -228,28 +144,9 @@ static void ggml_cpy_flt_cuda( const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02, const int nb03, const int ne10, const int ne11, const int ne12, const int nb10, const int nb11, const int nb12, const int nb13, cudaStream_t stream, char ** cdst_indirect, int & graph_cpynode_index) { - if constexpr ((std::is_same_v && std::is_same_v || - std::is_same_v && std::is_same_v) - && transpose){ - // printf("cuda cpy transpose ne=%d ne00=%d ne01=%d ne10=%d ne11=%d\n", ne, ne00, ne01, ne10, ne11); - // printf("cuda cpy transpose nb00=%d nb01=%d nb10=%d nb11=%d\n", nb00, nb01, nb10, nb11); - // if (ne00 == ne11 && ne01 == ne10 && nb00 == nb11 && nb10 == nb01){ //transpose - // if (transpose) { //transpose - // printf("cuda cpy transpose ne=%d ne00=%d ne01=%d ne10=%d ne11=%d\n", ne, ne00, ne01, ne10, ne11); - dim3 dimGrid( (ne01 + TILE_DIM - 1) / TILE_DIM, - (ne00 + TILE_DIM - 1) / TILE_DIM, - (ne/(ne00*ne01) + BLOCK_NM - 1) / BLOCK_NM ); - dim3 dimBlock(TILE_DIM, BLOCK_ROWS, 1); - cpy_flt_transpose<<>>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, cdst_indirect, graph_cpynode_index++); - } else{ // other - const int num_blocks = (ne + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; - cpy_flt><<>> - (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, cdst_indirect, graph_cpynode_index++); - } - // } else{ - // cpy_flt><<>> - // (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, cdst_indirect, graph_cpynode_index++); - // } + const int num_blocks = (ne + CUDA_CPY_BLOCK_SIZE - 1) / CUDA_CPY_BLOCK_SIZE; + cpy_flt><<>> + (cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, cdst_indirect, graph_cpynode_index++); } static void ggml_cpy_f32_q8_0_cuda( @@ -435,11 +332,7 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg CUDA_CHECK(cudaMemcpyAsync(src1_ddc, src0_ddc, ggml_nbytes(src0), cudaMemcpyDeviceToDevice, main_stream)); } } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32) { - if(src0->op_params[10] == 999){ - ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); - } else { - ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); - } + ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_BF16) { ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); } else if (src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F16) { @@ -470,11 +363,7 @@ void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, gg } else if (src0->type == GGML_TYPE_Q5_1 && src1->type == GGML_TYPE_F32) { ggml_cpy_q5_1_f32_cuda(src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16) { - if(src0->op_params[10] == 999){ - ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); - } else { - ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); - } + ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_BF16) { ggml_cpy_flt_cuda (src0_ddc, src1_ddc, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, main_stream, dest_ptrs_d, graph_cpynode_index); } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32) { diff --git a/ggml/src/ggml-cuda/cpy.cuh b/ggml/src/ggml-cuda/cpy.cuh index 211348b66a..0bd3c0c6f8 100644 --- a/ggml/src/ggml-cuda/cpy.cuh +++ b/ggml/src/ggml-cuda/cpy.cuh @@ -2,11 +2,6 @@ #define CUDA_CPY_BLOCK_SIZE 64 -const int TILE_DIM = 32; -const int BLOCK_ROWS = 8; -const int BLOCK_NM = 8; - - void ggml_cuda_cpy(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, ggml_tensor * src1, bool disable_indirection = false); void ggml_cuda_dup(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 03c8dca3e5..a792d6b888 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -3612,7 +3612,6 @@ struct ggml_tensor * ggml_transpose( result->op = GGML_OP_TRANSPOSE; result->src[0] = a; - result->op_params[10] = 999; // the transpose flag return result; } diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 1787e53eb5..9171957756 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -198,8 +198,6 @@ if (NOT LLAMA_SANITIZE_ADDRESS) endif() llama_build_and_test(test-gguf.cpp) llama_build_and_test(test-backend-ops.cpp) -llama_build_and_test(test-conv2d-implicit.cpp) -llama_build_and_test(test-transpose.cpp) llama_build_and_test(test-model-load-cancel.cpp LABEL "model") llama_build_and_test(test-autorelease.cpp LABEL "model") diff --git a/tests/test-conv2d-implicit.cpp b/tests/test-conv2d-implicit.cpp deleted file mode 100644 index 7b7a32d9f6..0000000000 --- a/tests/test-conv2d-implicit.cpp +++ /dev/null @@ -1,413 +0,0 @@ -#include "ggml.h" -#include "ggml-alloc.h" -#include "ggml-cpu.h" -#include "ggml-backend.h" - -#ifdef GGML_USE_CUDA -#include "ggml-cuda.h" -//#include -#endif - -#ifdef GGML_USE_METAL -#include "ggml-metal.h" -#endif - -#include -#include -#include -#include -#include -#include -#include -#include - -static void ggml_log_callback_default(ggml_log_level level, const char * text, void * user_data) { - (void) level; - (void) user_data; - fputs(text, stderr); - fflush(stderr); -} - -struct test_model { - struct ggml_tensor * a; - struct ggml_tensor * b; - ggml_backend_t backend = NULL; - ggml_backend_buffer_t buffer; - struct ggml_context * ctx; -}; - - - -void load_model(test_model & model, int ic, int oc, int iw, int ih, int kw = 3, int kh = 3, bool use_gpu = false ) { - // create data - int KW = kw, KH = kh, IC = ic, OC = oc; - int IW = iw, IH = ih, N = 1; - srand(time(NULL)); - - // printf(" input: IC = %d, OC = %d, IW = %d, IH = %d \n ", IC, OC, IW, IH); - - // Initialize adata - std::vector adata(KW * KH * IC * OC); - for (int i = 0; i < KW * KH * IC * OC; i++) { - // adata[i] = 2.f; - // adata[i] = (float)(i%KW)-1.f; - // adata[i] = (rand() % 255) / 255.0; - float r = -1.f + static_cast (rand()) /( static_cast (RAND_MAX/(1.f-(-1.f)))); - adata[i] = r; - } - - // Convert adata to fp16 format - std::vector hadata(KW * KH * IC * OC); - ggml_fp32_to_fp16_row(adata.data(), hadata.data(), KW * KH * IC * OC); - - // Initialize bdata - std::vector bdata(IW * IH * IC * N); - for (int i = 0; i < IW * IH * IC * N; i++) { - // bdata[i] = (float)(i%IW)/10.f; - // bdata[i] = 1.5f; - // bdata[i] = (rand() % 255) / 255.0; - float r = -1.f + static_cast (rand()) /( static_cast (RAND_MAX/(1.f-(-1.f)))); - bdata[i] = r; - } - - size_t buffer_size = 0; - { - // buffer_size += KW * KH * IC * OC * ggml_type_size(GGML_TYPE_F32); // tensor a - buffer_size += KW * KH * IC * OC * ggml_type_size(GGML_TYPE_F16); // tensor a - buffer_size += IW * IH * IC * N * ggml_type_size(GGML_TYPE_F32); // tensor b - buffer_size += 1024; // overhead - } - - // printf("%s: ggml tensor size = %d bytes\n", __func__, (int) sizeof(ggml_tensor)); - // printf("%s: backend buffer size = %0.2f MB\n", __func__, (buffer_size/ 1024.f/ 1024.f)); - - int num_tensors = 2; - struct ggml_init_params params { - /*.mem_size =*/ ggml_tensor_overhead() * num_tensors, - /*.mem_buffer =*/ NULL, - /*.no_alloc =*/ true, - }; - - // initialize the backend -#ifdef GGML_USE_CUDA - if (use_gpu) { - // fprintf(stderr, "%s: using CUDA backend\n", __func__); - model.backend = ggml_backend_cuda_init(0); - if (!model.backend) { - fprintf(stderr, "%s: ggml_backend_cuda_init() failed\n", __func__); - } - } -#endif - -#ifdef GGML_USE_METAL - if (use_gpu) { - fprintf(stderr, "%s: using Metal backend\n", __func__); - ggml_backend_metal_log_set_callback(ggml_log_callback_default, nullptr); - model.backend = ggml_backend_metal_init(); - if (!model.backend) { - fprintf(stderr, "%s: ggml_backend_metal_init() failed\n", __func__); - } - } -#endif - - if(!model.backend) { - // fallback to CPU backend - model.backend = ggml_backend_cpu_init(); - } - - model.buffer = ggml_backend_alloc_buffer(model.backend, buffer_size); - - // create context - model.ctx = ggml_init(params); - - // create tensors - model.a = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F16, KW, KH, IC, OC); - // model.a = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F32, KW, KH, IC, OC); - model.b = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F32, IW, IH, IC, N); - - // create a allocator - struct ggml_tallocr alloc = ggml_tallocr_new(model.buffer); - - // alloc memory - ggml_tallocr_alloc(&alloc, model.a); - - // load data to buffer - if(ggml_backend_is_cpu(model.backend)) { - memcpy(model.a->data, hadata.data(), ggml_nbytes(model.a)); - // memcpy(model.a->data, adata.data(), ggml_nbytes(model.a)); - } else { - ggml_backend_tensor_set(model.a, hadata.data(), 0, ggml_nbytes(model.a)); - // ggml_backend_tensor_set(model.a, adata.data(), 0, ggml_nbytes(model.a)); - } - - // alloc memory - ggml_tallocr_alloc(&alloc, model.b); - - if(ggml_backend_is_cpu(model.backend) -#ifdef GGML_USE_METAL - || ggml_backend_is_metal(model.backend) -#endif - ) { - memcpy(model.b->data, bdata.data(), ggml_nbytes(model.b)); - } else { - ggml_backend_tensor_set(model.b, bdata.data(), 0, ggml_nbytes(model.b)); - } -} - -typedef struct ggml_cgraph* (*build_graph_t)(const test_model& model); - -struct ggml_cgraph * build_graph_0(const test_model& model) { - static size_t buf_size = ggml_tensor_overhead()*GGML_DEFAULT_GRAPH_SIZE + ggml_graph_overhead(); - static std::vector buf(buf_size); - - struct ggml_init_params params0 = { - /*.mem_size =*/ buf_size, - /*.mem_buffer =*/ buf.data(), - /*.no_alloc =*/ true, // the tensors will be allocated later by ggml_gallocr_alloc_graph() - }; - - // create a temporally context to build the graph - struct ggml_context * ctx0 = ggml_init(params0); - - struct ggml_cgraph * gf = ggml_new_graph(ctx0); - - int s0 = 1; - int s1 = 1; - int p0 = 1; - int p1 = 1; - int d0 = 1; - int d1 = 1; - - - - // recalculate for avoid fragmentation - struct ggml_tensor* conv2d_res = ggml_conv_2d(ctx0, model.a, model.b, s0, s1, p0, p1, d0, d1); - ggml_set_name(conv2d_res, "conv2d_res"); - ggml_build_forward_expand(gf, conv2d_res); - // int64_t *ne = conv2d_res->ne; - // printf("conv2d: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - - - // struct ggml_tensor* wino_res = ggml_conv_2d_3x3(ctx0, model.a, model.b); - // ggml_set_name(wino_res, "wino_res"); - // ggml_build_forward_expand(gf, wino_res); - // ne = wino_res->ne; - // printf("wino: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - ggml_free(ctx0); - return gf; -} - -struct ggml_cgraph * build_graph_1(const test_model& model) { - static size_t buf_size = ggml_tensor_overhead()*GGML_DEFAULT_GRAPH_SIZE + ggml_graph_overhead(); - static std::vector buf(buf_size); - - struct ggml_init_params params0 = { - /*.mem_size =*/ buf_size, - /*.mem_buffer =*/ buf.data(), - /*.no_alloc =*/ true, // the tensors will be allocated later by ggml_gallocr_alloc_graph() - }; - - // create a temporally context to build the graph - struct ggml_context * ctx0 = ggml_init(params0); - - struct ggml_cgraph * gf = ggml_new_graph(ctx0); - - int s0 = 1; - int s1 = 1; - int p0 = 1; - int p1 = 1; - int d0 = 1; - int d1 = 1; - - - - // recalculate for avoid fragmentation - // struct ggml_tensor* conv2d_res = ggml_conv_2d(ctx0, model.a, model.b, s0, s1, p0, p1, d0, d1); - // ggml_set_name(conv2d_res, "conv2d_res"); - // ggml_build_forward_expand(gf, conv2d_res); - // int64_t *ne = conv2d_res->ne; - // printf("conv2d: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - - - // struct ggml_tensor* wino_res = ggml_conv_2d_implicitgemm(ctx0, model.a, model.b, s0, s1, p0, p1, d0, d1); - struct ggml_tensor* wino_res = ggml_conv_2d_direct(ctx0, model.a, model.b, s0, s1, p0, p1, d0, d1); - ggml_set_name(wino_res, "wino_res"); - ggml_build_forward_expand(gf, wino_res); - // ne = wino_res->ne; - // printf("wino: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - ggml_free(ctx0); - return gf; -} - - -std::vector compute_graph(const test_model & model, ggml_gallocr_t allocr, - build_graph_t build_graph, int iters, double *t) { - struct ggml_cgraph * gf = build_graph(model); - - - // allocate tensors - ggml_gallocr_alloc_graph(allocr, gf); - int n_threads = 1; - - if (ggml_backend_is_cpu(model.backend)) { - ggml_backend_cpu_set_n_threads(model.backend, n_threads); - } - -#ifdef GGML_USE_METAL - if (ggml_backend_is_metal(model.backend)) { - ggml_backend_metal_set_n_cb(model.backend, n_threads); - } -#endif - - - - ggml_backend_graph_compute(model.backend, gf); - - ggml_backend_synchronize(model.backend); - - int64_t start_time = ggml_time_us(); - - for(int iter=0; iter data(ggml_nelements(res)); - ggml_backend_tensor_get(res, data.data(), 0, ggml_nbytes(res)); - - *t = time_us/1000; - return data; - -} - - -int main(void) -{ - ggml_time_init(); - std::vector> configs = { - // std::make_tuple(64,64,48,64,3,3), - // std::make_tuple(320,320,104,152,3,3), - // std::make_tuple(640,640,52,76,3,3), - // std::make_tuple(640,640,104,152,3,3), - // std::make_tuple(960,320,104,152,3,3), - std::make_tuple(1280,1280,26,38,3,3), - // std::make_tuple(1280,1280,26,38,1,1), - // std::make_tuple(256,128,768,1024,3,3), - // std::make_tuple(128,3,768,1024,3,3), - // std::make_tuple(256,128,768,1024,1,1), - // std::make_tuple(512,256,384,512,1,1), - // std::make_tuple(1280,640,52,76,3,3), - // std::make_tuple(1920,1280,26,38,3,3), - // std::make_tuple(2560,1280,26,38,3,3), - // std::make_tuple(512,512,104,152,3,3), - // std::make_tuple(512,512,208,304,3,3), - // std::make_tuple(512,256,416,608,3,3), - // std::make_tuple(256,128,832,1216,3,3), - // std::make_tuple(256,256,832,1216,3,3), - // std::make_tuple(320,256,1024,1920) - }; - - int k = 0; - - for (auto c : configs){ - test_model model; - load_model(model, std::get<0>(c), std::get<1>(c), std::get<2>(c), - std::get<3>(c), std::get<4>(c), std::get<5>(c), true); - - ggml_gallocr_t allocr = NULL; - allocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); - - //create the worst case graph for memory usage estimation - struct ggml_cgraph * gf = build_graph_0(model); - - // compute the required memory - ggml_gallocr_reserve(allocr, gf); - size_t mem_size0 = ggml_gallocr_get_buffer_size(allocr, 0); - // fprintf(stderr, "%s: compute buffer size: %.2f MB\n", __func__, mem_size/1024.0f/1024.0f); - - - struct ggml_cgraph * gf_res_0 = NULL; - int iterations = 0; - - double run_time0; - std::vector im2col_data = compute_graph(model, allocr, build_graph_0, iterations, &run_time0); - - ggml_gallocr_free(allocr); - - allocr = NULL; - - allocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); - - //create the worst case graph for memory usage estimation - gf = build_graph_1(model); - - // compute the required memory - ggml_gallocr_reserve(allocr, gf); - size_t mem_size1 = ggml_gallocr_get_buffer_size(allocr, 0); - // fprintf(stderr, "%s: compute buffer size: %.2f MB\n", __func__, mem_size/1024.0f/1024.0f); - - - struct ggml_cgraph * gf_res_1 = NULL; - - double run_time1; - // std::vector wino_data = compute_graph(model, allocr, build_graph_1, iterations, &run_time1); - std::vector conv2d_data = compute_graph(model, allocr, build_graph_1, iterations, &run_time1); - - - ggml_gallocr_free(allocr); - - if(k==0) { - k = 1; - fprintf(stderr, "| (IC, OC, IW, IH) | im2col+GEMM TIME | im2col+GEMM VRAM | implicit GEMM TIME | implicit GEMM VRAM \n"); - fprintf(stderr, "| --- | --- | --- | --- | --- \n"); - } - - fprintf(stderr, " | (%d, %d, %d, %d) | %.2f ms | %.2f MB | %.2f ms | %.2f MB\n", - std::get<0>(c), std::get<1>(c), std::get<2>(c), std::get<3>(c), std::get<4>(c), std::get<5>(c), - run_time0, mem_size0/1024.0f/1024.0f, - run_time1, mem_size1/1024.0f/1024.0f - ); - - - // for(int i = 0; i < ggml_nelements(wino_res); i++) { - // for(int i = 0; i < 26*38; i++) { - for(int i = 0; i < conv2d_data.size(); i++) { - // float diff = fabs(conv2d_data[i] - wino_data[i]); - float diff = fabs(im2col_data[i] - conv2d_data[i]); - // if(diff > 0.5) { - printf("(%7.3f, %7.3f, %.2f, %d) \n", - im2col_data[i], conv2d_data[i], - diff, i); - // break; - // } - } - - ggml_free(model.ctx); - ggml_backend_buffer_free(model.buffer); - ggml_backend_free(model.backend); - ggml_gallocr_free(allocr); - - } - - // printf("\nPerforming test:\n"); - return 0; -} diff --git a/tests/test-transpose.cpp b/tests/test-transpose.cpp deleted file mode 100644 index 73263f3438..0000000000 --- a/tests/test-transpose.cpp +++ /dev/null @@ -1,375 +0,0 @@ -#include "ggml.h" -#include "ggml-alloc.h" -#include "ggml-cpu.h" -#include "ggml-backend.h" - -#ifdef GGML_USE_CUDA -#include "ggml-cuda.h" -//#include -#endif - -#ifdef GGML_USE_METAL -#include "ggml-metal.h" -#endif - -#include -#include -#include -#include -#include -#include -#include -#include - -static void ggml_log_callback_default(ggml_log_level level, const char * text, void * user_data) { - (void) level; - (void) user_data; - fputs(text, stderr); - fflush(stderr); -} - -struct test_model { - struct ggml_tensor * a; - struct ggml_tensor * b; - ggml_backend_t backend = NULL; - ggml_backend_buffer_t buffer; - struct ggml_context * ctx; -}; - - - -void load_model(test_model & model, int ic, int oc, int iw, int ih, int kw = 3, int kh = 3, bool use_gpu = false ) { - // create data - int KW = kw, KH = kh, IC = ic, OC = oc; - int IW = iw, IH = ih, N = 1; - srand(time(NULL)); - - // printf(" input: IC = %d, OC = %d, IW = %d, IH = %d \n ", IC, OC, IW, IH); - - // Initialize adata - std::vector adata(KW * KH * IC * OC); - for (int i = 0; i < KW * KH * IC * OC; i++) { - // adata[i] = 2.f; - adata[i] = (float)i; - // adata[i] = (rand() % 255) / 255.0; - // float r = -1.f + static_cast (rand()) /( static_cast (RAND_MAX/(1.f-(-1.f)))); - // adata[i] = r; - } - - // Convert adata to fp16 format - std::vector hadata(KW * KH * IC * OC); - ggml_fp32_to_fp16_row(adata.data(), hadata.data(), KW * KH * IC * OC); - - // Initialize bdata - std::vector bdata(IW * IH * IC * N); - for (int i = 0; i < IW * IH * IC * N; i++) { - // bdata[i] = (float)(i%IW)/10.f; - // bdata[i] = 1.5f; - bdata[i] = (float)(i+1); - // bdata[i] = (rand() % 255) / 255.0; - // float r = -1.f + static_cast (rand()) /( static_cast (RAND_MAX/(1.f-(-1.f)))); - // bdata[i] = r; - } - - // for(int i = 0; i < IH; i++) { - // // float diff = fabs(conv2d_data[i] - wino_data[i]); - // for(int j = 0; j < IW; j++) { - // printf("%.0f, ", bdata[i*IW+j]); - // } - // printf("\n"); - // } - for(int i = 0; i < KH; i++) { - // float diff = fabs(conv2d_data[i] - wino_data[i]); - for(int j = 0; j < KW; j++) { - printf("%.0f, ", adata[i*KW+j]); - } - printf("\n"); - } - printf(">>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>\n"); - - size_t buffer_size = 0; - { - // buffer_size += KW * KH * IC * OC * ggml_type_size(GGML_TYPE_F32); // tensor a - buffer_size += KW * KH * IC * OC * ggml_type_size(GGML_TYPE_F16); // tensor a - buffer_size += IW * IH * IC * N * ggml_type_size(GGML_TYPE_F32); // tensor b - buffer_size += 1024; // overhead - } - - // printf("%s: ggml tensor size = %d bytes\n", __func__, (int) sizeof(ggml_tensor)); - // printf("%s: backend buffer size = %0.2f MB\n", __func__, (buffer_size/ 1024.f/ 1024.f)); - - int num_tensors = 2; - struct ggml_init_params params { - /*.mem_size =*/ ggml_tensor_overhead() * num_tensors, - /*.mem_buffer =*/ NULL, - /*.no_alloc =*/ true, - }; - - // initialize the backend -#ifdef GGML_USE_CUDA - if (use_gpu) { - // fprintf(stderr, "%s: using CUDA backend\n", __func__); - model.backend = ggml_backend_cuda_init(0); - if (!model.backend) { - fprintf(stderr, "%s: ggml_backend_cuda_init() failed\n", __func__); - } - } -#endif - -#ifdef GGML_USE_METAL - if (use_gpu) { - fprintf(stderr, "%s: using Metal backend\n", __func__); - ggml_backend_metal_log_set_callback(ggml_log_callback_default, nullptr); - model.backend = ggml_backend_metal_init(); - if (!model.backend) { - fprintf(stderr, "%s: ggml_backend_metal_init() failed\n", __func__); - } - } -#endif - - if(!model.backend) { - // fallback to CPU backend - model.backend = ggml_backend_cpu_init(); - } - - model.buffer = ggml_backend_alloc_buffer(model.backend, buffer_size); - - // create context - model.ctx = ggml_init(params); - - // create tensors - model.a = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F16, KW, KH, IC, OC); - // model.a = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F32, KW, KH, IC, OC); - model.b = ggml_new_tensor_4d(model.ctx, GGML_TYPE_F32, IW, IH, IC, N); - - int64_t *ne = model.a->ne; - printf("before trans: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - - // create a allocator - struct ggml_tallocr alloc = ggml_tallocr_new(model.buffer); - - // alloc memory - ggml_tallocr_alloc(&alloc, model.a); - - // load data to buffer - if(ggml_backend_is_cpu(model.backend)) { - memcpy(model.a->data, hadata.data(), ggml_nbytes(model.a)); - // memcpy(model.a->data, adata.data(), ggml_nbytes(model.a)); - } else { - ggml_backend_tensor_set(model.a, hadata.data(), 0, ggml_nbytes(model.a)); - // ggml_backend_tensor_set(model.a, adata.data(), 0, ggml_nbytes(model.a)); - } - - // alloc memory - ggml_tallocr_alloc(&alloc, model.b); - - if(ggml_backend_is_cpu(model.backend) -#ifdef GGML_USE_METAL - || ggml_backend_is_metal(model.backend) -#endif - ) { - memcpy(model.b->data, bdata.data(), ggml_nbytes(model.b)); - } else { - ggml_backend_tensor_set(model.b, bdata.data(), 0, ggml_nbytes(model.b)); - } -} - -typedef struct ggml_cgraph* (*build_graph_t)(const test_model& model); - -struct ggml_cgraph * build_graph_0(const test_model& model) { - static size_t buf_size = ggml_tensor_overhead()*GGML_DEFAULT_GRAPH_SIZE + ggml_graph_overhead(); - static std::vector buf(buf_size); - - struct ggml_init_params params0 = { - /*.mem_size =*/ buf_size, - /*.mem_buffer =*/ buf.data(), - /*.no_alloc =*/ true, // the tensors will be allocated later by ggml_gallocr_alloc_graph() - }; - - // create a temporally context to build the graph - struct ggml_context * ctx0 = ggml_init(params0); - - struct ggml_cgraph * gf = ggml_new_graph(ctx0); - - int s0 = 1; - int s1 = 1; - int p0 = 1; - int p1 = 1; - int d0 = 1; - int d1 = 1; - - - - // recalculate for avoid fragmentation - // struct ggml_tensor* conv2d_res = ggml_cont(ctx0, ggml_transpose(ctx0, model.b)); - struct ggml_tensor* conv2d_res = ggml_cont(ctx0, ggml_transpose(ctx0, model.a)); - ggml_set_name(conv2d_res, "transpose_res"); - ggml_build_forward_expand(gf, conv2d_res); - int64_t *ne = conv2d_res->ne; - printf("conv2d: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - - - // struct ggml_tensor* wino_res = ggml_conv_2d_3x3(ctx0, model.a, model.b); - // ggml_set_name(wino_res, "wino_res"); - // ggml_build_forward_expand(gf, wino_res); - // ne = wino_res->ne; - // printf("wino: (%zu, %zu, %zu, %zu) \n", ne[0], ne[1], ne[2], ne[3]); - ggml_free(ctx0); - return gf; -} - - - -std::vector compute_graph(const test_model & model, ggml_gallocr_t allocr, - build_graph_t build_graph, int iters, double *t) { - struct ggml_cgraph * gf = build_graph(model); - - - // allocate tensors - ggml_gallocr_alloc_graph(allocr, gf); - int n_threads = 1; - - if (ggml_backend_is_cpu(model.backend)) { - ggml_backend_cpu_set_n_threads(model.backend, n_threads); - } - -#ifdef GGML_USE_METAL - if (ggml_backend_is_metal(model.backend)) { - ggml_backend_metal_set_n_cb(model.backend, n_threads); - } -#endif - - ggml_backend_synchronize(model.backend); - - ggml_backend_graph_compute(model.backend, gf); - - ggml_backend_synchronize(model.backend); - - int64_t start_time = ggml_time_us(); - - for(int iter=0; iter data(ggml_nelements(res)); - std::vector fdata(ggml_nelements(res)); - std::vector data(ggml_nelements(res)); - ggml_backend_tensor_get(res, fdata.data(), 0, ggml_nbytes(res)); - ggml_fp16_to_fp32_row(fdata.data(), data.data(), ggml_nelements(res)); - *t = time_us/1000; - return data; - -} - - -int main(void) -{ - ggml_time_init(); - std::vector> configs = { - // std::make_tuple(64,64,48,64,3,3), - // std::make_tuple(320,320,104,152,3,3), - // std::make_tuple(640,640,52,76,3,3), - // std::make_tuple(640,640,104,152,3,3), - // std::make_tuple(960,320,104,152,3,3), - // std::make_tuple(1,128,38,49,3,3), - std::make_tuple(1,1,38,49,38,49), - // std::make_tuple(1280,1280,26,38,1,1), - // std::make_tuple(256,128,768,1024,3,3), - // std::make_tuple(256,128,768,1024,1,1), - // std::make_tuple(1280,640,52,76,3,3), - // std::make_tuple(1920,1280,26,38,3,3), - // std::make_tuple(2560,1280,26,38,3,3), - // std::make_tuple(512,512,104,152,3,3), - // std::make_tuple(512,512,208,304,3,3), - // std::make_tuple(512,256,416,608,3,3), - // std::make_tuple(256,128,832,1216,3,3), - // std::make_tuple(256,256,832,1216,3,3), - // std::make_tuple(320,256,1024,1920) - }; - - int k = 0; - - for (auto c : configs){ - test_model model; - load_model(model, std::get<0>(c), std::get<1>(c), std::get<2>(c), - std::get<3>(c), std::get<4>(c), std::get<5>(c), true); - - ggml_gallocr_t allocr = NULL; - allocr = ggml_gallocr_new(ggml_backend_get_default_buffer_type(model.backend)); - - //create the worst case graph for memory usage estimation - struct ggml_cgraph * gf = build_graph_0(model); - - // compute the required memory - ggml_gallocr_reserve(allocr, gf); - size_t mem_size0 = ggml_gallocr_get_buffer_size(allocr, 0); - // fprintf(stderr, "%s: compute buffer size: %.2f MB\n", __func__, mem_size/1024.0f/1024.0f); - - - struct ggml_cgraph * gf_res_0 = NULL; - int iterations = 0; - - double run_time0; - std::vector im2col_data = compute_graph(model, allocr, build_graph_0, iterations, &run_time0); - - - - - - - - //create the worst case graph for memory usage estimation - - - - - - - - // for(int i = 0; i < ggml_nelements(wino_res); i++) { - // for(int i = 0; i < 26*38; i++) { - // for(int i = 0; i < std::get<2>(c); i++) { - // // float diff = fabs(conv2d_data[i] - wino_data[i]); - // for(int j = 0; j < std::get<3>(c); j++) { - // printf("%4.1f, ", im2col_data[i*std::get<3>(c)+j]); - // } - // printf("\n"); - // } - for(int i = 0; i < std::get<4>(c); i++) { - // float diff = fabs(conv2d_data[i] - wino_data[i]); - for(int j = 0; j < std::get<5>(c); j++) { - printf("%4.1f, ", im2col_data[i*std::get<5>(c)+j]); - } - printf("\n"); - } - - ggml_free(model.ctx); - ggml_backend_buffer_free(model.buffer); - ggml_backend_free(model.backend); - ggml_gallocr_free(allocr); - - } - - // printf("\nPerforming test:\n"); - return 0; -}