From 26ef7cdab87216015dfdd84a70b0ca05f3e531f1 Mon Sep 17 00:00:00 2001 From: Jeremy L Thompson Date: Tue, 22 Oct 2024 11:27:04 -0600 Subject: [PATCH] debug - truncate jit output --- backends/cuda-gen/ceed-cuda-gen-operator-build.cpp | 3 --- backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp | 4 ---- backends/cuda/ceed-cuda-compile.cpp | 6 +++--- backends/hip-gen/ceed-hip-gen-operator-build.cpp | 4 ---- backends/hip-ref/ceed-hip-ref-qfunction-load.cpp | 4 ---- backends/hip-ref/ceed-hip-ref-restriction.c | 1 - backends/hip/ceed-hip-compile.cpp | 6 +++--- interface/ceed-jit-tools.c | 2 -- 8 files changed, 6 insertions(+), 24 deletions(-) diff --git a/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp b/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp index 315db3844f..3019dde777 100644 --- a/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp +++ b/backends/cuda-gen/ceed-cuda-gen-operator-build.cpp @@ -896,9 +896,6 @@ extern "C" int CeedOperatorBuildKernel_Cuda_gen(CeedOperator op) { code << "// -----------------------------------------------------------------------------\n\n"; // View kernel for debugging - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated Operator Kernels:\n"); - CeedDebug(ceed, code.str().c_str()); - CeedCallBackend(CeedCompile_Cuda(ceed, code.str().c_str(), &data->module, 1, "T_1D", CeedIntMax(Q_1d, data->max_P_1d))); CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, operator_name.c_str(), &data->op)); diff --git a/backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp b/backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp index 2d5540ead8..0ef3dde204 100644 --- a/backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp +++ b/backends/cuda-ref/ceed-cuda-ref-qfunction-load.cpp @@ -105,10 +105,6 @@ extern "C" int CeedQFunctionBuildKernel_Cuda_ref(CeedQFunction qf) { code << " }\n"; code << "}\n"; - // View kernel for debugging - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated QFunction Kernels:\n"); - CeedDebug(ceed, code.str().c_str()); - // Compile kernel CeedCallBackend(CeedCompile_Cuda(ceed, code.str().c_str(), &data->module, 0)); CeedCallBackend(CeedGetKernel_Cuda(ceed, data->module, kernel_name.c_str(), &data->QFunction)); diff --git a/backends/cuda/ceed-cuda-compile.cpp b/backends/cuda/ceed-cuda-compile.cpp index c6aafbcbf6..20c57db2e8 100644 --- a/backends/cuda/ceed-cuda-compile.cpp +++ b/backends/cuda/ceed-cuda-compile.cpp @@ -121,6 +121,9 @@ int CeedCompile_Cuda(Ceed ceed, const char *source, CUmodule *module, const Ceed CeedCallNvrtc(ceed, nvrtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); // Compile kernel + CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- ATTEMPTING TO COMPILE JIT SOURCE ----------\n"); + CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); + CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- END OF JIT SOURCE ----------\n"); nvrtcResult result = nvrtcCompileProgram(prog, num_opts + num_jit_source_dirs + num_jit_defines, opts); for (CeedInt i = 0; i < num_jit_source_dirs; i++) { @@ -134,9 +137,6 @@ int CeedCompile_Cuda(Ceed ceed, const char *source, CUmodule *module, const Ceed char *log; size_t log_size; - CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); - CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); - CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); CeedCallNvrtc(ceed, nvrtcGetProgramLogSize(prog, &log_size)); CeedCallBackend(CeedMalloc(log_size, &log)); CeedCallNvrtc(ceed, nvrtcGetProgramLog(prog, log)); diff --git a/backends/hip-gen/ceed-hip-gen-operator-build.cpp b/backends/hip-gen/ceed-hip-gen-operator-build.cpp index f1a876ce26..05b4e27c24 100644 --- a/backends/hip-gen/ceed-hip-gen-operator-build.cpp +++ b/backends/hip-gen/ceed-hip-gen-operator-build.cpp @@ -896,10 +896,6 @@ extern "C" int CeedOperatorBuildKernel_Hip_gen(CeedOperator op) { code << "}\n"; code << "// -----------------------------------------------------------------------------\n\n"; - // View kernel for debugging - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated Operator Kernels:\n"); - CeedDebug(ceed, code.str().c_str()); - CeedInt block_sizes[3] = {0, 0, 0}; CeedInt num_elem; diff --git a/backends/hip-ref/ceed-hip-ref-qfunction-load.cpp b/backends/hip-ref/ceed-hip-ref-qfunction-load.cpp index 2311f8a332..0e58932439 100644 --- a/backends/hip-ref/ceed-hip-ref-qfunction-load.cpp +++ b/backends/hip-ref/ceed-hip-ref-qfunction-load.cpp @@ -109,10 +109,6 @@ extern "C" int CeedQFunctionBuildKernel_Hip_ref(CeedQFunction qf) { code << " }\n"; code << "}\n"; - // View kernel for debugging - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Generated QFunction Kernels:\n"); - CeedDebug(ceed, code.str().c_str()); - // Compile kernel CeedCallBackend(CeedCompile_Hip(ceed, code.str().c_str(), &data->module, 1, "BLOCK_SIZE", ceed_Hip->opt_block_size)); CeedCallBackend(CeedGetKernel_Hip(ceed, data->module, kernel_name.c_str(), &data->QFunction)); diff --git a/backends/hip-ref/ceed-hip-ref-restriction.c b/backends/hip-ref/ceed-hip-ref-restriction.c index ca1d19d7a6..7642e09376 100644 --- a/backends/hip-ref/ceed-hip-ref-restriction.c +++ b/backends/hip-ref/ceed-hip-ref-restriction.c @@ -71,7 +71,6 @@ static inline int CeedElemRestrictionSetupCompile_Hip(CeedElemRestriction rstr) "// AtPoints restriction source\n#include \n\n" "// Standard restriction source\n#include \n"; - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "----- Loading Restriction Kernel Source Complete! -----\n"); CeedCallBackend(CeedCompile_Hip(ceed, restriction_kernel_source, &impl->module, 6, "RSTR_ELEM_SIZE", elem_size, "RSTR_NUM_ELEM", num_elem, "RSTR_NUM_COMP", num_comp, "RSTR_NUM_NODES", impl->num_nodes, "RSTR_COMP_STRIDE", comp_stride, "USE_DETERMINISTIC", is_deterministic ? 1 : 0)); diff --git a/backends/hip/ceed-hip-compile.cpp b/backends/hip/ceed-hip-compile.cpp index 20f2eb0e2a..dface44ef6 100644 --- a/backends/hip/ceed-hip-compile.cpp +++ b/backends/hip/ceed-hip-compile.cpp @@ -123,6 +123,9 @@ int CeedCompile_Hip(Ceed ceed, const char *source, hipModule_t *module, const Ce CeedCallHiprtc(ceed, hiprtcCreateProgram(&prog, code.str().c_str(), NULL, 0, NULL, NULL)); // Compile kernel + CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- ATTEMPTING TO COMPILE JIT SOURCE ----------\n"); + CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); + CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- END OF JIT SOURCE ----------\n"); hiprtcResult result = hiprtcCompileProgram(prog, num_opts + num_jit_source_dirs + num_jit_defines, opts); for (CeedInt i = 0; i < num_jit_source_dirs; i++) { @@ -136,9 +139,6 @@ int CeedCompile_Hip(Ceed ceed, const char *source, hipModule_t *module, const Ce size_t log_size; char *log; - CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); - CeedDebug(ceed, "Source:\n%s\n", code.str().c_str()); - CeedDebug256(ceed, CEED_DEBUG_COLOR_ERROR, "---------- CEED JIT SOURCE FAILED TO COMPILE ----------\n"); CeedChk_hiprtc(ceed, hiprtcGetProgramLogSize(prog, &log_size)); CeedCallBackend(CeedMalloc(log_size, &log)); CeedCallHiprtc(ceed, hiprtcGetProgramLog(prog, log)); diff --git a/interface/ceed-jit-tools.c b/interface/ceed-jit-tools.c index d89f713dfe..14f1babb87 100644 --- a/interface/ceed-jit-tools.c +++ b/interface/ceed-jit-tools.c @@ -130,8 +130,6 @@ int CeedLoadSourceToInitializedBuffer(Ceed ceed, const char *source_file_path, C CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "---------- Ceed JiT ----------\n"); CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Current source file: "); CeedDebug(ceed, "%s\n", source_file_path); - CeedDebug256(ceed, CEED_DEBUG_COLOR_SUCCESS, "Current buffer:\n"); - CeedDebug(ceed, "%s\n", *buffer); // Read file to temporary buffer source_file = fopen(source_file_path, "rb");