//+------------------------------------------------------------------+ //| BenchSubrangeAB_AllInOne.mq5 | //| Foco: SOLO el caso interesante -> m_reconstruir=true con subrango | //| real (fs != total_cols), comparando: | //| CPU (patron validado: out=X + overwrite del rango) | //| GPU-A (mandar TODO + offsets, leer completo, sin reconstruccion | //| manual porque el resto nunca se toco) | //| GPU-B (mandar SOLO el rango, reconstruir con out=X + overwrite) | //| | //| El caso fs==total_cols y el caso !reconstruir NO se benchmarkean | //| aca a proposito: ya sabemos que van por la ruta trivial (0 copias | //| intermedias, sin ambiguedad de estrategia). | //+------------------------------------------------------------------+ #property script_show_inputs #ifndef SCALERBYLEO_CL_DEVICE #define SCALERBYLEO_CL_DEVICE (CL_USE_GPU_ONLY) #endif //+------------------------------------------------------------------+ static const string MinMaxKernelSrcTemplate = "#define MATRIX_IDX(row, col, cols) ((row) * (cols) + (col)) \r\n" "\r\n" "__kernel void max_min_scaler_transform_matrix( \r\n" " __global const real_t* min_values, \r\n" " __global const real_t* diff_inv, \r\n" " __global real_t* matrix, \r\n" " const uint cols, \r\n" " const uint istart \r\n" ") \r\n" "{ \r\n" " const uint row = get_global_id(0); \r\n" " const uint col = get_global_id(1); \r\n" " const uint matrix_idx = MATRIX_IDX(row, istart + col, cols); \r\n" " matrix[matrix_idx] = (matrix[matrix_idx] - min_values[col]) * diff_inv[col]; \r\n" "} \r\n" "\r\n" "__kernel void max_min_scaler_transform_matrix_exact( \r\n" " __global const real_t* min_values, \r\n" " __global const real_t* diff_inv, \r\n" " __global real_t* matrix, \r\n" " const uint cols \r\n" ") \r\n" "{ \r\n" " const uint row = get_global_id(0); \r\n" " const uint col = get_global_id(1); \r\n" " const uint matrix_idx = MATRIX_IDX(row, col, cols); \r\n" " matrix[matrix_idx] = (matrix[matrix_idx] - min_values[col]) * diff_inv[col]; \r\n" "} \r\n"; //+------------------------------------------------------------------+ class CCLCtxShared { private: static int s_ref_count; public: static int s_context; static bool ContextInit(void) { if(s_context != INVALID_HANDLE) { s_ref_count++; return true; } ::ResetLastError(); if((s_context = CLContextCreate(SCALERBYLEO_CL_DEVICE)) == INVALID_HANDLE) return false; s_ref_count = 1; return true; } static void ContextRelease(void) { if(s_context == INVALID_HANDLE) return; if(--s_ref_count <= 0) { CLContextFree(s_context); s_context = INVALID_HANDLE; s_ref_count = 0; } } }; int CCLCtxShared::s_context = INVALID_HANDLE; int CCLCtxShared::s_ref_count = 0; //+------------------------------------------------------------------+ template class CCLProgramManager { public: static int s_program; static int s_kernel_full; static int s_kernel_exact; static bool s_ready; static bool Init(void) { if(s_ready) return true; if(!CCLCtxShared::ContextInit()) return false; const string real_t_def = (sizeof(TFlt) == sizeof(float)) ? "#define real_t float \r\n" : "#define real_t double \r\n"; const string src = real_t_def + MinMaxKernelSrcTemplate; string build_log; ::ResetLastError(); if((s_program = CLProgramCreate(CCLCtxShared::s_context, src, build_log)) == INVALID_HANDLE) { Print("CLProgramCreate failed: ", build_log, " err=", GetLastError()); CCLCtxShared::ContextRelease(); return false; } if((s_kernel_full = CLKernelCreate(s_program, "max_min_scaler_transform_matrix")) == INVALID_HANDLE) { Print("CLKernelCreate(full) failed, err=", GetLastError()); return false; } if((s_kernel_exact = CLKernelCreate(s_program, "max_min_scaler_transform_matrix_exact")) == INVALID_HANDLE) { Print("CLKernelCreate(exact) failed, err=", GetLastError()); return false; } s_ready = true; return true; } static void Deinit(void) { if(!s_ready) return; if(s_kernel_full != INVALID_HANDLE) { CLKernelFree(s_kernel_full); s_kernel_full = INVALID_HANDLE; } if(s_kernel_exact != INVALID_HANDLE) { CLKernelFree(s_kernel_exact); s_kernel_exact = INVALID_HANDLE; } if(s_program != INVALID_HANDLE) { CLProgramFree(s_program); s_program = INVALID_HANDLE; } CCLCtxShared::ContextRelease(); s_ready = false; } }; template int CCLProgramManager::s_program = INVALID_HANDLE; template int CCLProgramManager::s_kernel_full = INVALID_HANDLE; template int CCLProgramManager::s_kernel_exact = INVALID_HANDLE; template bool CCLProgramManager::s_ready = false; //+------------------------------------------------------------------+ //| MinMax scaler: ruta trivial + dos estrategias para subrango real | //+------------------------------------------------------------------+ template class MaxMinScalerCL { private: bool m_loaded_scaler; ulong m_i_start; vector m_min_vals; vector m_diff_inv; int m_buffer_h_min_vals; int m_buffer_h_diff_inv; bool EnsureAuxBuffers(void) { const uint bytes = uint(m_min_vals.Size() * sizeof(TFlt)); if(m_buffer_h_min_vals == INVALID_HANDLE) if((m_buffer_h_min_vals = CLBufferCreate(CCLCtxShared::s_context, bytes, CL_MEM_READ_ONLY)) == INVALID_HANDLE) return false; if(!CLBufferWrite(m_buffer_h_min_vals, 0, m_min_vals)) return false; if(m_buffer_h_diff_inv == INVALID_HANDLE) if((m_buffer_h_diff_inv = CLBufferCreate(CCLCtxShared::s_context, bytes, CL_MEM_READ_ONLY)) == INVALID_HANDLE) return false; if(!CLBufferWrite(m_buffer_h_diff_inv, 0, m_diff_inv)) return false; return true; } public: MaxMinScalerCL(void) : m_loaded_scaler(false), m_i_start(0), m_buffer_h_min_vals(INVALID_HANDLE), m_buffer_h_diff_inv(INVALID_HANDLE) {} ~MaxMinScalerCL(void) { if(m_buffer_h_min_vals != INVALID_HANDLE) CLBufferFree(m_buffer_h_min_vals); if(m_buffer_h_diff_inv != INVALID_HANDLE) CLBufferFree(m_buffer_h_diff_inv); } bool Fit(const matrix &X, const ulong i_start, const ulong fs) { if(!CCLProgramManager::Init()) return false; m_i_start = i_start; m_min_vals.Resize(fs); m_diff_inv.Resize(fs); for(ulong c = 0; c < fs; c++) { TFlt mn = X[0][i_start + c]; TFlt mx = X[0][i_start + c]; for(ulong r = 1; r < X.Rows(); r++) { const TFlt v = X[r][i_start + c]; if(v < mn) mn = v; if(v > mx) mx = v; } const TFlt diff = (mx - mn); m_min_vals[c] = mn; m_diff_inv[c] = (diff == 0) ? TFlt(0) : TFlt(1.0 / double(diff)); } m_loaded_scaler = true; return EnsureAuxBuffers(); } //--- SUBRANGO ESTRATEGIA A: mandar TODO + offsets, leer completo directo (0 copias CPU extra) bool SubrangeStrategyA(matrix &X, matrix &out) { if(!m_loaded_scaler) return false; const ulong rows = X.Rows(); const ulong total_cols = X.Cols(); const ulong fs = m_min_vals.Size(); int matrix_h = INVALID_HANDLE; const uint bytes = uint(sizeof(TFlt) * total_cols * rows); matrix_h = CLBufferCreate(CCLCtxShared::s_context, bytes, CL_MEM_READ_WRITE); if(matrix_h == INVALID_HANDLE) return false; bool ok = CLBufferWrite(matrix_h, 0, X); if(ok) ok = CLSetKernelArgMem(CCLProgramManager::s_kernel_full, 0, m_buffer_h_min_vals) && CLSetKernelArgMem(CCLProgramManager::s_kernel_full, 1, m_buffer_h_diff_inv) && CLSetKernelArgMem(CCLProgramManager::s_kernel_full, 2, matrix_h) && CLSetKernelArg(CCLProgramManager::s_kernel_full, 3, uint(fs)) && CLSetKernelArg(CCLProgramManager::s_kernel_full, 4, uint(m_i_start)); if(ok) { const uint work_off[2] = {0, 0}; uint work_size[2] = { (uint)rows, (uint)fs }; ok = CLExecute(CCLProgramManager::s_kernel_full, 2, work_off, work_size); } if(ok) { out.Resize(rows, total_cols); ok = CLBufferRead(matrix_h, 0, out); } CLBufferFree(matrix_h); return ok; } //--- SUBRANGO ESTRATEGIA B: mandar SOLO el rango, reconstruir con out=X + overwrite bool SubrangeStrategyB(matrix &X, matrix &out) { if(!m_loaded_scaler) return false; const ulong rows = X.Rows(); const ulong fs = m_min_vals.Size(); matrix narrow; narrow.Resize(rows, fs); for(ulong c = 0; c < fs; c++) for(ulong r = 0; r < rows; r++) narrow[r][c] = X[r][m_i_start + c]; int matrix_h = INVALID_HANDLE; const uint bytes = uint(sizeof(TFlt) * fs * rows); matrix_h = CLBufferCreate(CCLCtxShared::s_context, bytes, CL_MEM_READ_WRITE); bool ok = (matrix_h != INVALID_HANDLE) && CLBufferWrite(matrix_h, 0, narrow); if(ok) ok = CLSetKernelArgMem(CCLProgramManager::s_kernel_exact, 0, m_buffer_h_min_vals) && CLSetKernelArgMem(CCLProgramManager::s_kernel_exact, 1, m_buffer_h_diff_inv) && CLSetKernelArgMem(CCLProgramManager::s_kernel_exact, 2, matrix_h) && CLSetKernelArg(CCLProgramManager::s_kernel_exact, 3, uint(fs)); if(ok) { const uint work_off[2] = {0, 0}; uint work_size[2] = { (uint)rows, (uint)fs }; ok = CLExecute(CCLProgramManager::s_kernel_exact, 2, work_off, work_size); } if(!ok) { CLBufferFree(matrix_h); return false; } matrix res; res.Resize(rows, fs); ok = CLBufferRead(matrix_h, 0, res); CLBufferFree(matrix_h); if(!ok) return false; out = X; for(ulong c = 0; c < fs; c++) for(ulong r = 0; r < rows; r++) out[r][m_i_start + c] = res[r][c]; return true; } }; //+------------------------------------------------------------------+ //| Version CPU pura (patron validado: out=X + overwrite del rango) | //+------------------------------------------------------------------+ template class MaxMinScalerCPU { private: ulong m_i_start; vector m_min_vals; vector m_diff_inv; public: MaxMinScalerCPU(void) : m_i_start(0) {} bool Fit(const matrix &X, const ulong i_start, const ulong fs) { m_i_start = i_start; m_min_vals.Resize(fs); m_diff_inv.Resize(fs); for(ulong c = 0; c < fs; c++) { TFlt mn = X[0][i_start + c]; TFlt mx = X[0][i_start + c]; for(ulong r = 1; r < X.Rows(); r++) { const TFlt v = X[r][i_start + c]; if(v < mn) mn = v; if(v > mx) mx = v; } const TFlt diff = (mx - mn); m_min_vals[c] = mn; m_diff_inv[c] = (diff == 0) ? TFlt(0) : TFlt(1.0 / double(diff)); } return true; } bool Transform(const matrix &X, matrix &out) { const ulong rows = X.Rows(); const ulong fs = m_min_vals.Size(); out = X; for(ulong c = 0; c < fs; c++) for(ulong r = 0; r < rows; r++) out[r][m_i_start + c] = TFlt((double(X[r][m_i_start + c]) - double(m_min_vals[c])) * double(m_diff_inv[c])); return true; } }; //+------------------------------------------------------------------+ //| Config del barrido: SOLO subrango real, fracciones chicas a 90% | //+------------------------------------------------------------------+ int g_sizes[] = {500, 2000, 5000, 9600}; double g_fs_fracs[] = {0.01, 0.05, 0.10, 0.25, 0.50, 0.75, 0.90}; int g_repeats = 5; //+------------------------------------------------------------------+ ulong RunOnce_CPU(matrix &X, const ulong i_start, const ulong fs) { MaxMinScalerCPU sc; sc.Fit(X, i_start, fs); matrix out; const ulong t0 = GetMicrosecondCount(); sc.Transform(X, out); return GetMicrosecondCount() - t0; } //+------------------------------------------------------------------+ ulong RunOnce_A(matrix &X, const ulong i_start, const ulong fs) { MaxMinScalerCL sc; sc.Fit(X, i_start, fs); matrix out; const ulong t0 = GetMicrosecondCount(); sc.SubrangeStrategyA(X, out); return GetMicrosecondCount() - t0; } //+------------------------------------------------------------------+ ulong RunOnce_B(matrix &X, const ulong i_start, const ulong fs) { MaxMinScalerCL sc; sc.Fit(X, i_start, fs); matrix out; const ulong t0 = GetMicrosecondCount(); sc.SubrangeStrategyB(X, out); return GetMicrosecondCount() - t0; } //+------------------------------------------------------------------+ ulong Median(ulong &arr[]) { const int n = ArraySize(arr); ArraySort(arr); return arr[n / 2]; } //+------------------------------------------------------------------+ void OnStart() { Print("=== Bench SUBRANGO REAL (reconstruir=true, fs != total_cols): CPU vs GPU-A vs GPU-B ==="); PrintFormat("%-6s %-6s %-6s %-12s %-12s %-12s %-10s %-10s", "size", "fs", "frac", "CPU(us)", "GPU_A(us)", "GPU_B(us)", "A_vs_CPU", "B_vs_CPU"); for(int si = 0; si < ArraySize(g_sizes); si++) { const int size = g_sizes[si]; matrix X; X.Resize(size, size); X.Random(0.0, 20.0); for(int fi = 0; fi < ArraySize(g_fs_fracs); fi++) { ulong fs = (ulong)MathMax(1, MathRound(size * g_fs_fracs[fi])); if(fs >= (ulong)size) fs = size - 1; // forzamos subrango real, nunca fs==total_cols aca const ulong i_start = 0; ulong cpu_t[], a_t[], b_t[]; ArrayResize(cpu_t, g_repeats); ArrayResize(a_t, g_repeats); ArrayResize(b_t, g_repeats); for(int rep = 0; rep < g_repeats; rep++) { cpu_t[rep] = RunOnce_CPU(X, i_start, fs); a_t[rep] = RunOnce_A(X, i_start, fs); b_t[rep] = RunOnce_B(X, i_start, fs); } const ulong cpu_med = Median(cpu_t); const ulong a_med = Median(a_t); const ulong b_med = Median(b_t); const double a_ratio = (a_med > 0) ? (double)cpu_med / (double)a_med : 0.0; const double b_ratio = (b_med > 0) ? (double)cpu_med / (double)b_med : 0.0; PrintFormat("%-6d %-6d %-6.2f %-12I64u %-12I64u %-12I64u %-10.2f %-10.2f", size, fs, g_fs_fracs[fi], cpu_med, a_med, b_med, a_ratio, b_ratio); } } Print("=== Fin ==="); Print("Ratio > 1.0 = GPU mas rapida que CPU. Comparar tambien GPU_A vs GPU_B directamente en la tabla."); CCLProgramManager::Deinit(); } //+------------------------------------------------------------------+