From 66ae6d6be881527d01626302b1c5f4fde23d9d61 Mon Sep 17 00:00:00 2001 From: slaren Date: Wed, 29 Nov 2023 18:32:01 +0100 Subject: [PATCH] update im2col test params, show callstack with GGML_ASSERT on CUDA failures --- src/ggml-cuda.cu | 6 +++--- tests/test-backend-ops.cpp | 32 ++++++++++++++++++-------------- 2 files changed, 21 insertions(+), 17 deletions(-) diff --git a/src/ggml-cuda.cu b/src/ggml-cuda.cu index 2d377ed01..9881dd346 100644 --- a/src/ggml-cuda.cu +++ b/src/ggml-cuda.cu @@ -189,7 +189,7 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size"); fprintf(stderr, "\nCUDA error %d at %s:%d: %s\n", err_, __FILE__, __LINE__, \ cudaGetErrorString(err_)); \ fprintf(stderr, "current device: %d\n", id); \ - exit(1); \ + GGML_ASSERT(!"CUDA error"); \ } \ } while (0) @@ -203,7 +203,7 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size"); fprintf(stderr, "\ncuBLAS error %d at %s:%d: %s\n", \ err_, __FILE__, __LINE__, cublasGetStatusString(err_)); \ fprintf(stderr, "current device: %d\n", id); \ - exit(1); \ + GGML_ASSERT(!"cuBLAS error"); \ } \ } while (0) #else @@ -215,7 +215,7 @@ static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size"); cudaGetDevice(&id); \ fprintf(stderr, "\ncuBLAS error %d at %s:%d\n", err_, __FILE__, __LINE__); \ fprintf(stderr, "current device: %d\n", id); \ - exit(1); \ + GGML_ASSERT(!"cuBLAS error"); \ } \ } while (0) #endif // CUDART_VERSION >= 11 diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 224d4a0eb..f7d81dc0d 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -699,35 +699,39 @@ struct test_alibi : public test_case { // GGML_OP_IM2COL struct test_im2col : public test_case { - const ggml_type type_a; - const ggml_type type_b; - const std::array ne_a; - const std::array ne_b; + const ggml_type type_input; + const ggml_type type_kernel; + const std::array ne_input; + const std::array ne_kernel; + // stride const int s0; const int s1; + // padding const int p0; const int p1; + // dilatation const int d0; const int d1; + // mode const bool is_2D; std::string vars() override { - return VARS_TO_STR11(type_a, type_b, ne_a, ne_b, s0, s1, p0, p1, d0, d1, is_2D); + return VARS_TO_STR11(type_input, type_kernel, ne_input, ne_kernel, s0, s1, p0, p1, d0, d1, is_2D); } - test_im2col(ggml_type type_a = GGML_TYPE_F16, ggml_type type_b = GGML_TYPE_F32, - std::array ne_a = {10, 10, 10, 10}, - std::array ne_b = {10, 10, 10, 10}, + test_im2col(ggml_type type_input = GGML_TYPE_F32, ggml_type type_kernel = GGML_TYPE_F16, + std::array ne_input = {10, 10, 3, 1}, // [input_width, input_height, input_channels, 1] + std::array ne_kernel = {3, 3, 3, 1}, // [kernel_width, kernel_height, input_channels, 1] int s0 = 1, int s1 = 1, - int p0 = 0, int p1 = 0, + int p0 = 1, int p1 = 1, int d0 = 1, int d1 = 1, - bool is_2D = false) - : type_a(type_a), type_b(type_b), ne_a(ne_a), ne_b(ne_b), s0(s0), s1(s1), p0(p0), p1(p1), d0(d0), d1(d1), is_2D(is_2D) {} + bool is_2D = true) + : type_input(type_input), type_kernel(type_kernel), ne_input(ne_input), ne_kernel(ne_kernel), s0(s0), s1(s1), p0(p0), p1(p1), d0(d0), d1(d1), is_2D(is_2D) {} ggml_tensor * build_graph(ggml_context * ctx) override { - ggml_tensor * a = ggml_new_tensor(ctx, type_a, 4, ne_a.data()); - ggml_tensor * b = ggml_new_tensor(ctx, type_b, 4, ne_b.data()); - ggml_tensor * out = ggml_im2col(ctx, a, b, s0, s1, p0, p1, d0, d1, is_2D); + ggml_tensor * input = ggml_new_tensor(ctx, type_input, 4, ne_input.data()); + ggml_tensor * kernel = ggml_new_tensor(ctx, type_kernel, 4, ne_kernel.data()); + ggml_tensor * out = ggml_im2col(ctx, kernel, input, s0, s1, p0, p1, d0, d1, is_2D); return out; } };