Lines Matching refs:var_suffix
225 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelFieldData_Hip_gen() local
226 …std::string P_name = (is_tensor ? "P_1d" : "P") + var_suffix, Q_name = is_tensor ? "Q_1d… in CeedOperatorBuildKernelFieldData_Hip_gen()
255 code << tab << "const CeedInt dim" << var_suffix << " = " << dim << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
260 …code << tab << "const CeedInt P" << var_suffix << " = " << (basis == CEED_BASIS_NONE ? Q : P) << "… in CeedOperatorBuildKernelFieldData_Hip_gen()
264 code << tab << "const CeedInt num_comp" << var_suffix << " = " << num_comp << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
297 code << tab << "CeedScalar *s_B" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
303 code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
305 …code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Hip_gen()
306 …e << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
336 code << tab << "CeedScalar *s_B" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
342 code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
344 …code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Hip_gen()
345 …e << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
356 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
358 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
360 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Hip_gen()
361 …e << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
372 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
374 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
376 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Hip_gen()
377 …e << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
383 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
385 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
387 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << P_name << "*" << Q_name << (is_… in CeedOperatorBuildKernelFieldData_Hip_gen()
388 << (is_tensor ? "" : var_suffix) << "];\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
389 …ame << ", " << Q_name << (is_tensor ? "" : "*dim") << (is_tensor ? "" : var_suffix) << ">(data, G." in CeedOperatorBuildKernelFieldData_Hip_gen()
390 << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Hip_gen()
414 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelRestriction_Hip_gen() local
415 std::string P_name = (is_all_tensor ? "P_1d" : "P") + var_suffix; in CeedOperatorBuildKernelRestriction_Hip_gen()
440 code << tab << "CeedScalar *r_e" << var_suffix << " = " << buffer_name << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
444 …code << tab << "CeedScalar r_e" << var_suffix << "[num_comp" << var_suffix << "*" << P_name << "];… in CeedOperatorBuildKernelRestriction_Hip_gen()
447 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
454 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
456 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
458 …d" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", comp_stride" << var_suffix… in CeedOperatorBuildKernelRestriction_Hip_gen()
459 …P_name << ">(data, l_size" << var_suffix << ", elem, indices.inputs[" << i << "], d" << var_suffix… in CeedOperatorBuildKernelRestriction_Hip_gen()
474 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelRestriction_Hip_gen()
475 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
476 … "ReadLVecStrided" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", " << P_nam… in CeedOperatorBuildKernelRestriction_Hip_gen()
477 …<< var_suffix << "_0, strides" << var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, d"… in CeedOperatorBuildKernelRestriction_Hip_gen()
478 << var_suffix << ");\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
485 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
503 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
505 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
507 …d" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", comp_stride" << var_suffix… in CeedOperatorBuildKernelRestriction_Hip_gen()
508 …_name << ">(data, l_size" << var_suffix << ", elem, indices.outputs[" << i << "], r_e" << var_suff… in CeedOperatorBuildKernelRestriction_Hip_gen()
523 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelRestriction_Hip_gen()
524 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelRestriction_Hip_gen()
525 …"WriteLVecStrided" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", " << P_nam… in CeedOperatorBuildKernelRestriction_Hip_gen()
526 …<< var_suffix << "_0, strides" << var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, r_… in CeedOperatorBuildKernelRestriction_Hip_gen()
556 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelBasis_Hip_gen() local
557 …std::string P_name = (is_tensor ? "P_1d" : "P") + var_suffix, Q_name = is_tensor ? "Q_1d" … in CeedOperatorBuildKernelBasis_Hip_gen()
582 code << tab << "CeedScalar *r_q" << var_suffix << " = r_e" << var_suffix << ";\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
589 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (dim >= 3 ? Q… in CeedOperatorBuildKernelBasis_Hip_gen()
590 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
591 << ", s_B" << var_suffix << ", r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
598 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelBasis_Hip_gen()
599 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Hip_gen()
600 << var_suffix << ", s_B" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
607 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (dim >= 3 ? Q… in CeedOperatorBuildKernelBasis_Hip_gen()
608 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
609 << ", s_B" << var_suffix << ", r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
614 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << Q_name << "];… in CeedOperatorBuildKernelBasis_Hip_gen()
615 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
616 << ", s_B" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
624 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelBasis_Hip_gen()
626 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Hip_gen()
627 …<< var_suffix << ", s_B" << var_suffix << ", s_G" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
631 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelBasis_Hip_gen()
632 …code << tab << function_name << "<num_comp" << var_suffix << ", dim" << var_suffix << ", " << P_na… in CeedOperatorBuildKernelBasis_Hip_gen()
633 …<< ", OP_T_1D>(data, r_e" << var_suffix << ", s_G" << var_suffix << ", r_q" << var_suffix << ");\n… in CeedOperatorBuildKernelBasis_Hip_gen()
645 …code << tab << "CeedScalar r_q" << var_suffix << "[" << (is_all_tensor && (dim >= 3) ? Q_name : "1… in CeedOperatorBuildKernelBasis_Hip_gen()
648 …ab << function_name << "<" << P_name << ", " << Q_name << ">(data, W, r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
661 code << tab << "CeedScalar *r_e" << var_suffix << " = r_q" << var_suffix << ";\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
664 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
668 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
669 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
677 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Hip_gen()
678 << var_suffix << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
682 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
686 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
687 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
692 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Hip_gen()
693 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
702 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Hip_gen()
703 …<< var_suffix << ", s_B" << var_suffix << ", s_G" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Hip_gen()
707 …code << tab << function_name << "<num_comp" << var_suffix << ", dim" << var_suffix << ", " << P_na… in CeedOperatorBuildKernelBasis_Hip_gen()
708 …<< ", OP_T_1D>(data, r_q" << var_suffix << ", s_G" << var_suffix << ", r_e" << var_suffix << ");\n… in CeedOperatorBuildKernelBasis_Hip_gen()
742 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
750 code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
752 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelQFunction_Hip_gen()
759 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3… in CeedOperatorBuildKernelQFunction_Hip_gen()
760 …or (CeedInt i = 0; i < num_comp" << var_suffix << "*" << (max_dim >= 3 ? Q_name : "1") << "; i++) … in CeedOperatorBuildKernelQFunction_Hip_gen()
763 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelQFunction_Hip_gen()
770 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3… in CeedOperatorBuildKernelQFunction_Hip_gen()
771 …or (CeedInt i = 0; i < num_comp" << var_suffix << "*" << (max_dim >= 3 ? Q_name : "1") << "; i++) … in CeedOperatorBuildKernelQFunction_Hip_gen()
775 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << Q_name << "];… in CeedOperatorBuildKernelQFunction_Hip_gen()
776 …<< tab << "for (CeedInt i = 0; i < num_comp" << var_suffix << "*" << Q_name << "; i++) r_q" << var… in CeedOperatorBuildKernelQFunction_Hip_gen()
778 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Hip_gen()
809 std::string var_suffix = "_in_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
810 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Hip_gen()
819 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
820 code << tab << "ReadPoint<num_comp" << var_suffix << ", comp_stride" << var_suffix in CeedOperatorBuildKernelQFunction_Hip_gen()
821 …lem, p, max_num_points, indices.inputs[" << i << "], d" << var_suffix << ", r_s" << var_suffix << … in CeedOperatorBuildKernelQFunction_Hip_gen()
824 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
825 …code << tab << "InterpAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_points, " <… in CeedOperatorBuildKernelQFunction_Hip_gen()
826 << ">(data, i, r_c" << var_suffix << ", r_x, r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
829 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Hip_gen()
830 …code << tab << "GradAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_points, " << … in CeedOperatorBuildKernelQFunction_Hip_gen()
831 << ">(data, i, r_c" << var_suffix << ", r_x, r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
834 code << tab << "CeedScalar r_s" << var_suffix << "[1];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
835 code << tab << "r_s" << var_suffix << "[0] = 1.0;\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
848 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
856 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
859 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
862 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Hip_gen()
884 std::string var_suffix = "_in_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
895 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
911 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelQFunction_Hip_gen()
912 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
913 … << tab << "ReadEVecSliceStrided3d<num_comp" << var_suffix << ", " << Q_name << ", strides" << var… in CeedOperatorBuildKernelQFunction_Hip_gen()
914 …<< var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, q, d" << var_suffix << ", r_s" <<… in CeedOperatorBuildKernelQFunction_Hip_gen()
921 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
923 … code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
926 …code << tab << "ReadEVecSliceStandard3d<num_comp" << var_suffix << ", comp_stride" << var_suffix <… in CeedOperatorBuildKernelQFunction_Hip_gen()
927 …<< var_suffix << ", elem, q, indices.inputs[" << i << "], d" << var_suffix << ", r_s" << var_suffi… in CeedOperatorBuildKernelQFunction_Hip_gen()
932 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
933 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "; j++) {\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
935 … code << tab << "r_s" << var_suffix << "[j] = r_q" << var_suffix << "[q + j*" << Q_name << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
940 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Hip_gen()
941 …<< tab << "GradColloSlice3d<num_comp" << var_suffix << ", " << Q_name << ", OP_T_1D>(data, q, r_q"… in CeedOperatorBuildKernelQFunction_Hip_gen()
942 << var_suffix << ", r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
945 code << tab << "CeedScalar r_s" << var_suffix << "[1];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
946 code << tab << "r_s" << var_suffix << "[0] = r_q" << var_suffix << "[q];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
959 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
967 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
970 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
973 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Hip_gen()
1046 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
1047 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Hip_gen()
1063 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1064 code << tab << "WritePoint<num_comp" << var_suffix << ", comp_stride" << var_suffix in CeedOperatorBuildKernelQFunction_Hip_gen()
1066 << ", r_s" << var_suffix << ", d" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1072 …code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "; j++) r_s" << var_suffix << "… in CeedOperatorBuildKernelQFunction_Hip_gen()
1075 …code << tab << "InterpTransposeAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_po… in CeedOperatorBuildKernelQFunction_Hip_gen()
1076 << ">(data, i, r_s" << var_suffix << ", r_x, r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1081 …ab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "*dim" << var_suffix << "; j++) r_s" << … in CeedOperatorBuildKernelQFunction_Hip_gen()
1084 …code << tab << "GradTransposeAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_poin… in CeedOperatorBuildKernelQFunction_Hip_gen()
1085 << ">(data, i, r_s" << var_suffix << ", r_x, r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1102 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Hip_gen() local
1103 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Hip_gen()
1112 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << " ; j++) {\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1114 … code << tab << "r_q" << var_suffix << "[q + j*" << Q_name << "] = r_s" << var_suffix << "[j];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1119 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << " ; j++) {\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1121 … code << tab << "r_q" << var_suffix << "[q + j*" << Q_name << "] = r_s" << var_suffix << "[j];\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1126 … "GradColloSliceTranspose3d<num_comp" << var_suffix << ", " << Q_name << ", OP_T_1D>(data, q, r_s"… in CeedOperatorBuildKernelQFunction_Hip_gen()
1127 << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Hip_gen()
1958 std::string var_suffix = "_in_" + std::to_string(f); in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen() local
1963 …ode << tab << "CeedScalar r_e" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3 … in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
1966 … code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_in_" << active_field_index << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2043 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen() local
2050 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2052 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2053 …d" << max_dim << "d_Assembly<num_comp" << var_suffix << ", comp_stride" << var_suffix << ", P_1d" … in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2054 … << ">(data, l_size" << var_suffix << ", elem, n, r_e" << var_suffix << ", values_array);\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2057 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen() local
2064 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2066 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2067 …ard" << max_dim << "d_Single<num_comp" << var_suffix << ", comp_stride" << var_suffix << ", P_1d" … in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()
2068 …<< ">(data, l_size" << var_suffix << ", elem, n, indices.outputs[" << i << "], r_e" << var_suffix … in CeedOperatorBuildKernelAssemblyAtPoints_Hip_gen()