1syntax = "proto3"; 2 3package xla.gpu; 4 5import "tensorflow/compiler/xla/stream_executor/dnn.proto"; 6import "tensorflow/compiler/xla/xla_data.proto"; 7 8// Backend configs for XLA:GPU. 9// 10// These are metadata that the GPU backend attaches to HloInstructions and later 11// uses during e.g. codegen. 12// 13// Remember that proto3 doesn't give clients a way to tell the difference 14// between a field not being present and a field having the default value. 15// Choose your defaults carefully. 16// 17// No guarantee is made about the stability of these protos. 18// 19// See HloInstruction::backend_config() for more info. 20 21// Backend config for a convolution that runs through cudnn. 22message CudnnConvBackendConfig { 23 reserved 1, 2; 24 25 // Opaque algorithm number and tuning knobs chosen for this conv. 26 stream_executor.dnn.AlgorithmProto algorithm = 6; 27 28 // The scaling factor multiplied with the convolution result. 29 double conv_result_scale = 4; 30 31 // Below are the fields related to cuDNN's fused convolution. Refer to 32 // GpuConvParams for their meanings. 33 34 // The requested activation (e.g. relu) after the convolution. It is with type 35 // stream_executor::dnn::ActivationMode. 36 int64 activation_mode = 3; 37 38 // The scaling factor multiplied with the side input. If no side input buffer 39 // is provided, this field must be 0. 40 double side_input_scale = 5; 41} 42 43// Backend config for the GEMM operation running through cuBLAS. 44message GemmBackendConfig { 45 // Opaque optional algorithm number. No chosen number indicates that a 46 // different cuBLAS API will be used, which does not allow for choosing an 47 // algorithm. 48 oneof algorithm { 49 int64 selected_algorithm = 1; 50 } 51 52 double alpha_real = 2; 53 double alpha_imag = 9; 54 55 double beta = 3; 56 57 xla.DotDimensionNumbers dot_dimension_numbers = 7; 58 59 xla.PrecisionConfig precision_config = 12; 60 61 // cublasLt matmul epilogue. 62 enum Epilogue { 63 DEFAULT = 0; 64 BIAS = 1; 65 } 66 67 Epilogue epilogue = 13; 68} 69 70// Backend config for bitcast operation generated from MLIR MHLO dialect. 71message BitcastBackendConfig { 72 LayoutProto source_layout = 1; 73 LayoutProto result_layout = 2; 74} 75