################################################################################ # # This program and the accompanying materials are made available under the # terms of the Apache License, Version 2.0 which is available at # https://www.apache.org/licenses/LICENSE-2.0. # # See the NOTICE file distributed with this work for additional # information regarding copyright ownership. # Unless required by applicable law or agreed to in writing, software # distributed under the License is distributed on an "AS IS" BASIS, WITHOUT # WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the # License for the specific language governing permissions and limitations # under the License. # # SPDX-License-Identifier: Apache-2.0 ################################################################################ # TemplateProcessing.cmake - Complete Rewrite with Comprehensive Type Handling # This version generates ALL type variants to prevent undefined references # # ═══════════════════════════════════════════════════════════════ # BINARY SIZE WARNING: Template instantiations create MASSIVE binaries # ═══════════════════════════════════════════════════════════════ # # MEASURED SIZES: # - Normal build (no functrace): ~200MB # - Functrace build (SD_GCC_FUNCTRACE=ON): ~3.3GB (!) # # ROOT CAUSE: Template instantiations for ALL type combinations # - Each operation × each type combination = thousands of instantiations # - With debug symbols + functrace: exceeds 2GB relocation limit # # HOW TO REDUCE BINARY SIZE: # 1. **Disable functrace**: -Dlibnd4j.calltrace=OFF (90% reduction) # 2. **Reduce types**: Use -Dlibnd4j.datatypes="float32;double;int32;int64" # 3. **Split files**: Chunk sizes below control file splitting # 4. **Selective rendering**: Enable SD_SELECTIVE_TYPES in types.h # # For debugging without functrace: Use GDB with core dumps (see docs) # ═══════════════════════════════════════════════════════════════ set(CUSTOMOPS_GENERIC_SOURCES "" CACHE INTERNAL "Template-generated source files") # Chunk sizing for functrace builds: adapt chunk sizes so total compile memory is roughly 60% of # system RAM, even when high parallelism (buildthreads ~14) is used. set(USE_MULTI_PASS_GENERATION "ON" CACHE STRING "Use multi-pass generation (ON/OFF/AUTO)") if(SD_GCC_FUNCTRACE) set(_functrace_jobs 0) if(DEFINED CMAKE_BUILD_PARALLEL_LEVEL AND NOT CMAKE_BUILD_PARALLEL_LEVEL STREQUAL "") set(_functrace_jobs ${CMAKE_BUILD_PARALLEL_LEVEL}) elseif(DEFINED SD_PARALLEL_COMPILE_JOBS AND NOT SD_PARALLEL_COMPILE_JOBS STREQUAL "0") set(_functrace_jobs ${SD_PARALLEL_COMPILE_JOBS}) endif() if(_functrace_jobs LESS 1) include(ProcessorCount) ProcessorCount(_functrace_detected_jobs) if(_functrace_detected_jobs GREATER 0) set(_functrace_jobs ${_functrace_detected_jobs}) else() set(_functrace_jobs 4) endif() endif() cmake_host_system_information(RESULT _functrace_total_kb QUERY TOTAL_PHYSICAL_MEMORY) if(_functrace_total_kb) math(EXPR _functrace_total_mb "${_functrace_total_kb} / 1024") else() set(_functrace_total_mb 0) endif() if(_functrace_total_mb LESS 1024) set(_functrace_total_mb 1024) endif() # PERFORMANCE OPTIMIZATION: Use 75% of RAM for faster builds (increased from 60%) # Assumes system has adequate swap if needed math(EXPR _functrace_budget_mb "${_functrace_total_mb} * 3 / 4") math(EXPR _per_job_mb "${_functrace_budget_mb} / ${_functrace_jobs}") if(_per_job_mb LESS 512) set(_per_job_mb 512) endif() # Empirical: chunk target 8 → ~225MB, so ~28MB per instantiation slot. # PERFORMANCE OPTIMIZATION: Increase chunk sizes to reduce file count # OLD: Limited to 6-48 per chunk # NEW: Allow up to 96 per chunk for faster builds (fewer files = less overhead) math(EXPR _chunk_target "${_per_job_mb} / 28") if(_chunk_target LESS 12) set(_chunk_target 12) # Increased from 6 elseif(_chunk_target GREATER 96) set(_chunk_target 96) # Increased from 48 endif() # Direct instantiation chunks - also increased math(EXPR _direct_chunk "${_chunk_target} * 13 / 20") if(_direct_chunk LESS 8) set(_direct_chunk 8) # Increased from 4 endif() if(_direct_chunk GREATER _chunk_target) set(_direct_chunk ${_chunk_target}) endif() set(CHUNK_TARGET_INSTANTIATIONS "${_chunk_target}" CACHE STRING "Adaptive chunk size for functrace builds" FORCE) set(MULTI_PASS_CHUNK_SIZE "${_direct_chunk}" CACHE STRING "Adaptive direct chunk size for functrace builds" FORCE) if(NOT DEFINED CHUNK_MAX_INSTANTIATIONS OR CHUNK_MAX_INSTANTIATIONS LESS CHUNK_TARGET_INSTANTIATIONS) set(CHUNK_MAX_INSTANTIATIONS "${CHUNK_TARGET_INSTANTIATIONS}" CACHE STRING "Maximum template instantiations per chunk file (aligned with target)" FORCE) endif() message(STATUS "Functrace chunk sizing: target=${CHUNK_TARGET_INSTANTIATIONS}, direct=${MULTI_PASS_CHUNK_SIZE}, jobs=${_functrace_jobs}, per-job-limit≈${_per_job_mb}MB (60% of ${_functrace_total_mb}MB)") message(STATUS "⚠️ WARNING: Functrace builds still create ~3GB binaries (may exceed linker limits)") else() if(NOT DEFINED CHUNK_TARGET_INSTANTIATIONS) set(CHUNK_TARGET_INSTANTIATIONS "5" CACHE STRING "Target template instantiations per chunk file (1-20)") endif() if(NOT DEFINED CHUNK_MAX_INSTANTIATIONS) set(CHUNK_MAX_INSTANTIATIONS "10" CACHE STRING "Maximum template instantiations per chunk file") endif() if(NOT DEFINED MULTI_PASS_CHUNK_SIZE) set(MULTI_PASS_CHUNK_SIZE "20" CACHE STRING "Chunk size for direct instantiation files") endif() endif() # Enable selective rendering diagnostics in debug builds if(CMAKE_BUILD_TYPE STREQUAL "Debug") set(SRCORE_ENABLE_DIAGNOSTICS ON) endif() # ============================================================================ # COMPREHENSIVE TYPE EQUIVALENCE SYSTEM # ============================================================================ # Get ALL type variants for a given type - returns complete equivalence class function(get_all_type_variants type all_variants_var) set(all_variants "") string(STRIP "${type}" normalized_type) # Special handling for float16/half types if("${normalized_type}" MATCHES "float16|half|__half") # Only use float16 in all cases - never use __half or half as those are CUDA-specific set(all_variants "float16") set(${all_variants_var} "${all_variants}" PARENT_SCOPE) return() endif() # 8-bit signed integers set(INT8_CLASS "int8_t;SignedChar;signed char;char") foreach(variant ${INT8_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${INT8_CLASS}) break() endif() endforeach() # 8-bit unsigned integers set(UINT8_CLASS "uint8_t;UnsignedChar;unsigned char") if(NOT all_variants) foreach(variant ${UINT8_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${UINT8_CLASS}) break() endif() endforeach() endif() # 16-bit signed integers set(INT16_CLASS "int16_t;short;short int;signed short") if(NOT all_variants) foreach(variant ${INT16_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${INT16_CLASS}) break() endif() endforeach() endif() # 16-bit unsigned integers set(UINT16_CLASS "uint16_t;unsigned short;unsigned short int") if(NOT all_variants) foreach(variant ${UINT16_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${UINT16_CLASS}) break() endif() endforeach() endif() # 32-bit signed integers if(CMAKE_SIZEOF_VOID_P EQUAL 8) set(INT32_CLASS "int32_t;int;signed int;Int32Type") else() set(INT32_CLASS "int32_t;int;signed int;long;Int32Type") endif() if(NOT all_variants) foreach(variant ${INT32_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${INT32_CLASS}) break() endif() endforeach() endif() # 32-bit unsigned integers if(CMAKE_SIZEOF_VOID_P EQUAL 8) set(UINT32_CLASS "uint32_t;unsigned int;unsigned") else() set(UINT32_CLASS "uint32_t;unsigned int;unsigned;unsigned long") endif() if(NOT all_variants) foreach(variant ${UINT32_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${UINT32_CLASS}) break() endif() endforeach() endif() # 64-bit signed integers if(CMAKE_SIZEOF_VOID_P EQUAL 8) set(INT64_CLASS "int64_t;long long;long;sd::LongType;LongType") else() set(INT64_CLASS "int64_t;long long;sd::LongType;LongType") endif() if(NOT all_variants) foreach(variant ${INT64_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${INT64_CLASS}) break() endif() endforeach() endif() # 64-bit unsigned integers if(CMAKE_SIZEOF_VOID_P EQUAL 8) set(UINT64_CLASS "uint64_t;unsigned long long;unsigned long;sd::UnsignedLong;UnsignedLong;size_t") else() set(UINT64_CLASS "uLongType;unsigned long long;sd::UnsignedLong;UnsignedLong") endif() if(NOT all_variants) foreach(variant ${UINT64_CLASS}) if("${normalized_type}" STREQUAL "${variant}") set(all_variants ${UINT64_CLASS}) break() endif() endforeach() endif() # Non-aliased types if(NOT all_variants) if("${normalized_type}" STREQUAL "float") set(all_variants "float") elseif("${normalized_type}" STREQUAL "double") set(all_variants "double") elseif("${normalized_type}" STREQUAL "bfloat16") set(all_variants "bfloat16") elseif("${normalized_type}" STREQUAL "bool") set(all_variants "bool") else() # Unknown type - use as-is set(all_variants "${normalized_type}") endif() endif() # Remove duplicates if(all_variants) list(REMOVE_DUPLICATES all_variants) endif() set(${all_variants_var} "${all_variants}" PARENT_SCOPE) endfunction() # Get all types that should have cross-type conversions function(get_conversion_capable_types types_var) set(types "bool;int8_t;uint8_t;int16_t;uint16_t;int32_t;uint32_t;int64_t;uint64_t;float16;bfloat16;float;double") # Add common aliases that might appear list(APPEND types "SignedChar;UnsignedChar;short;unsigned short;int;unsigned int;long;unsigned long;long long;unsigned long long") set(${types_var} "${types}" PARENT_SCOPE) endfunction() # Deduplication helper function(add_unique_instantiation instantiation dedupe_set_var content_var) set(dedupe_set ${${dedupe_set_var}}) set(content "${${content_var}}") list(FIND dedupe_set "${instantiation}" found_idx) if(found_idx EQUAL -1) string(APPEND content "${instantiation}\n") list(APPEND dedupe_set "${instantiation}") endif() set(${dedupe_set_var} "${dedupe_set}" PARENT_SCOPE) set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # Memory configuration # ============================================================================ function(configure_memory_chunking) # CLEAN BUILD FIX: Chunking configuration is now set in Options.cmake based on build flags # This eliminates the need for CACHE FORCE which only worked for incremental builds # For clean builds, the correct values are set FROM THE START in Options.cmake # This function now only handles: # 1. Validation of the configured values # 2. AUTO mode for USE_MULTI_PASS_GENERATION # Validate chunking configuration message(STATUS "📊 Using chunking configuration: CHUNK_TARGET=${CHUNK_TARGET_INSTANTIATIONS}, MULTI_PASS=${MULTI_PASS_CHUNK_SIZE}") # Handle USE_MULTI_PASS_GENERATION AUTO mode if(USE_MULTI_PASS_GENERATION STREQUAL "AUTO") cmake_host_system_information(RESULT AVAILABLE_MEMORY QUERY AVAILABLE_PHYSICAL_MEMORY) if(AVAILABLE_MEMORY LESS 3000 OR DEFINED ENV{CI}) set(USE_MULTI_PASS_GENERATION ON CACHE STRING "Use multi-pass generation" FORCE) message(STATUS "Auto-detected low memory/CI: enabling multi-pass") else() set(USE_MULTI_PASS_GENERATION OFF CACHE STRING "Use multi-pass generation" FORCE) endif() endif() set(CHUNK_TARGET_INSTANTIATIONS_SET TRUE CACHE INTERNAL "Memory auto-detection completed") endfunction() # ============================================================================ # Main entry point # ============================================================================ function(execute_template_processing_with_selective_rendering) message(STATUS "🔧 Starting Template Processing with Comprehensive Type Handling...") # Step 1: Include selective rendering core include(${CMAKE_CURRENT_LIST_DIR}/SelectiveRenderingCore.cmake OPTIONAL RESULT_VARIABLE SR_FOUND) # Step 2: Setup selective rendering if available if(SR_FOUND) setup_selective_rendering_unified_safe() # Verify selective rendering worked if(NOT DEFINED UNIFIED_COMBINATIONS_2 OR NOT DEFINED UNIFIED_COMBINATIONS_3) message(WARNING "⚠️ Selective rendering failed to initialize, falling back to full matrix") setup_fallback_combinations() endif() else() message(STATUS "⚠️ Selective rendering not available, using full type matrix") setup_fallback_combinations() endif() # Step 3: Configure memory chunking configure_memory_chunking() # Step 4: Display filtering results if available if(DEFINED UNIFIED_COMBINATIONS_2 AND DEFINED UNIFIED_COMBINATIONS_3) list(LENGTH UNIFIED_COMBINATIONS_2 filtered_2_count) list(LENGTH UNIFIED_COMBINATIONS_3 filtered_3_count) message(STATUS "🎯 Selective Rendering Results:") message(STATUS " - 2-type combinations: ${filtered_2_count}") message(STATUS " - 3-type combinations: ${filtered_3_count}") endif() message(STATUS "🔧 Template processing: Multi-pass=${USE_MULTI_PASS_GENERATION}, Chunks=${CHUNK_TARGET_INSTANTIATIONS}") # Step 5: Execute template generation execute_template_generation() # Step 6: Propagate results to parent scope set(CUSTOMOPS_GENERIC_SOURCES ${CUSTOMOPS_GENERIC_SOURCES} PARENT_SCOPE) endfunction() # ============================================================================ # Template handlers with comprehensive type coverage # ============================================================================ # Handler for specials_single templates function(handle_specials_single t1 content_var) set(content "${${content_var}}") set(dedupe_set "") # Normalize type normalize_to_canonical_type("${t1}" norm_t1) # Skip redundant type aliases if(NOT t1 STREQUAL norm_t1) set(${content_var} "${content}" PARENT_SCOPE) return() endif() string(APPEND content "namespace sd {\n") # Instantiate only the specific member functions needed (class instantiation would duplicate these) add_unique_instantiation("template void sd::SpecialMethods<${t1}>::concatCpuGeneric(const std::vector&, NDArray&, const sd::LongType);" dedupe_set content) add_unique_instantiation("template void sd::SpecialMethods<${t1}>::sortGeneric(NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void sd::SpecialMethods<${t1}>::sortTadGeneric(NDArray*, sd::LongType*, int, bool);" dedupe_set content) string(APPEND content "}\n") set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for specials_double templates - includes cross-type conversions function(handle_specials_double t1 t2 content_var) set(content "${${content_var}}") set(dedupe_set "") # Map C++ types to their canonical form to detect same-type conversions normalize_to_canonical_type("${t1}" canonical_t1) normalize_to_canonical_type("${t2}" canonical_t2) # Skip if they're the same type after normalization if(canonical_t1 STREQUAL canonical_t2) # Same type - no TypeCast needed, only sorting operations string(APPEND content "namespace sd {\n") # DoubleMethods for sorting (same-type sorting is valid) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortByKey(sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortByValue(sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortTadByKey(sd::NDArray*, sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortTadByValue(sd::NDArray*, sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) string(APPEND content "}\n") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Rest of validation for actual type conversions... # [existing validation code] string(APPEND content "namespace sd {\n") # Only generate TypeCast for actual conversions - using sd::LongType as canonical form add_unique_instantiation("template void TypeCast::convertGeneric<${t1}, ${t2}>(void**, void*, sd::LongType, void*);" dedupe_set content) # DoubleMethods instantiations for sorting operations add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortByKey(sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortByValue(sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortTadByKey(sd::NDArray*, sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) add_unique_instantiation("template void DoubleMethods<${t1}, ${t2}>::sortTadByValue(sd::NDArray*, sd::NDArray*, sd::NDArray*, bool);" dedupe_set content) string(APPEND content "}\n") set(${content_var} "${content}" PARENT_SCOPE) endfunction() function(normalize_to_canonical_type cpp_type canonical_var) # Use a local variable to avoid modifying the input parameter directly set(local_cpp_type "${cpp_type}") # Strip any namespace prefixes (e.g., sd::) string(REGEX REPLACE "^[a-zA-Z0-9_]+::" "" local_cpp_type "${local_cpp_type}") # 1. DEFINE TYPE ALIASES # ======================== # Group all possible spellings for each canonical type in lists. set(types_int64 "int64_t" "long long" "LongType") set(types_uint64 "uint64_t" "unsigned long long" "UnsignedLong") set(types_int32 "int32_t" "int" "Int32Type" "signed int" "signed") set(types_uint32 "uint32_t" "unsigned int" "unsigned") set(types_int16 "int16_t" "short" "short int" "signed short" "signed short int") set(types_uint16 "uint16_t" "unsigned short" "unsigned short int") set(types_int8 "int8_t" "signed char" "char" "SignedChar") set(types_uint8 "uint8_t" "unsigned char" "UnsignedChar") set(types_float16 "float16" "half" "__half") set(types_bfloat16 "bfloat16") set(types_float "float") set(types_double "double") set(types_bool "bool") set(types_string "std::string" "utf8") set(types_u16string "std::u16string" "utf16") set(types_u32string "std::u32string" "utf32") # 2. HANDLE PLATFORM-DEPENDENT TYPES # ================================== # The 'long' type is 32-bit on 32-bit systems and some 64-bit systems (like Windows), # but 64-bit on most other 64-bit systems (like Linux). if(CMAKE_SIZEOF_VOID_P EQUAL 8) # For 64-bit architectures, assume 'long' is 64-bit. list(APPEND types_int64 "long") list(APPEND types_uint64 "unsigned long") else() # For 32-bit architectures, 'long' is 32-bit. list(APPEND types_int32 "long") list(APPEND types_uint32 "unsigned long") endif() # 3. BUILD REGEX PATTERNS FROM ALIASES # ====================================== # This creates a regex like "^(alias1|alias2|alias3)$" for each type list. foreach(type_name "int64" "uint64" "int32" "uint32" "int16" "uint16" "int8" "uint8" "float16" "bfloat16" "float" "double" "bool" "string" "u16string" "u32string") string(REPLACE ";" "|" patterns_${type_name} "${types_${type_name}}") set(patterns_${type_name} "^(${patterns_${type_name}})$") endforeach() # 4. MATCH AND SET CANONICAL TYPE # ================================= set(canonical "") if(local_cpp_type MATCHES "${patterns_int64}") set(canonical "sd::LongType") elseif(local_cpp_type MATCHES "${patterns_uint64}") set(canonical "sd::UnsignedLong") elseif(local_cpp_type MATCHES "${patterns_int32}") set(canonical "int32_t") elseif(local_cpp_type MATCHES "${patterns_uint32}") set(canonical "uint32_t") elseif(local_cpp_type MATCHES "${patterns_int16}") set(canonical "int16_t") elseif(local_cpp_type MATCHES "${patterns_uint16}") set(canonical "uint16_t") elseif(local_cpp_type MATCHES "${patterns_int8}") set(canonical "int8_t") elseif(local_cpp_type MATCHES "${patterns_uint8}") set(canonical "uint8_t") elseif(local_cpp_type MATCHES "${patterns_float16}") set(canonical "float16") elseif(local_cpp_type MATCHES "${patterns_bfloat16}") set(canonical "bfloat16") elseif(local_cpp_type MATCHES "${patterns_float}") set(canonical "float") elseif(local_cpp_type MATCHES "${patterns_double}") set(canonical "double") elseif(local_cpp_type MATCHES "${patterns_bool}") set(canonical "bool") elseif(local_cpp_type MATCHES "${patterns_string}") set(canonical "std::string") elseif(local_cpp_type MATCHES "${patterns_u16string}") set(canonical "std::u16string") elseif(local_cpp_type MATCHES "${patterns_u32string}") set(canonical "std::u32string") else() # Default to the original (namespace-stripped) type if no match is found set(canonical "${local_cpp_type}") endif() set(${canonical_var} "${canonical}" PARENT_SCOPE) endfunction() # Handler for broadcast_int templates function(handle_broadcast_int t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") foreach(type ${type_variants}) if(is_cuda) add_unique_instantiation("template class BroadcastInt<${type}>;" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::BroadcastInt<${type}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${type}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${type}>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() endforeach() set(${content_var} "${content}" PARENT_SCOPE) endfunction() function(handle_random t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize type normalize_to_canonical_type("${t1}" norm_t1) # Skip redundant type aliases if(NOT t1 STREQUAL norm_t1) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaSingle(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaDouble(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void const* vx, sd::LongType const* xShapeBuffer, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaTriple(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void const* vx, sd::LongType const* xShapeBuffer, void const* vy, sd::LongType const* yShapeBuffer, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) else() add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void const* x, sd::LongType const* xShapeBuffer, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void const* x, sd::LongType const* xShapeBuffer, void const* y, sd::LongType const* yShapeBuffer, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() function(handle_reduce_same t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Generate instantiations for the actual type used (including 'long') # Don't skip aliases because NativeOpExecutioner.cpp uses 'long' explicitly if(is_cuda) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execReduce(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, const sd::LongType* dXShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* vreductionBuffer, void* z, const sd::LongType* dZShapeInfo, const sd::LongType* hZShapeInfo, const sd::LongType* dims);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execReduceScalar(dim3 launchDims, cudaStream_t* stream, int opNum, const void* x, const sd::LongType* xShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* z, const sd::LongType* zShapeInfo, const sd::LongType* hZShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, void* reductionBuffer, const sd::LongType* tadOnlyShapeInfo);" dedupe_set content) add_unique_instantiation("template ${t1} functions::reduce::ReduceSameFunction<${t1}>::execScalarCudaLegacy(int opNum, void const* vx, sd::LongType const* xShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo);" dedupe_set content) else() add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::exec(int opNum, sd::memory::Workspace* workspace, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo, sd::LongType const* dimension);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) add_unique_instantiation("template ${t1} functions::reduce::ReduceSameFunction<${t1}>::execScalar(int opNum, const void* x, const sd::LongType* xShapeInfo, void* extraParams);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for broadcast_bool templates function(handle_broadcast_bool t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Only process if t2 is bool if(NOT "${t2}" STREQUAL "bool") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Normalize t1 normalize_to_canonical_type("${t1}" norm_t1) # Skip string types if(norm_t1 MATCHES "std::string|std::u16string|std::u32string") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) if(type1_enum STREQUAL "") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Validate pair _internal_srcore_is_valid_pair("${type1_enum}" "BOOL" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::execBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, void const *x, sd::LongType const *xShapeInfo, void const *y, sd::LongType const *yShapeInfo, void *result, sd::LongType const *resultShapeInfo, void *extraParams, sd::LongType *dimension, sd::LongType dimensionLength, sd::LongType const *tadOnlyShapeInfo, sd::LongType const *tadOffsets, sd::LongType const *tadOnlyShapeInfoZ, sd::LongType const *tadOffsetsZ);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::execBroadcast(dim3 launchDims, cudaStream_t *stream, const int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo, void *extraParams);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::execInverseBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, void *extraParams, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadOnlyShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, void *extraParams, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo, void *extraParams);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastBool<${t1}, bool>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, void *extraParams, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for broadcast (3-type) templates function(handle_broadcast_int t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") if(is_cuda) add_unique_instantiation("template class BroadcastInt<${t1}>;" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for pairwise templates # Handler for scalar templates function(handle_scalar t1 t2 t3 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) normalize_to_canonical_type("${t3}" norm_t3) # Skip any combination where ANY type isn't in its canonical form # This prevents duplicates like (unsigned short int, uint16_t, uint32_t) when (uint16_t, uint16_t, uint32_t) exists if(NOT (t1 STREQUAL norm_t1 AND t2 STREQUAL norm_t2 AND t3 STREQUAL norm_t3)) # At least one type is a non-canonical alias - skip to avoid duplicates set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) map_cpp_to_enum("${norm_t3}" type3_enum) # Validation _internal_srcore_is_valid_pair("${type1_enum}" "${type3_enum}" valid_13) _internal_srcore_is_valid_pair("${type2_enum}" "${type3_enum}" valid_23) _internal_srcore_is_valid_triple("${type1_enum}" "${type2_enum}" "${type3_enum}" valid_triple) if(NOT valid_13 OR NOT valid_23 OR NOT valid_triple) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations using NORMALIZED types (which are same as originals now due to check above) if(is_cuda) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::executeCudaShaped(dim3& launchDims, cudaStream_t* stream, int opNum, const void *vx, const sd::LongType *xShapeInfo, const sd::LongType *hxShapeInfo, void *vz, const sd::LongType *zShapeInfo, const sd::LongType *hzShapeInfo, const void *vscalar, void *vextraParams);" dedupe_set content) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::executeCudaAlongDimension(dim3& launchDims, cudaStream_t* stream, int opNum, const void *vx, const sd::LongType *xShapeInfo, void *vz, const sd::LongType *zShapeInfo, const void *vscalars, void *vextraParams, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) else() add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::transform(int opNum, const void *x, const sd::LongType *xShapeInfo, void *result, const sd::LongType *resultShapeInfo, const void *scalar, void *extraParams, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::transform(int opNum, const void *x, const sd::LongType *xShapeInfo, void *extraParams, void *z, const sd::LongType *zShapeInfo, const void *scalars, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetsZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for random templates function(handle_random t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize type normalize_to_canonical_type("${t1}" norm_t1) # Skip redundant type aliases if(NOT t1 STREQUAL norm_t1) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaSingle(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaDouble(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void const* vx, sd::LongType const* xShapeBuffer, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::executeCudaTriple(dim3& launchDims, cudaStream_t* stream, int opNum, sd::Pointer stateHost, void const* vx, sd::LongType const* xShapeBuffer, void const* vy, sd::LongType const* yShapeBuffer, void* vz, sd::LongType const* zShapeBuffer, void* vextraArguments);" dedupe_set content) else() add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void const* x, sd::LongType const* xShapeBuffer, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) add_unique_instantiation("template void functions::random::RandomFunction<${t1}>::execTransform(int opNum, sd::Pointer state, void const* x, sd::LongType const* xShapeBuffer, void const* y, sd::LongType const* yShapeBuffer, void* z, sd::LongType const* zShapeBuffer, void* extraArguments);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for reduce_same templates function(handle_reduce_same t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Generate ONLY for types.h typedefs (sd::LongType, sd::UnsignedLong) if(is_cuda) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execReduce(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, const sd::LongType* dXShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* vreductionBuffer, void* z, const sd::LongType* dZShapeInfo, const sd::LongType* hZShapeInfo, const sd::LongType* dims);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execReduceScalar(dim3 launchDims, cudaStream_t* stream, int opNum, const void* x, const sd::LongType* xShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* z, const sd::LongType* zShapeInfo, const sd::LongType* hZShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, void* reductionBuffer, const sd::LongType* tadOnlyShapeInfo);" dedupe_set content) add_unique_instantiation("template ${t1} functions::reduce::ReduceSameFunction<${t1}>::execScalarCudaLegacy(int opNum, void const* vx, sd::LongType const* xShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo);" dedupe_set content) else() add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::exec(int opNum, sd::memory::Workspace* workspace, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo, sd::LongType const* dimension);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceSameFunction<${t1}>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) add_unique_instantiation("template ${t1} functions::reduce::ReduceSameFunction<${t1}>::execScalar(int opNum, const void* x, const sd::LongType* xShapeInfo, void* extraParams);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for reduce3 templates function(handle_reduce3 t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") if(is_cuda) # Existing CUDA instantiations - these are correct add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, int postProcessOrNot, sd::LongType* allocationPointer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets, sd::LongType const* yTadOnlyShapeInfo, sd::LongType const* yTadOffsets);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execAll(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* dimension, long long int dimensionLength, int postProcessOrNot, sd::LongType* allocationPointer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets, sd::LongType const* yTadOnlyShapeInfo, sd::LongType const* yTadOffsets);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execScalar(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* allocationPointer, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo);" dedupe_set content) else() # CPU instantiations - all signatures from the implementation file # exec with 11 parameters (dimension, start, stop) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType start, sd::LongType stop);" dedupe_set content) # exec with 12 parameters (includes tadShapeInfo and tadOffsets) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType const* tadShapeInfo, sd::LongType const* tadOffsets, sd::LongType start, sd::LongType stop);" dedupe_set content) # execScalar add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) # execAll with full TAD information for both x and y add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execAll(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType const* xTadShapeInfo, sd::LongType const* xOffsets, sd::LongType const* yTadShapeInfo, sd::LongType const* yOffsets, sd::LongType start, sd::LongType stop);" dedupe_set content) # Note: The following are template methods with OpType and are typically instantiated elsewhere: # - execScalar # - exec (both overloads) # - execAll endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for indexreduce templates # ============================================================================ # TYPE NORMALIZATION FUNCTION # ============================================================================ # ============================================================================ # PAIRWISE HANDLER WITH NORMALIZATION # ============================================================================ function(handle_pairwise t1 t2 t3 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") message(STATUS "DEBUG: handle_pairwise called with t1='${t1}', t2='${t2}', t3='${t3}'") # Normalize types first normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) normalize_to_canonical_type("${t3}" norm_t3) message(STATUS "DEBUG: handle_pairwise normalized: '${norm_t1}', '${norm_t2}', '${norm_t3}'") # Skip any combination where ANY type isn't in its canonical form # This prevents duplicates like (unsigned short int, uint16_t, uint32_t) when (uint16_t, uint16_t, uint32_t) exists if(NOT (t1 STREQUAL norm_t1 AND t2 STREQUAL norm_t2 AND t3 STREQUAL norm_t3)) # At least one type is a non-canonical alias - skip to avoid duplicates set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map normalized types to enum names for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) map_cpp_to_enum("${norm_t3}" type3_enum) # Validation checks... _internal_srcore_is_valid_pair("${type1_enum}" "${type2_enum}" valid_12) _internal_srcore_is_valid_pair("${type2_enum}" "${type3_enum}" valid_23) _internal_srcore_is_valid_triple("${type1_enum}" "${type2_enum}" "${type3_enum}" valid_123) if(NOT valid_12 OR NOT valid_23 OR NOT valid_123) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations using NORMALIZED types (which are same as originals now due to check above) if(is_cuda) add_unique_instantiation("template void functions::pairwise_transforms::PairWiseTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::executeCudaShaped(dim3& launchDims, cudaStream_t* stream, int opNum, const void *vx, const sd::LongType *xShapeInfo, const void *vy, const sd::LongType *yShapeInfo, void *vz, const sd::LongType *zShapeInfo, void *vextraParams);" dedupe_set content) else() add_unique_instantiation("template void functions::pairwise_transforms::PairWiseTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo, void *extraParams, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::pairwise_transforms::PairWiseTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, sd::LongType xStride, const void *y, sd::LongType yStride, void *z, sd::LongType resultStride, void *extraParams, sd::LongType len, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # SCALAR HANDLER WITH NORMALIZATION # ============================================================================ function(handle_scalar t1 t2 t3 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) normalize_to_canonical_type("${t3}" norm_t3) # Skip any combination where ANY type isn't in its canonical form # This prevents duplicates like (unsigned short int, uint16_t, uint32_t) when (uint16_t, uint16_t, uint32_t) exists if(NOT (t1 STREQUAL norm_t1 AND t2 STREQUAL norm_t2 AND t3 STREQUAL norm_t3)) # At least one type is a non-canonical alias - skip to avoid duplicates set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) map_cpp_to_enum("${norm_t3}" type3_enum) # Validation _internal_srcore_is_valid_pair("${type1_enum}" "${type3_enum}" valid_13) _internal_srcore_is_valid_pair("${type2_enum}" "${type3_enum}" valid_23) _internal_srcore_is_valid_triple("${type1_enum}" "${type2_enum}" "${type3_enum}" valid_triple) if(NOT valid_13 OR NOT valid_23 OR NOT valid_triple) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations using NORMALIZED types (which are same as originals now due to check above) if(is_cuda) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::executeCudaShaped(dim3& launchDims, cudaStream_t* stream, int opNum, const void *vx, const sd::LongType *xShapeInfo, const sd::LongType *hxShapeInfo, void *vz, const sd::LongType *zShapeInfo, const sd::LongType *hzShapeInfo, const void *vscalar, void *vextraParams);" dedupe_set content) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::executeCudaAlongDimension(dim3& launchDims, cudaStream_t* stream, int opNum, const void *vx, const sd::LongType *xShapeInfo, void *vz, const sd::LongType *zShapeInfo, const void *vscalars, void *vextraParams, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) else() add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::transform(int opNum, const void *x, const sd::LongType *xShapeInfo, void *result, const sd::LongType *resultShapeInfo, const void *scalar, void *extraParams, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::scalar::ScalarTransform<${norm_t1}, ${norm_t2}, ${norm_t3}>::transform(int opNum, const void *x, const sd::LongType *xShapeInfo, void *extraParams, void *z, const sd::LongType *zShapeInfo, const void *scalars, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetsZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # BROADCAST HANDLER WITH NORMALIZATION # ============================================================================ function(handle_broadcast t1 t2 t3 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) normalize_to_canonical_type("${t3}" norm_t3) # Skip any combination where ANY type isn't in its canonical form # This prevents duplicates like (unsigned short int, uint16_t, uint32_t) when (uint16_t, uint16_t, uint32_t) exists if(NOT (t1 STREQUAL norm_t1 AND t2 STREQUAL norm_t2 AND t3 STREQUAL norm_t3)) # At least one type is a non-canonical alias - skip to avoid duplicates set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) map_cpp_to_enum("${norm_t3}" type3_enum) # Validation _internal_srcore_is_valid_triple("${type1_enum}" "${type2_enum}" "${type3_enum}" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations using NORMALIZED types (which are same as originals now due to check above) if(is_cuda) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadOnlyShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execInverseBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadOnlyShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LoopKind::Kind loopKind, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # REDUCE3 HANDLER WITH NORMALIZATION # ============================================================================ function(handle_reduce3 t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) # Skip redundant type aliases if(norm_t1 STREQUAL norm_t2) if(NOT t1 STREQUAL t2) set(${content_var} "${content}" PARENT_SCOPE) return() endif() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) _internal_srcore_is_valid_pair("${type1_enum}" "${type2_enum}" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, int postProcessOrNot, sd::LongType* allocationPointer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets, sd::LongType const* yTadOnlyShapeInfo, sd::LongType const* yTadOffsets);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execAll(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* dimension, long long int dimensionLength, int postProcessOrNot, sd::LongType* allocationPointer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets, sd::LongType const* yTadOnlyShapeInfo, sd::LongType const* yTadOffsets);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execScalar(dim3 launchDims, cudaStream_t* stream, int opNum, void const* vx, sd::LongType const* xShapeInfo, void const* vy, sd::LongType const* yShapeInfo, void* extraParams, void* vz, sd::LongType const* zShapeInfo, sd::LongType* allocationPointer, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo);" dedupe_set content) else() add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::exec(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType const* tadShapeInfo, sd::LongType const* tadOffsets, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::reduce3::Reduce3<${t1}, ${t2}>::execAll(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void const* y, sd::LongType const* yShapeInfo, void* z, sd::LongType const* zShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, sd::LongType const* xTadShapeInfo, sd::LongType const* xOffsets, sd::LongType const* yTadShapeInfo, sd::LongType const* yOffsets, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # INDEXREDUCE HANDLER WITH NORMALIZATION # ============================================================================ function(handle_indexreduce t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) # Skip redundant combinations for index type (t2) # Index type should always be int64_t after normalization if(norm_t2 STREQUAL "int64_t" AND NOT t2 STREQUAL "int64_t") # Skip LongType, sd::LongType, long, long long variants set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) _internal_srcore_is_valid_pair("${type1_enum}" "${type2_enum}" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::indexreduce::IndexReduce<${t1}, ${t2}>::executeIndexReduce(dim3 launchDims, cudaStream_t* stream, const int opNum, void const* dx, sd::LongType const* xShapeInfo, sd::LongType xRank, void* extraParams, void* result, sd::LongType const* zShapeInfo, sd::LongType zRank, sd::LongType* dimension, sd::LongType dimensionLength, int postProcessOrNot, sd::LongType* allocationBuffer, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets);" dedupe_set content) add_unique_instantiation("template void functions::indexreduce::IndexReduce<${t1}, ${t2}>::executeIndexReduceScalar(dim3 launchDims, cudaStream_t* stream, const int opNum, void const* dx, sd::LongType const* xShapeInfo, sd::LongType xRank, void* extraParams, void* result, sd::LongType const* zShapeInfo, sd::LongType zRank, sd::LongType* dimension, sd::LongType dimensionLength, int postProcessOrNot, sd::LongType* allocationBuffer, void* reductionBuffer, sd::LongType const* tadOnlyShapeInfo, sd::LongType const* tadOffsets);" dedupe_set content) else() add_unique_instantiation("template class sd::IndexReductionLoops<${t1}, ${t2}>;" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for reduce_float templates function(handle_reduce_float t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) # Only process if t2 is a float type if(NOT norm_t2 MATCHES "float|double|float16|bfloat16") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Skip redundant type alias combinations if(norm_t1 STREQUAL norm_t2 AND NOT t1 STREQUAL t2) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) _internal_srcore_is_valid_pair("${type1_enum}" "${type2_enum}" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::reduce::ReduceFloatFunction<${t1}, ${t2}>::execReduce(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, const sd::LongType* dXShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* vreductionBuffer, void* z, const sd::LongType* dZShapeInfo, const sd::LongType* hZShapeInfo, const sd::LongType* dims);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceFloatFunction<${t1}, ${t2}>::execReduceScalar(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, const sd::LongType* xShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* z, const sd::LongType* dZShapeInfo, const sd::LongType* hZShapeInfo, sd::LongType* dimension, long long int dimensionLength, void* reductionBuffer, const sd::LongType* tadOnlyShapeInfo);" dedupe_set content) else() add_unique_instantiation("template void functions::reduce::ReduceFloatFunction<${t1}, ${t2}>::exec(int opNum, sd::memory::Workspace* workspace, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo, sd::LongType const* dimension);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceFloatFunction<${t1}, ${t2}>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for reduce_bool templates function(handle_reduce_bool t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Only process if t2 is bool if(NOT t2 STREQUAL "bool") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Normalize t1 but don't skip aliases - NativeOpExecutioner uses 'long' explicitly normalize_to_canonical_type("${t1}" norm_t1) # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) _internal_srcore_is_valid_pair("${type1_enum}" "BOOL" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template void functions::reduce::ReduceBoolFunction<${t1}, bool>::execReduce(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, sd::LongType* dXShapeInfo, sd::LongType* hXShapeInfo, void* extraParams, void* vreductionBuffer, void* z, sd::LongType* dZShapeInfo, sd::LongType* hZShapeInfo, sd::LongType* dims);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceBoolFunction<${t1}, bool>::execReduceScalar(dim3 launchDims, cudaStream_t* stream, const int opNum, const void* x, const sd::LongType* xShapeInfo, const sd::LongType* hXShapeInfo, void* extraParams, void* z, const sd::LongType* zShapeInfo, const sd::LongType* hZShapeInfo, sd::LongType* dimension, sd::LongType dimensionLength, void* reductionBuffer, const sd::LongType* tadOnlyShapeInfo);" dedupe_set content) else() add_unique_instantiation("template void functions::reduce::ReduceBoolFunction<${t1}, bool>::exec(int opNum, sd::memory::Workspace* workspace, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo, sd::LongType const* dimension);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceBoolFunction<${t1}, bool>::execScalar(int opNum, void const* x, sd::LongType const* xShapeInfo, void* extraParams, void* z, sd::LongType const* zShapeInfo);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for reduce_long templates function(handle_reduce_long t1 t2 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types to canonical forms for validation normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) # Only process if t2 is a 64-bit integer type (maps to INT64 or UINT64) map_cpp_to_enum("${norm_t2}" type2_enum) if(NOT (type2_enum STREQUAL "INT64" OR type2_enum STREQUAL "UINT64")) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) message(STATUS "DEBUG handle_reduce_long: t1='${t1}' norm_t1='${norm_t1}' type1_enum='${type1_enum}'") message(STATUS "DEBUG handle_reduce_long: t2='${t2}' norm_t2='${norm_t2}' type2_enum='${type2_enum}'") _internal_srcore_is_valid_pair("${type1_enum}" "${type2_enum}" is_valid) message(STATUS "DEBUG handle_reduce_long: is_valid='${is_valid}' for pair (${type1_enum}, ${type2_enum})") if(NOT is_valid) message(STATUS "DEBUG handle_reduce_long: REJECTED pair (${type1_enum}, ${type2_enum})") set(${content_var} "${content}" PARENT_SCOPE) return() endif() message(STATUS "DEBUG handle_reduce_long: ACCEPTED pair (${type1_enum}, ${type2_enum}), generating instantiations") # Generate instantiations using types from types.h (sd::LongType = long long) # Generate ONLY for types.h typedefs (sd::LongType, sd::UnsignedLong) if(is_cuda) # CUDA reduce_long implementations (if needed) else() add_unique_instantiation("template void functions::reduce::ReduceLongFunction<${t1}, ${t2}>::exec(int, sd::memory::Workspace*, void const*, sd::LongType const*, void*, void*, sd::LongType const*, sd::LongType*);" dedupe_set content) add_unique_instantiation("template void functions::reduce::ReduceLongFunction<${t1}, ${t2}>::execScalar(int, void const*, sd::LongType const*, void*, void*, sd::LongType const*);" dedupe_set content) add_unique_instantiation("template class sd::ReductionLongLoops<${t1}, ${t2}>;" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # ============================================================================ # Legacy genCompilation function # ============================================================================ function(genCompilation FL_ITEM generated_sources_var) get_filename_component(FILE_ITEM_WE ${FL_ITEM} NAME_WE) set(EXTENSION "cpp") if(FL_ITEM MATCHES "cu.in$") set(EXTENSION "cu") endif() file(READ ${FL_ITEM} CONTENT_FL) set(SD_FLOAT_TYPES_GEN 0) set(SD_INTEGER_TYPES_GEN 0) set(SD_COMMON_TYPES_GEN 0) set(SD_PAIRWISE_TYPES_GEN 0) set(RANGE_STOP -1) string(REGEX MATCHALL "#cmakedefine[ \t]+SD_(INTEGER|COMMON|FLOAT|PAIRWISE)_TYPES_GEN" TYPE_MATCHES ${CONTENT_FL}) # Use selective rendering type count if available if(DEFINED UNIFIED_TYPE_COUNT AND UNIFIED_TYPE_COUNT GREATER 0) math(EXPR SD_INTEGER_TYPES_END "${UNIFIED_TYPE_COUNT} - 1") math(EXPR SD_COMMON_TYPES_END "${UNIFIED_TYPE_COUNT} - 1") math(EXPR SD_FLOAT_TYPES_END "${UNIFIED_TYPE_COUNT} - 1") math(EXPR SD_PAIRWISE_TYPES_END "${UNIFIED_TYPE_COUNT} - 1") else() set(SD_INTEGER_TYPES_END 7) set(SD_COMMON_TYPES_END 12) set(SD_FLOAT_TYPES_END 3) set(SD_PAIRWISE_TYPES_END 12) endif() foreach(TYPEX ${TYPE_MATCHES}) set(STOP -1) if(TYPEX MATCHES "SD_INTEGER_TYPES_GEN$") set(SD_INTEGER_TYPES_GEN 1) set(STOP ${SD_INTEGER_TYPES_END}) endif() if(TYPEX MATCHES "SD_COMMON_TYPES_GEN$") set(SD_COMMON_TYPES_GEN 1) set(STOP ${SD_COMMON_TYPES_END}) endif() if(TYPEX MATCHES "SD_FLOAT_TYPES_GEN$") set(SD_FLOAT_TYPES_GEN 1) set(STOP ${SD_FLOAT_TYPES_END}) endif() if(TYPEX MATCHES "SD_PAIRWISE_TYPES_GEN$") set(SD_PAIRWISE_TYPES_GEN 1) set(STOP ${SD_PAIRWISE_TYPES_END}) endif() if(STOP GREATER RANGE_STOP) set(RANGE_STOP ${STOP}) endif() endforeach() file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/compilation_units") set(local_generated_sources ${${generated_sources_var}}) if(RANGE_STOP GREATER -1) foreach(FL_TYPE_INDEX RANGE 0 ${RANGE_STOP}) set(CURRENT_SD_FLOAT_TYPES_GEN ${SD_FLOAT_TYPES_GEN}) set(CURRENT_SD_INTEGER_TYPES_GEN ${SD_INTEGER_TYPES_GEN}) set(CURRENT_SD_COMMON_TYPES_GEN ${SD_COMMON_TYPES_GEN}) set(CURRENT_SD_PAIRWISE_TYPES_GEN ${SD_PAIRWISE_TYPES_GEN}) if(FL_TYPE_INDEX GREATER ${SD_FLOAT_TYPES_END}) set(CURRENT_SD_FLOAT_TYPES_GEN 0) endif() if(FL_TYPE_INDEX GREATER ${SD_INTEGER_TYPES_END}) set(CURRENT_SD_INTEGER_TYPES_GEN 0) endif() if(FL_TYPE_INDEX GREATER ${SD_COMMON_TYPES_END}) set(CURRENT_SD_COMMON_TYPES_GEN 0) endif() set(SD_FLOAT_TYPES_GEN ${CURRENT_SD_FLOAT_TYPES_GEN}) set(SD_INTEGER_TYPES_GEN ${CURRENT_SD_INTEGER_TYPES_GEN}) set(SD_COMMON_TYPES_GEN ${CURRENT_SD_COMMON_TYPES_GEN}) set(SD_PAIRWISE_TYPES_GEN ${CURRENT_SD_PAIRWISE_TYPES_GEN}) set(GENERATED_SOURCE "${CMAKE_BINARY_DIR}/compilation_units/${FILE_ITEM_WE}_${FL_TYPE_INDEX}.${EXTENSION}") configure_file("${FL_ITEM}" "${GENERATED_SOURCE}" @ONLY) list(APPEND local_generated_sources ${GENERATED_SOURCE}) endforeach() endif() set(${generated_sources_var} ${local_generated_sources} PARENT_SCOPE) endfunction() # ============================================================================ # Direct instantiation file creation # ============================================================================ function(setup_type_mapping) # Setup type index to C++ type mapping if(DEFINED UNIFIED_ACTIVE_TYPES) set(idx 0) foreach(type ${UNIFIED_ACTIVE_TYPES}) if(type STREQUAL "BOOL") set(SRCORE_TYPE_CPP_${idx} "bool" PARENT_SCOPE) elseif(type STREQUAL "INT8") set(SRCORE_TYPE_CPP_${idx} "int8_t;SignedChar;signed char;char" PARENT_SCOPE) elseif(type STREQUAL "UINT8") set(SRCORE_TYPE_CPP_${idx} "uint8_t;UnsignedChar;unsigned char" PARENT_SCOPE) elseif(type STREQUAL "INT16") set(SRCORE_TYPE_CPP_${idx} "int16_t;short;short int;signed short" PARENT_SCOPE) elseif(type STREQUAL "UINT16") set(SRCORE_TYPE_CPP_${idx} "uint16_t;unsigned short;unsigned short int" PARENT_SCOPE) elseif(type STREQUAL "INT32") set(SRCORE_TYPE_CPP_${idx} "int32_t;int;signed int;Int32Type" PARENT_SCOPE) elseif(type STREQUAL "UINT32") set(SRCORE_TYPE_CPP_${idx} "uint32_t;unsigned int;unsigned" PARENT_SCOPE) elseif(type STREQUAL "INT64") # Use sd::LongType (long long) from types.h set(SRCORE_TYPE_CPP_${idx} "sd::LongType" PARENT_SCOPE) elseif(type STREQUAL "UINT64") # Use sd::UnsignedLong (uint64_t) from types.h set(SRCORE_TYPE_CPP_${idx} "sd::UnsignedLong" PARENT_SCOPE) elseif(type STREQUAL "FLOAT32") set(SRCORE_TYPE_CPP_${idx} "float" PARENT_SCOPE) elseif(type STREQUAL "DOUBLE") set(SRCORE_TYPE_CPP_${idx} "double" PARENT_SCOPE) elseif(type STREQUAL "HALF") set(SRCORE_TYPE_CPP_${idx} "float16" PARENT_SCOPE) elseif(type STREQUAL "BFLOAT16") set(SRCORE_TYPE_CPP_${idx} "bfloat16" PARENT_SCOPE) elseif(type STREQUAL "UTF8") set(SRCORE_TYPE_CPP_${idx} "std::string" PARENT_SCOPE) set(SRCORE_IS_STRING_${idx} TRUE PARENT_SCOPE) elseif(type STREQUAL "UTF16") set(SRCORE_TYPE_CPP_${idx} "std::u16string" PARENT_SCOPE) set(SRCORE_IS_STRING_${idx} TRUE PARENT_SCOPE) elseif(type STREQUAL "UTF32") set(SRCORE_TYPE_CPP_${idx} "std::u32string" PARENT_SCOPE) set(SRCORE_IS_STRING_${idx} TRUE PARENT_SCOPE) else() message(WARNING "Unknown type in UNIFIED_ACTIVE_TYPES: ${type}") set(SRCORE_TYPE_CPP_${idx} "" PARENT_SCOPE) endif() math(EXPR idx "${idx} + 1") endforeach() endif() endfunction() function(create_direct_instantiation_file template_file combinations output_dir generated_sources_var) get_filename_component(template_name ${template_file} NAME_WE) file(READ "${template_file}" template_content) string(REGEX MATCHALL "#include[^\n]*" includes "${template_content}") # Determine file extension set(IS_CUDA_FILE FALSE) if(template_file MATCHES "\\.cu\\.in$") set(IS_CUDA_FILE TRUE) set(file_extension "cu") else() set(file_extension "cpp") endif() # Build file header set(file_header "") foreach(inc ${includes}) string(APPEND file_header "${inc}\n") endforeach() # Add extra includes for type conversions if(template_name MATCHES ".*specials_double.*") string(APPEND file_header "#include \n") endif() string(APPEND file_header "\n// Direct instantiations - comprehensive type variant generation\n\n") # Setup type mapping setup_type_mapping() # Process combinations set(chunk_content "${file_header}") set(instantiation_count 0) set(chunk_index 0) set(local_generated_sources ${${generated_sources_var}}) set(total_instantiations 0) set(processed_normalized_combinations "") # Track normalized combinations to avoid duplicates foreach(combination ${combinations}) string(REPLACE "," ";" parts "${combination}") list(LENGTH parts parts_count) # Extract types from indices set(t1 "") set(t2 "") set(t3 "") set(has_string FALSE) if(parts_count GREATER_EQUAL 1) list(GET parts 0 idx1) if(DEFINED SRCORE_TYPE_CPP_${idx1}) set(raw_t1 "${SRCORE_TYPE_CPP_${idx1}}") # Normalize all types to use typedefs from types.h (sd::LongType, sd::UnsignedLong) normalize_to_canonical_type("${raw_t1}" t1) if(DEFINED SRCORE_IS_STRING_${idx1} AND SRCORE_IS_STRING_${idx1}) set(has_string TRUE) endif() endif() endif() if(parts_count GREATER_EQUAL 2) list(GET parts 1 idx2) if(DEFINED SRCORE_TYPE_CPP_${idx2}) set(raw_t2 "${SRCORE_TYPE_CPP_${idx2}}") # Normalize all types to use typedefs from types.h (sd::LongType, sd::UnsignedLong) normalize_to_canonical_type("${raw_t2}" t2) if(DEFINED SRCORE_IS_STRING_${idx2} AND SRCORE_IS_STRING_${idx2}) set(has_string TRUE) endif() endif() endif() if(parts_count GREATER_EQUAL 3) list(GET parts 2 idx3) if(DEFINED SRCORE_TYPE_CPP_${idx3}) set(raw_t3 "${SRCORE_TYPE_CPP_${idx3}}") # Normalize all types to use typedefs from types.h (sd::LongType, sd::UnsignedLong) normalize_to_canonical_type("${raw_t3}" t3) if(DEFINED SRCORE_IS_STRING_${idx3} AND SRCORE_IS_STRING_${idx3}) set(has_string TRUE) endif() endif() endif() # Skip string combinations if(has_string) continue() endif() # Skip if we couldn't map the types if(NOT t1) continue() endif() # Skip if we've already processed this normalized combination # This prevents duplicates when 'long' and 'int64_t' both normalize to 'int64_t' set(normalized_combo "${t1}") if(t2) set(normalized_combo "${normalized_combo},${t2}") endif() if(t3) set(normalized_combo "${normalized_combo},${t3}") endif() list(FIND processed_normalized_combinations "${normalized_combo}" already_processed) if(already_processed GREATER -1) continue() endif() list(APPEND processed_normalized_combinations "${normalized_combo}") # Dispatch to appropriate handler dispatch_to_handler("${template_name}" "${t1}" "${t2}" "${t3}" "${parts_count}" chunk_content ${IS_CUDA_FILE}) math(EXPR instantiation_count "${instantiation_count} + 1") math(EXPR total_instantiations "${total_instantiations} + 1") # Write chunk if limit reached if(instantiation_count GREATER_EQUAL ${MULTI_PASS_CHUNK_SIZE}) set(chunk_file "${output_dir}/${template_name}_direct_${chunk_index}.${file_extension}") file(WRITE "${chunk_file}" "${chunk_content}") list(APPEND local_generated_sources "${chunk_file}") set(chunk_content "${file_header}") set(instantiation_count 0) math(EXPR chunk_index "${chunk_index} + 1") endif() endforeach() # Write remaining content if(instantiation_count GREATER 0) # If call tracing is enabled, don't create huge "_final" files # Instead, write them as regular chunks to prevent relocation overflow if(SD_GCC_FUNCTRACE) set(chunk_file "${output_dir}/${template_name}_direct_${chunk_index}.${file_extension}") else() set(chunk_file "${output_dir}/${template_name}_direct_final.${file_extension}") endif() file(WRITE "${chunk_file}" "${chunk_content}") list(APPEND local_generated_sources "${chunk_file}") endif() set(${generated_sources_var} ${local_generated_sources} PARENT_SCOPE) if(IS_CUDA_FILE) message(STATUS "✅ Generated CUDA instantiation files for ${template_name}: ${total_instantiations} instantiations") else() message(STATUS "✅ Generated CPU instantiation files for ${template_name}: ${total_instantiations} instantiations") endif() endfunction() # Handler for broadcast_int templates - excludes bool and string types function(handle_broadcast_int t1 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize type normalize_to_canonical_type("${t1}" norm_t1) # Skip bool and string types if(norm_t1 STREQUAL "bool" OR norm_t1 MATCHES "std::string|std::u16string|std::u32string") set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Skip redundant type aliases (only keep canonical form) if(NOT t1 STREQUAL norm_t1) # This is an alias, skip if we're not the canonical form set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations if(is_cuda) add_unique_instantiation("template class BroadcastInt<${t1}>;" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::BroadcastInt<${t1}>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Handler for broadcast (3-type) templates function(handle_broadcast t1 t2 t3 content_var is_cuda) set(content "${${content_var}}") set(dedupe_set "") # Normalize types normalize_to_canonical_type("${t1}" norm_t1) normalize_to_canonical_type("${t2}" norm_t2) normalize_to_canonical_type("${t3}" norm_t3) # Skip any combination where ANY type isn't in its canonical form # This prevents duplicates like (unsigned short int, uint16_t, uint32_t) when (uint16_t, uint16_t, uint32_t) exists if(NOT (t1 STREQUAL norm_t1 AND t2 STREQUAL norm_t2 AND t3 STREQUAL norm_t3)) # At least one type is a non-canonical alias - skip to avoid duplicates set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Map to enum for validation map_cpp_to_enum("${norm_t1}" type1_enum) map_cpp_to_enum("${norm_t2}" type2_enum) map_cpp_to_enum("${norm_t3}" type3_enum) # Validation _internal_srcore_is_valid_triple("${type1_enum}" "${type2_enum}" "${type3_enum}" is_valid) if(NOT is_valid) set(${content_var} "${content}" PARENT_SCOPE) return() endif() # Generate instantiations using NORMALIZED types (which are same as originals now due to check above) if(is_cuda) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadOnlyShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execInverseBroadcast(dim3 launchDims, cudaStream_t *stream, int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadOnlyShapeInfo, const sd::LongType *tadOffsets, const sd::LongType *tadOnlyShapeInfoZ, const sd::LongType *tadOffsetsZ);" dedupe_set content) else() add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LoopKind::Kind loopKind, sd::LongType start, sd::LongType stop);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::exec(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *z, const sd::LongType *zShapeInfo);" dedupe_set content) add_unique_instantiation("template void functions::broadcast::Broadcast<${norm_t1}, ${norm_t2}, ${norm_t3}>::execInverse(int opNum, const void *x, const sd::LongType *xShapeInfo, const void *y, const sd::LongType *yShapeInfo, void *result, const sd::LongType *resultShapeInfo, sd::LongType *dimension, sd::LongType dimensionLength, const sd::LongType *tadShapeInfo, const sd::LongType *tadOffset, const sd::LongType *tadShapeInfoZ, const sd::LongType *tadOffsetZ, sd::LongType start, sd::LongType stop);" dedupe_set content) endif() set(${content_var} "${content}" PARENT_SCOPE) endfunction() # Helper function to map C++ type to enum (local to this function) function(map_cpp_to_enum cpp_type enum_var) set(enum_name "") if(cpp_type MATCHES "int8_t|SignedChar|signed char|char") set(enum_name "INT8") elseif(cpp_type MATCHES "uint8_t|UnsignedChar|unsigned char") set(enum_name "UINT8") elseif(cpp_type MATCHES "int16_t|short") set(enum_name "INT16") elseif(cpp_type MATCHES "uint16_t|unsigned short") set(enum_name "UINT16") elseif(cpp_type MATCHES "int32_t|int|Int32Type|signed int") set(enum_name "INT32") elseif(cpp_type MATCHES "uint32_t|unsigned int|unsigned") set(enum_name "UINT32") elseif(cpp_type MATCHES "int64_t|long long|long|LongType|sd::LongType") set(enum_name "INT64") elseif(cpp_type MATCHES "uint64_t|unsigned long long|unsigned long|UnsignedLong|sd::UnsignedLong") set(enum_name "UINT64") elseif(cpp_type STREQUAL "float") set(enum_name "FLOAT32") elseif(cpp_type STREQUAL "double") set(enum_name "DOUBLE") elseif(cpp_type STREQUAL "float16") set(enum_name "HALF") elseif(cpp_type STREQUAL "bfloat16") set(enum_name "BFLOAT16") elseif(cpp_type STREQUAL "bool") set(enum_name "BOOL") endif() set(${enum_var} "${enum_name}" PARENT_SCOPE) endfunction() function(dispatch_to_handler template_name t1 t2 t3 parts_count content_var is_cuda) set(content "${${content_var}}") # DEBUG: Log every call message(STATUS "DEBUG dispatch_to_handler:") message(STATUS " template_name: '${template_name}'") message(STATUS " parts_count: ${parts_count}") message(STATUS " t1: '${t1}'") message(STATUS " t2: '${t2}'") message(STATUS " t3: '${t3}'") # Create a template-specific deduplication set name if(is_cuda) set(dedupe_property_name "DEDUPE_SET_${template_name}_CUDA") else() set(dedupe_property_name "DEDUPE_SET_${template_name}_CPU") endif() get_property(template_dedupe_set GLOBAL PROPERTY ${dedupe_property_name}) if(NOT template_dedupe_set) set(template_dedupe_set "") endif() # Parse type lists if(t1) string(REPLACE ";" "\\;" t1_escaped "${t1}") string(REPLACE "\\;" ";" t1_list "${t1_escaped}") else() set(t1_list "") endif() if(NOT t2 STREQUAL "") string(REPLACE ";" "\\;" t2_escaped "${t2}") string(REPLACE "\\;" ";" t2_list "${t2_escaped}") else() set(t2_list "") endif() if(NOT t3 STREQUAL "") string(REPLACE ";" "\\;" t3_escaped "${t3}") string(REPLACE "\\;" ";" t3_list "${t3_escaped}") else() set(t3_list "") endif() # Single type templates if(template_name MATCHES ".*specials_single.*") foreach(type IN LISTS t1_list) normalize_to_canonical_type("${type}" norm_type) set(combo_key "${norm_type}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) handle_specials_single("${norm_type}" content) list(APPEND template_dedupe_set "${combo_key}") endif() endforeach() # Two type templates elseif(template_name MATCHES ".*specials_double.*" AND parts_count GREATER_EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) set(combo_key "${norm_v1},${norm_v2}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "${v2_enum}" is_valid) if(is_valid) handle_specials_double("${norm_v1}" "${norm_v2}" content) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() elseif(template_name MATCHES ".*broadcast_int.*") foreach(type IN LISTS t1_list) normalize_to_canonical_type("${type}" norm_type) set(combo_key "${norm_type}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) if(NOT norm_type STREQUAL "bool" AND NOT norm_type MATCHES "std::string|std::u16string|std::u32string") handle_broadcast_int("${norm_type}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endforeach() elseif(template_name MATCHES ".*broadcast_bool.*" AND parts_count EQUAL 2) foreach(v2 IN LISTS t2_list) if(v2 STREQUAL "bool") foreach(v1 IN LISTS t1_list) normalize_to_canonical_type("${v1}" norm_v1) set(combo_key "${norm_v1},bool") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) if(NOT norm_v1 MATCHES "std::string|std::u16string|std::u32string") handle_broadcast_bool("${norm_v1}" "bool" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endforeach() endif() endforeach() # Three type templates elseif(template_name MATCHES ".*broadcast.*" AND NOT template_name MATCHES ".*broadcast_bool.*|.*broadcast_int.*" AND parts_count EQUAL 3) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) foreach(v3 IN LISTS t3_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) normalize_to_canonical_type("${v3}" norm_v3) set(combo_key "${norm_v1},${norm_v2},${norm_v3}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) map_cpp_to_enum("${norm_v3}" v3_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "" AND NOT v3_enum STREQUAL "") _internal_srcore_is_valid_triple("${v1_enum}" "${v2_enum}" "${v3_enum}" is_valid) if(is_valid) handle_broadcast("${norm_v1}" "${norm_v2}" "${norm_v3}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() endforeach() elseif(template_name MATCHES "pairwise" AND parts_count EQUAL 3) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) foreach(v3 IN LISTS t3_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) normalize_to_canonical_type("${v3}" norm_v3) set(combo_key "${norm_v1},${norm_v2},${norm_v3}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) map_cpp_to_enum("${norm_v3}" v3_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "" AND NOT v3_enum STREQUAL "") _internal_srcore_is_valid_triple("${v1_enum}" "${v2_enum}" "${v3_enum}" is_valid) if(is_valid) handle_pairwise("${norm_v1}" "${norm_v2}" "${norm_v3}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() endforeach() elseif(template_name MATCHES ".*scalar.*" AND parts_count EQUAL 3) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) foreach(v3 IN LISTS t3_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) normalize_to_canonical_type("${v3}" norm_v3) set(combo_key "${norm_v1},${norm_v2},${norm_v3}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) map_cpp_to_enum("${norm_v3}" v3_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "" AND NOT v3_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "${v3_enum}" valid_13) _internal_srcore_is_valid_pair("${v2_enum}" "${v3_enum}" valid_23) _internal_srcore_is_valid_triple("${v1_enum}" "${v2_enum}" "${v3_enum}" valid_triple) if(valid_13 AND valid_23 AND valid_triple) handle_scalar("${norm_v1}" "${norm_v2}" "${norm_v3}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() endforeach() # Single type handlers elseif(template_name MATCHES ".*random.*") foreach(type IN LISTS t1_list) normalize_to_canonical_type("${type}" norm_type) set(combo_key "${norm_type}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) handle_random("${norm_type}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endforeach() elseif(template_name MATCHES ".*reduce_same.*") foreach(type IN LISTS t1_list) normalize_to_canonical_type("${type}" norm_type) set(combo_key "${norm_type}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) handle_reduce_same("${type}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endforeach() # Two type handlers with validation elseif(template_name MATCHES ".*reduce3.*" AND parts_count EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) set(combo_key "${norm_v1},${norm_v2}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "${v2_enum}" is_valid) if(is_valid) handle_reduce3("${norm_v1}" "${norm_v2}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() elseif(template_name MATCHES ".*indexreduc.*" AND parts_count EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) set(combo_key "${norm_v1},${norm_v2}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "${v2_enum}" is_valid) if(is_valid) handle_indexreduce("${norm_v1}" "${norm_v2}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endforeach() endforeach() elseif(template_name MATCHES ".*reduce_float.*" AND parts_count EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) normalize_to_canonical_type("${v1}" norm_v1) normalize_to_canonical_type("${v2}" norm_v2) if(norm_v2 MATCHES "float|double|float16|bfloat16") set(combo_key "${norm_v1},${norm_v2}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) map_cpp_to_enum("${norm_v2}" v2_enum) if(NOT v1_enum STREQUAL "" AND NOT v2_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "${v2_enum}" is_valid) if(is_valid) handle_reduce_float("${norm_v1}" "${norm_v2}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endif() endforeach() endforeach() elseif(template_name MATCHES ".*reduce_bool.*" AND parts_count EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) if(v2 STREQUAL "bool") normalize_to_canonical_type("${v1}" norm_v1) set(combo_key "${norm_v1},bool") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) if(NOT v1_enum STREQUAL "") _internal_srcore_is_valid_pair("${v1_enum}" "BOOL" is_valid) if(is_valid) handle_reduce_bool("${v1}" "bool" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endif() endforeach() endforeach() elseif(template_name MATCHES ".*reduce_long.*" AND parts_count EQUAL 2) foreach(v1 IN LISTS t1_list) foreach(v2 IN LISTS t2_list) normalize_to_canonical_type("${v2}" norm_v2) if(norm_v2 STREQUAL "sd::LongType" OR norm_v2 STREQUAL "sd::UnsignedLong") normalize_to_canonical_type("${v1}" norm_v1) set(combo_key "${norm_v1},${norm_v2}") list(FIND template_dedupe_set "${combo_key}" found_idx) if(found_idx EQUAL -1) map_cpp_to_enum("${norm_v1}" v1_enum) if(NOT v1_enum STREQUAL "") # Check if valid pair for either INT64 or UINT64 _internal_srcore_is_valid_pair("${v1_enum}" "INT64" is_valid_int64) _internal_srcore_is_valid_pair("${v1_enum}" "UINT64" is_valid_uint64) if(is_valid_int64 OR is_valid_uint64) handle_reduce_long("${v1}" "${v2}" content ${is_cuda}) list(APPEND template_dedupe_set "${combo_key}") endif() endif() endif() endif() endforeach() endforeach() else() if(CMAKE_BUILD_TYPE STREQUAL "Debug") message(STATUS "No handler for template: ${template_name} with ${parts_count} types") endif() endif() # Save the updated template-specific dedupe set set_property(GLOBAL PROPERTY ${dedupe_property_name} "${template_dedupe_set}") set(${content_var} "${content}" PARENT_SCOPE) endfunction() function(setup_fallback_combinations) # Fallback when selective rendering is not available set(UNIFIED_COMBINATIONS_1 "0;1;2;3;4;5;6;7;8;9;10;11;12" PARENT_SCOPE) set(combos_2 "") foreach(i RANGE 0 12) foreach(j RANGE 0 12) list(APPEND combos_2 "${i},${j}") endforeach() endforeach() set(UNIFIED_COMBINATIONS_2 ${combos_2} PARENT_SCOPE) set(combos_3 "") foreach(i RANGE 0 12) foreach(j RANGE 0 12) foreach(k RANGE 0 12) list(APPEND combos_3 "${i},${j},${k}") endforeach() endforeach() endforeach() set(UNIFIED_COMBINATIONS_3 ${combos_3} PARENT_SCOPE) # Setup default active types set(UNIFIED_ACTIVE_TYPES "BOOL;INT8;UINT8;INT16;UINT16;INT32;UINT32;INT64;UINT64;FLOAT32;DOUBLE;HALF;BFLOAT16" PARENT_SCOPE) set(UNIFIED_TYPE_COUNT 13 PARENT_SCOPE) endfunction() # ============================================================================ # CUDA template processing # ============================================================================ function(process_cuda_comb_templates output_dir generated_sources_var) message(STATUS "🔧 Processing CUDA combination templates...") # Find all CUDA combination templates file(GLOB_RECURSE CUDA_COMB_TEMPLATES "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cuda/comb_compilation_units/*.cu.in") if(NOT CUDA_COMB_TEMPLATES) message(STATUS "⚠️ No CUDA combination templates found") return() endif() if(NOT DEFINED UNIFIED_COMBINATIONS_2 OR NOT DEFINED UNIFIED_COMBINATIONS_3) message(FATAL_ERROR "❌ CUDA processing requires selective rendering combinations!") endif() set(local_generated_sources ${${generated_sources_var}}) list(LENGTH CUDA_COMB_TEMPLATES template_count) message(STATUS "📋 Found ${template_count} CUDA combination templates") # Use selective rendering combinations if(DEFINED UNIFIED_COMBINATIONS_2 AND DEFINED UNIFIED_COMBINATIONS_3) set(COMBINATIONS_1 "") set(COMBINATIONS_2 ${UNIFIED_COMBINATIONS_2}) set(COMBINATIONS_3 ${UNIFIED_COMBINATIONS_3}) # Generate 1-type combinations from 2-type combinations foreach(combination ${COMBINATIONS_2}) string(REPLACE "," ";" combo_parts "${combination}") list(GET combo_parts 0 comb1) list(FIND COMBINATIONS_1 "${comb1}" found_idx) if(found_idx EQUAL -1) list(APPEND COMBINATIONS_1 "${comb1}") endif() list(GET combo_parts 1 comb2) list(FIND COMBINATIONS_1 "${comb2}" found_idx) if(found_idx EQUAL -1) list(APPEND COMBINATIONS_1 "${comb2}") endif() endforeach() foreach(combination ${COMBINATIONS_3}) string(REPLACE "," ";" combo_parts "${combination}") foreach(part ${combo_parts}) list(FIND COMBINATIONS_1 "${part}" found_idx) if(found_idx EQUAL -1) list(APPEND COMBINATIONS_1 "${part}") endif() endforeach() endforeach() list(LENGTH COMBINATIONS_1 combo_1_count) list(LENGTH COMBINATIONS_2 combo_2_count) list(LENGTH COMBINATIONS_3 combo_3_count) message(STATUS "🎯 Using selective CUDA combinations: 1-type=${combo_1_count}, 2-type=${combo_2_count}, 3-type=${combo_3_count}") endif() foreach(TEMPLATE_FILE ${CUDA_COMB_TEMPLATES}) get_filename_component(template_name ${TEMPLATE_FILE} NAME_WE) # Determine which combination set to use based on template name if(template_name MATCHES ".*_template_1\$") set(combinations_to_use ${COMBINATIONS_1}) message(STATUS "🔄 Processing 1-type CUDA template: ${template_name}") elseif(template_name MATCHES ".*_template_2\$") set(combinations_to_use ${COMBINATIONS_2}) message(STATUS "🔄 Processing 2-type CUDA template: ${template_name}") elseif(template_name MATCHES ".*_template_3\$") set(combinations_to_use ${COMBINATIONS_3}) message(STATUS "🔄 Processing 3-type CUDA template: ${template_name}") else() # Default to 3-type for templates without explicit numbering set(combinations_to_use ${COMBINATIONS_3}) message(STATUS "🔄 Processing CUDA template (defaulting to 3-type): ${template_name}") endif() create_direct_instantiation_file("${TEMPLATE_FILE}" "${combinations_to_use}" "${output_dir}" local_generated_sources) endforeach() set(${generated_sources_var} ${local_generated_sources} PARENT_SCOPE) list(LENGTH local_generated_sources cuda_generated_count) message(STATUS "✅ CUDA combination template processing complete: ${cuda_generated_count} files generated") endfunction() # ============================================================================ # Main template generation execution # ============================================================================ function(execute_template_generation) message(STATUS "🔧 STARTING UNIFIED CPU/CUDA TEMPLATE GENERATION") set(ALL_GENERATED_SOURCES "") # Use the combinations provided by selective rendering if(DEFINED UNIFIED_COMBINATIONS_3) set(COMBINATIONS_3 ${UNIFIED_COMBINATIONS_3}) else() message(WARNING "⚠️ No 3-type combinations provided by selective rendering") set(COMBINATIONS_3 "") endif() if(DEFINED UNIFIED_COMBINATIONS_2) set(COMBINATIONS_2 ${UNIFIED_COMBINATIONS_2}) else() message(WARNING "⚠️ No 2-type combinations provided by selective rendering") set(COMBINATIONS_2 "") endif() # Derive 1-type combinations from 2-type combinations set(COMBINATIONS_1 "") foreach(combination ${COMBINATIONS_2}) string(REPLACE "," ";" combo_parts "${combination}") list(GET combo_parts 0 comb1) list(FIND COMBINATIONS_1 "${comb1}" found_idx) if(found_idx EQUAL -1) list(APPEND COMBINATIONS_1 "${comb1}") endif() list(GET combo_parts 1 comb2) list(FIND COMBINATIONS_1 "${comb2}" found_idx) if(found_idx EQUAL -1) list(APPEND COMBINATIONS_1 "${comb2}") endif() endforeach() # Report what we're working with list(LENGTH COMBINATIONS_1 combo_1_count) list(LENGTH COMBINATIONS_2 combo_2_count) list(LENGTH COMBINATIONS_3 combo_3_count) message(STATUS "📊 Template combinations from selective rendering:") message(STATUS " - 1-type: ${combo_1_count} combinations") message(STATUS " - 2-type: ${combo_2_count} combinations") message(STATUS " - 3-type: ${combo_3_count} combinations") if(SD_CUDA) message(STATUS "🚀 Processing CUDA templates...") # Process legacy CUDA compilation units file(GLOB_RECURSE CUDA_COMPILATION_UNITS ./include/loops/cuda/compilation_units/*.cu.in ./include/ops/impl/compilation_units/*.cpp.in) foreach(FL_ITEM ${CUDA_COMPILATION_UNITS}) genCompilation(${FL_ITEM} ALL_GENERATED_SOURCES) endforeach() set(CUDA_INST_DIR "${CMAKE_BINARY_DIR}/cuda_instantiations") file(MAKE_DIRECTORY "${CUDA_INST_DIR}") # Process CUDA combination templates process_cuda_comb_templates("${CUDA_INST_DIR}" ALL_GENERATED_SOURCES) else() message(STATUS "🖥️ Processing CPU templates...") # Process legacy CPU compilation units file(GLOB_RECURSE REGULAR_COMPILATION_UNITS ./include/ops/declarable/helpers/cpu/compilation_units/*.cpp.in ./include/helpers/cpu/loops/*.cpp.in) file(GLOB_RECURSE EXCLUDED_OLD_TEMPLATES ./include/loops/cpu/compilation_units/scalar_p.cpp.in ./include/loops/cpu/compilation_units/broadcast_p.cpp.in) file(GLOB_RECURSE REMAINING_COMPILATION_UNITS ./include/loops/cpu/compilation_units/*.cpp.in) if(EXCLUDED_OLD_TEMPLATES) list(REMOVE_ITEM REMAINING_COMPILATION_UNITS ${EXCLUDED_OLD_TEMPLATES}) endif() list(APPEND REGULAR_COMPILATION_UNITS ${REMAINING_COMPILATION_UNITS}) foreach(FL_ITEM ${REGULAR_COMPILATION_UNITS}) genCompilation(${FL_ITEM} ALL_GENERATED_SOURCES) endforeach() set(CPU_INST_DIR "${CMAKE_BINARY_DIR}/cpu_instantiations") file(MAKE_DIRECTORY "${CPU_INST_DIR}") # Process all CPU template files # 3-type templates (use COMBINATIONS_3) set(TEMPLATES_3 "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/pairwise_instantiation_template_3.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/scalar_instantiation_template_3.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/broadcast_instantiation_template_3.cpp.in") foreach(TEMPLATE_FILE ${TEMPLATES_3}) if(EXISTS "${TEMPLATE_FILE}") message(STATUS " Processing 3-type template: ${TEMPLATE_FILE}") create_direct_instantiation_file("${TEMPLATE_FILE}" "${COMBINATIONS_3}" "${CPU_INST_DIR}" ALL_GENERATED_SOURCES) else() message(WARNING " Template file not found: ${TEMPLATE_FILE}") endif() endforeach() # 2-type templates (use COMBINATIONS_2) set(TEMPLATES_2 "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/broadcast_bool_instantiation_template_2.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/indexreduction_instantiation_template_2.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/reduce3_instantiation_template_2.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/reduce_float_instantiation_template_2.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/reduce_bool_instantiation_template_2.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/reduce_long_instantiation_template_2.cpp.in") foreach(TEMPLATE_FILE ${TEMPLATES_2}) if(EXISTS "${TEMPLATE_FILE}") message(STATUS " Processing 2-type template: ${TEMPLATE_FILE}") create_direct_instantiation_file("${TEMPLATE_FILE}" "${COMBINATIONS_2}" "${CPU_INST_DIR}" ALL_GENERATED_SOURCES) else() message(WARNING " Template file not found: ${TEMPLATE_FILE}") endif() endforeach() # 1-type templates (use COMBINATIONS_1) set(TEMPLATES_1 "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/broadcast_int_instantiation_template_1.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/random_instantiation_template_1.cpp.in" "${CMAKE_CURRENT_SOURCE_DIR}/include/loops/cpu/comb_compilation_units/reduce_same_instantiation_template_1.cpp.in") foreach(TEMPLATE_FILE ${TEMPLATES_1}) if(EXISTS "${TEMPLATE_FILE}") message(STATUS " Processing 1-type template: ${TEMPLATE_FILE}") create_direct_instantiation_file("${TEMPLATE_FILE}" "${COMBINATIONS_1}" "${CPU_INST_DIR}" ALL_GENERATED_SOURCES) else() message(WARNING " Template file not found: ${TEMPLATE_FILE}") endif() endforeach() # Process specials templates set(SPECIALS_DOUBLE_TEMPLATE "${CMAKE_CURRENT_SOURCE_DIR}/include/ops/impl/compilation_units/specials_double.cpp.in") if(EXISTS "${SPECIALS_DOUBLE_TEMPLATE}") message(STATUS " Processing specials_double template") create_direct_instantiation_file("${SPECIALS_DOUBLE_TEMPLATE}" "${COMBINATIONS_2}" "${CPU_INST_DIR}" ALL_GENERATED_SOURCES) endif() set(SPECIALS_SINGLE_TEMPLATE "${CMAKE_CURRENT_SOURCE_DIR}/include/ops/impl/compilation_units/specials_single.cpp.in") if(EXISTS "${SPECIALS_SINGLE_TEMPLATE}") message(STATUS " Processing specials_single template") create_direct_instantiation_file("${SPECIALS_SINGLE_TEMPLATE}" "${COMBINATIONS_1}" "${CPU_INST_DIR}" ALL_GENERATED_SOURCES) endif() endif() # Final reporting set(CUSTOMOPS_GENERIC_SOURCES ${ALL_GENERATED_SOURCES} CACHE INTERNAL "Template-generated source files" FORCE) list(LENGTH ALL_GENERATED_SOURCES total_count) if(SD_CUDA) message(STATUS "✅ CUDA template processing complete: ${total_count} files generated") else() message(STATUS "✅ CPU template processing complete: ${total_count} files generated") endif() # Export to parent scope set(CUSTOMOPS_GENERIC_SOURCES ${ALL_GENERATED_SOURCES} PARENT_SCOPE) endfunction() # ============================================================================ # Main entry point for external calls # ============================================================================ function(setup_template_processing) # Check if selective rendering has already been executed get_property(cached_sources CACHE CUSTOMOPS_GENERIC_SOURCES PROPERTY VALUE) if(cached_sources) # Use cached results set(CUSTOMOPS_GENERIC_SOURCES ${cached_sources} PARENT_SCOPE) list(LENGTH cached_sources cached_count) if(SD_CUDA) message(STATUS "🔄 Using cached CUDA template processing results (${cached_count} files)") else() message(STATUS "🔄 Using cached CPU template processing results (${cached_count} files)") endif() else() # Execute selective rendering + template processing execute_template_processing_with_selective_rendering() set(CUSTOMOPS_GENERIC_SOURCES ${CUSTOMOPS_GENERIC_SOURCES} PARENT_SCOPE) endif() endfunction() message(STATUS "✅ Enhanced TemplateProcessing.cmake loaded with comprehensive type handling - functions defined, execution deferred")