Lines Matching refs:var_suffix
198 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelFieldData_Cuda_gen() local
199 …std::string P_name = (is_tensor ? "P_1d" : "P") + var_suffix, Q_name = is_tensor ? "Q_1… in CeedOperatorBuildKernelFieldData_Cuda_gen()
228 code << tab << "const CeedInt dim" << var_suffix << " = " << dim << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
233 …code << tab << "const CeedInt P" << var_suffix << " = " << (basis == CEED_BASIS_NONE ? Q : P) << "… in CeedOperatorBuildKernelFieldData_Cuda_gen()
237 code << tab << "const CeedInt num_comp" << var_suffix << " = " << num_comp << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
270 code << tab << "CeedScalar *s_B" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
276 code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
278 …code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Cuda_gen()
279 …e << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
309 code << tab << "CeedScalar *s_B" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
315 code << tab << "CeedScalar *s_B" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
317 …code << tab << "__shared__ CeedScalar s_B" << var_suffix << "[" << P_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Cuda_gen()
318 …e << ", " << Q_name << ">(data, B." << option_name << "[" << i << "], s_B" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
329 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
331 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
333 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Cuda_gen()
334 …e << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
345 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
347 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
349 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << Q_name << "*" << Q_name << "];\… in CeedOperatorBuildKernelFieldData_Cuda_gen()
350 …e << ", " << Q_name << ">(data, G." << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
356 code << tab << "CeedScalar *s_G" << var_suffix << " = " << reuse_var << ";\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
358 code << tab << "CeedScalar *s_G" << var_suffix << " = NULL;\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
360 …code << tab << "__shared__ CeedScalar s_G" << var_suffix << "[" << P_name << "*" << Q_name << (is_… in CeedOperatorBuildKernelFieldData_Cuda_gen()
361 << (is_tensor ? "" : var_suffix) << "];\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
362 …ame << ", " << Q_name << (is_tensor ? "" : "*dim") << (is_tensor ? "" : var_suffix) << ">(data, G." in CeedOperatorBuildKernelFieldData_Cuda_gen()
363 << option_name << "[" << i << "], s_G" << var_suffix << ");\n"; in CeedOperatorBuildKernelFieldData_Cuda_gen()
387 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelRestriction_Cuda_gen() local
388 std::string P_name = (is_all_tensor ? "P_1d" : "P") + var_suffix; in CeedOperatorBuildKernelRestriction_Cuda_gen()
413 code << tab << "CeedScalar *r_e" << var_suffix << " = " << buffer_name << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
417 …code << tab << "CeedScalar r_e" << var_suffix << "[num_comp" << var_suffix << "*" << P_name << "];… in CeedOperatorBuildKernelRestriction_Cuda_gen()
420 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
427 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
429 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
431 …d" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", comp_stride" << var_suffix… in CeedOperatorBuildKernelRestriction_Cuda_gen()
432 …P_name << ">(data, l_size" << var_suffix << ", elem, indices.inputs[" << i << "], d" << var_suffix… in CeedOperatorBuildKernelRestriction_Cuda_gen()
447 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelRestriction_Cuda_gen()
448 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
449 … "ReadLVecStrided" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", " << P_nam… in CeedOperatorBuildKernelRestriction_Cuda_gen()
450 …<< var_suffix << "_0, strides" << var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, d"… in CeedOperatorBuildKernelRestriction_Cuda_gen()
451 << var_suffix << ");\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
458 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
476 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
478 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
480 …d" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", comp_stride" << var_suffix… in CeedOperatorBuildKernelRestriction_Cuda_gen()
481 …_name << ">(data, l_size" << var_suffix << ", elem, indices.outputs[" << i << "], r_e" << var_suff… in CeedOperatorBuildKernelRestriction_Cuda_gen()
496 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelRestriction_Cuda_gen()
497 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelRestriction_Cuda_gen()
498 …"WriteLVecStrided" << (is_all_tensor ? max_dim : 1) << "d<num_comp" << var_suffix << ", " << P_nam… in CeedOperatorBuildKernelRestriction_Cuda_gen()
499 …<< var_suffix << "_0, strides" << var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, r_… in CeedOperatorBuildKernelRestriction_Cuda_gen()
529 std::string var_suffix = (is_input ? "_in_" : "_out_") + std::to_string(i); in CeedOperatorBuildKernelBasis_Cuda_gen() local
530 …std::string P_name = (is_tensor ? "P_1d" : "P") + var_suffix, Q_name = is_tensor ? "Q_1d" … in CeedOperatorBuildKernelBasis_Cuda_gen()
555 code << tab << "CeedScalar *r_q" << var_suffix << " = r_e" << var_suffix << ";\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
562 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (dim >= 3 ? Q… in CeedOperatorBuildKernelBasis_Cuda_gen()
563 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
564 << ", s_B" << var_suffix << ", r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
571 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelBasis_Cuda_gen()
572 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Cuda_gen()
573 << var_suffix << ", s_B" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
580 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (dim >= 3 ? Q… in CeedOperatorBuildKernelBasis_Cuda_gen()
581 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
582 << ", s_B" << var_suffix << ", r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
587 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << Q_name << "];… in CeedOperatorBuildKernelBasis_Cuda_gen()
588 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
589 << ", s_B" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
597 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelBasis_Cuda_gen()
599 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Cuda_gen()
600 …<< var_suffix << ", s_B" << var_suffix << ", s_G" << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
604 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelBasis_Cuda_gen()
605 …code << tab << function_name << "<num_comp" << var_suffix << ", dim" << var_suffix << ", " << P_na… in CeedOperatorBuildKernelBasis_Cuda_gen()
606 …<< ", OP_T_1D>(data, r_e" << var_suffix << ", s_G" << var_suffix << ", r_q" << var_suffix << ");\n… in CeedOperatorBuildKernelBasis_Cuda_gen()
618 …code << tab << "CeedScalar r_q" << var_suffix << "[" << (is_all_tensor && (dim >= 3) ? Q_name : "1… in CeedOperatorBuildKernelBasis_Cuda_gen()
621 …ab << function_name << "<" << P_name << ", " << Q_name << ">(data, W, r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
634 code << tab << "CeedScalar *r_e" << var_suffix << " = r_q" << var_suffix << ";\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
637 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
641 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
642 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
650 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Cuda_gen()
651 << var_suffix << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
655 code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_scratch;\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
659 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
660 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
665 … function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << ", OP_T_1D>(dat… in CeedOperatorBuildKernelBasis_Cuda_gen()
666 << ", s_B" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
675 …code << tab << function_name << "<num_comp" << var_suffix << ", " << P_name << ", " << Q_name << "… in CeedOperatorBuildKernelBasis_Cuda_gen()
676 …<< var_suffix << ", s_B" << var_suffix << ", s_G" << var_suffix << ", r_e" << var_suffix << ");\n"; in CeedOperatorBuildKernelBasis_Cuda_gen()
680 …code << tab << function_name << "<num_comp" << var_suffix << ", dim" << var_suffix << ", " << P_na… in CeedOperatorBuildKernelBasis_Cuda_gen()
681 …<< ", OP_T_1D>(data, r_q" << var_suffix << ", s_G" << var_suffix << ", r_e" << var_suffix << ");\n… in CeedOperatorBuildKernelBasis_Cuda_gen()
715 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
723 code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
725 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelQFunction_Cuda_gen()
732 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3… in CeedOperatorBuildKernelQFunction_Cuda_gen()
733 …or (CeedInt i = 0; i < num_comp" << var_suffix << "*" << (max_dim >= 3 ? Q_name : "1") << "; i++) … in CeedOperatorBuildKernelQFunction_Cuda_gen()
736 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << (is_all_tenso… in CeedOperatorBuildKernelQFunction_Cuda_gen()
743 …code << tab << "CeedScalar r_c" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3… in CeedOperatorBuildKernelQFunction_Cuda_gen()
744 …or (CeedInt i = 0; i < num_comp" << var_suffix << "*" << (max_dim >= 3 ? Q_name : "1") << "; i++) … in CeedOperatorBuildKernelQFunction_Cuda_gen()
748 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*" << Q_name << "];… in CeedOperatorBuildKernelQFunction_Cuda_gen()
749 …<< tab << "for (CeedInt i = 0; i < num_comp" << var_suffix << "*" << Q_name << "; i++) r_q" << var… in CeedOperatorBuildKernelQFunction_Cuda_gen()
751 …code << tab << "CeedScalar r_q" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Cuda_gen()
782 std::string var_suffix = "_in_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
783 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Cuda_gen()
792 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
793 code << tab << "ReadPoint<num_comp" << var_suffix << ", comp_stride" << var_suffix in CeedOperatorBuildKernelQFunction_Cuda_gen()
794 …lem, p, max_num_points, indices.inputs[" << i << "], d" << var_suffix << ", r_s" << var_suffix << … in CeedOperatorBuildKernelQFunction_Cuda_gen()
797 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
798 …code << tab << "InterpAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_points, " <… in CeedOperatorBuildKernelQFunction_Cuda_gen()
799 << ">(data, i, r_c" << var_suffix << ", r_x, r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
802 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Cuda_gen()
803 …code << tab << "GradAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_points, " << … in CeedOperatorBuildKernelQFunction_Cuda_gen()
804 << ">(data, i, r_c" << var_suffix << ", r_x, r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
807 code << tab << "CeedScalar r_s" << var_suffix << "[1];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
808 code << tab << "r_s" << var_suffix << "[0] = 1.0;\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
821 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
829 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
832 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
835 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Cuda_gen()
857 std::string var_suffix = "_in_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
868 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
884 …code << tab << "const CeedInt strides" << var_suffix << "_0 = " << strides[0] << ", strides" << va… in CeedOperatorBuildKernelQFunction_Cuda_gen()
885 << ", strides" << var_suffix << "_2 = " << strides[2] << ";\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
886 … << tab << "ReadEVecSliceStrided3d<num_comp" << var_suffix << ", " << Q_name << ", strides" << var… in CeedOperatorBuildKernelQFunction_Cuda_gen()
887 …<< var_suffix << "_1, strides" << var_suffix << "_2>(data, elem, q, d" << var_suffix << ", r_s" <<… in CeedOperatorBuildKernelQFunction_Cuda_gen()
894 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
896 … code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
899 …code << tab << "ReadEVecSliceStandard3d<num_comp" << var_suffix << ", comp_stride" << var_suffix <… in CeedOperatorBuildKernelQFunction_Cuda_gen()
900 …<< var_suffix << ", elem, q, indices.inputs[" << i << "], d" << var_suffix << ", r_s" << var_suffi… in CeedOperatorBuildKernelQFunction_Cuda_gen()
905 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
906 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "; j++) {\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
908 … code << tab << "r_s" << var_suffix << "[j] = r_q" << var_suffix << "[q + j*" << Q_name << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
913 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Cuda_gen()
914 …<< tab << "GradColloSlice3d<num_comp" << var_suffix << ", " << Q_name << ", OP_T_1D>(data, q, r_q"… in CeedOperatorBuildKernelQFunction_Cuda_gen()
915 << var_suffix << ", r_s" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
918 code << tab << "CeedScalar r_s" << var_suffix << "[1];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
919 code << tab << "r_s" << var_suffix << "[0] = r_q" << var_suffix << "[q];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
932 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
940 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
943 code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
946 …code << tab << "CeedScalar r_s" << var_suffix << "[num_comp" << var_suffix << "*dim" << var_suffix… in CeedOperatorBuildKernelQFunction_Cuda_gen()
1019 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
1020 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1036 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1037 code << tab << "WritePoint<num_comp" << var_suffix << ", comp_stride" << var_suffix in CeedOperatorBuildKernelQFunction_Cuda_gen()
1039 << ", r_s" << var_suffix << ", d" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1045 …code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "; j++) r_s" << var_suffix << "… in CeedOperatorBuildKernelQFunction_Cuda_gen()
1048 …code << tab << "InterpTransposeAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_po… in CeedOperatorBuildKernelQFunction_Cuda_gen()
1049 << ">(data, i, r_s" << var_suffix << ", r_x, r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1054 …ab << "for (CeedInt j = 0; j < num_comp" << var_suffix << "*dim" << var_suffix << "; j++) r_s" << … in CeedOperatorBuildKernelQFunction_Cuda_gen()
1057 …code << tab << "GradTransposeAtPoints" << max_dim << "d<num_comp" << var_suffix << ", max_num_poin… in CeedOperatorBuildKernelQFunction_Cuda_gen()
1058 << ">(data, i, r_s" << var_suffix << ", r_x, r_c" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1075 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelQFunction_Cuda_gen() local
1076 std::string P_name = "P_1d" + var_suffix; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1085 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << " ; j++) {\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1087 … code << tab << "r_q" << var_suffix << "[q + j*" << Q_name << "] = r_s" << var_suffix << "[j];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1092 code << tab << "for (CeedInt j = 0; j < num_comp" << var_suffix << " ; j++) {\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1094 … code << tab << "r_q" << var_suffix << "[q + j*" << Q_name << "] = r_s" << var_suffix << "[j];\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1099 … "GradColloSliceTranspose3d<num_comp" << var_suffix << ", " << Q_name << ", OP_T_1D>(data, q, r_s"… in CeedOperatorBuildKernelQFunction_Cuda_gen()
1100 << var_suffix << ", r_q" << var_suffix << ");\n"; in CeedOperatorBuildKernelQFunction_Cuda_gen()
1953 std::string var_suffix = "_in_" + std::to_string(f); in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen() local
1958 …ode << tab << "CeedScalar r_e" << var_suffix << "[num_comp" << var_suffix << "*" << (max_dim >= 3 … in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
1961 … code << tab << "CeedScalar *r_e" << var_suffix << " = r_e_in_" << active_field_index << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2038 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen() local
2045 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2047 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2048 …d" << max_dim << "d_Assembly<num_comp" << var_suffix << ", comp_stride" << var_suffix << ", P_1d" … in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2049 … << ">(data, l_size" << var_suffix << ", elem, n, r_e" << var_suffix << ", values_array);\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2052 std::string var_suffix = "_out_" + std::to_string(i); in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen() local
2059 code << tab << "const CeedInt l_size" << var_suffix << " = " << l_size << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2061 code << tab << "const CeedInt comp_stride" << var_suffix << " = " << comp_stride << ";\n"; in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2062 …ard" << max_dim << "d_Single<num_comp" << var_suffix << ", comp_stride" << var_suffix << ", P_1d" … in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()
2063 …<< ">(data, l_size" << var_suffix << ", elem, n, indices.outputs[" << i << "], r_e" << var_suffix … in CeedOperatorBuildKernelAssemblyAtPoints_Cuda_gen()