Home
last modified time | relevance | path

Searched refs:crd (Results 1 – 25 of 56) sorted by relevance

123

/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/common/
Datomcrdtocv_impl.cu27 __global__ void Record_Box_Map_Times(int atom_numbers, const float *crd, const float *old_crd, floa… in Record_Box_Map_Times() argument
32 if (crd[3 * i + 0] - old_crd[3 * i + 0] > half_box[0]) { in Record_Box_Map_Times()
34 } else if (crd[3 * i + 0] - old_crd[3 * i + 0] < -half_box[0]) { in Record_Box_Map_Times()
37 if (crd[3 * i + 1] - old_crd[3 * i + 1] > half_box[1]) { in Record_Box_Map_Times()
39 } else if (crd[3 * i + 1] - old_crd[3 * i + 1] < -half_box[1]) { in Record_Box_Map_Times()
42 if (crd[3 * i + 2] - old_crd[3 * i + 2] > half_box[2]) { in Record_Box_Map_Times()
44 } else if (crd[3 * i + 2] - old_crd[3 * i + 2] < -half_box[2]) { in Record_Box_Map_Times()
50 __global__ void gen_nowarp_crd(int atom_numbers, const float *crd, float *box, int *box_map_times, … in gen_nowarp_crd() argument
53 nowarp_crd[3 * i + 0] = static_cast<float>(box_map_times[3 * i + 0]) * box[0] + crd[3 * i + 0]; in gen_nowarp_crd()
54 nowarp_crd[3 * i + 1] = static_cast<float>(box_map_times[3 * i + 1]) * box[1] + crd[3 * i + 1]; in gen_nowarp_crd()
[all …]
Dcrd_to_uint_crd_impl.cu20 …lobal__ void Crd_To_Uint_Crd(const int atom_numbers, const VECTOR *scale_factor, const VECTOR *crd, in Crd_To_Uint_Crd() argument
24 uint_crd[atom_i].uint_x = crd[atom_i].x * scale_factor[0].x; in Crd_To_Uint_Crd()
25 uint_crd[atom_i].uint_y = crd[atom_i].y * scale_factor[0].y; in Crd_To_Uint_Crd()
26 uint_crd[atom_i].uint_z = crd[atom_i].z * scale_factor[0].z; in Crd_To_Uint_Crd()
38 VECTOR *crd = const_cast<VECTOR *>(reinterpret_cast<const VECTOR *>(crd_f)); in CrdToUintCrd() local
45 atom_numbers, crd_to_uint_crd_cof, crd, uint_crd); in CrdToUintCrd()
Dgetcenter_impl.cu21 … const int *center_atoms, const VECTOR *crd, VECTOR *center_of_geometry) { in GetCenterOfGeometryKernel() argument
25 VECTOR temp = center_numbers_inverse * crd[atom_i]; in GetCenterOfGeometryKernel()
34 VECTOR *crd = const_cast<VECTOR *>(reinterpret_cast<const VECTOR *>(crd_f)); in GetCenterOfGeometry() local
37 center_numbers, center_numbers_inverse, center_atoms, crd, center_of_geometry); in GetCenterOfGeometry()
Dget_center_of_mass_impl.cu20 __global__ void Get_Center_Of_Mass(int residue_numbers, int *start, int *end, VECTOR *crd, float *a… in Get_Center_Of_Mass() argument
26 com_lin = com_lin + atom_mass[atom_i] * crd[atom_i]; in Get_Center_Of_Mass()
36 VECTOR *crd = reinterpret_cast<VECTOR *>(crd_f); in GetCenterOfMass() local
38 …Get_Center_Of_Mass<<<20, 32, 0, stream>>>(residue_numbers, start, end, crd, atom_mass, residue_mas… in GetCenterOfMass()
Dcrd_to_uint_crd_quarter_impl.cu19 …void Crd_To_Uint_Crd_Quarter(const int atom_numbers, const VECTOR *scale_factor, const VECTOR *crd, in Crd_To_Uint_Crd_Quarter() argument
24 VECTOR temp = crd[atom_i]; in Crd_To_Uint_Crd_Quarter()
41 VECTOR *crd = const_cast<VECTOR *>(reinterpret_cast<const VECTOR *>(crd_f)); in CrdToUintCrdQuarter() local
48 atom_numbers, crd_to_uint_crd_cof, crd, uint_crd); in CrdToUintCrdQuarter()
Dmap_center_of_mass_impl.cu21 … float *scaler, VECTOR *center_of_mass, VECTOR *box_length, VECTOR *no_wrap_crd, VECTOR *crd) { in Map_Center_Of_Mass() argument
34 crd[atom_i] = no_wrap_crd[atom_i] + trans_vec; in Map_Center_Of_Mass()
41 VECTOR *crd = reinterpret_cast<VECTOR *>(crd_f); in MapCenterOfMass() local
46 no_wrap_crd, crd); in MapCenterOfMass()
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/neighbor_list/
Dneighbor_list_impl.cu34 …tic __global__ void Crd_To_Uint_Crd(const int atom_numbers, float *scale_factor, const VECTOR *crd, in Crd_To_Uint_Crd() argument
39 VECTOR temp = crd[atom_i]; in Crd_To_Uint_Crd()
55 static __global__ void Crd_Periodic_Map(const int atom_numbers, VECTOR *crd, const float *box_lengt… in Crd_Periodic_Map() argument
58 if (crd[atom_i].x >= 0) { in Crd_Periodic_Map()
59 if (crd[atom_i].x < box_length[0]) { in Crd_Periodic_Map()
61 crd[atom_i].x = crd[atom_i].x - box_length[0]; in Crd_Periodic_Map()
64 crd[atom_i].x = crd[atom_i].x + box_length[0]; in Crd_Periodic_Map()
67 if (crd[atom_i].y >= 0) { in Crd_Periodic_Map()
68 if (crd[atom_i].y < box_length[1]) { in Crd_Periodic_Map()
70 crd[atom_i].y = crd[atom_i].y - box_length[1]; in Crd_Periodic_Map()
[all …]
Dneighbor_list_impl.cuh61 int *atom_in_grid_serial, GRID_BUCKET *bucket, float *crd, float *old_crd,
75 int *atom_in_grid_serial, GRID_BUCKET *bucket, float *crd, float *old_crd,
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/nvtit/
Dmd_iteration_gradient_descent_impl.cu21 __global__ void MD_Iteration_Gradient_Descent(const int atom_numbers, VECTOR *crd, VECTOR *frc, in MD_Iteration_Gradient_Descent() argument
25 crd[i].x = crd[i].x + learning_rate * frc[i].x; in MD_Iteration_Gradient_Descent()
26 crd[i].y = crd[i].y + learning_rate * frc[i].y; in MD_Iteration_Gradient_Descent()
27 crd[i].z = crd[i].z + learning_rate * frc[i].z; in MD_Iteration_Gradient_Descent()
35 void MDIterationGradientDescent(const int atom_numbers, float *crd, float *frc, const float learnin… in MDIterationGradientDescent() argument
37 VECTOR *d_crd = reinterpret_cast<VECTOR *>(crd); in MDIterationGradientDescent()
Dmd_iteration_leap_frog_impl.cu24 __global__ void MD_Iteration_Leap_Frog(const int atom_numbers, VECTOR *vel, VECTOR *crd, VECTOR *fr… in MD_Iteration_Leap_Frog() argument
36 crd[i].x = crd[i].x + dt * vel[i].x; in MD_Iteration_Leap_Frog()
37 crd[i].y = crd[i].y + dt * vel[i].y; in MD_Iteration_Leap_Frog()
38 crd[i].z = crd[i].z + dt * vel[i].z; in MD_Iteration_Leap_Frog()
46 void MDIterationLeapFrog(const int atom_numbers, float *vel, float *crd, float *frc, float *acc, in MDIterationLeapFrog() argument
49 VECTOR *d_crd = reinterpret_cast<VECTOR *>(crd); in MDIterationLeapFrog()
Dmd_iteration_leap_frog_with_max_vel_impl.cu21 …teration_Leap_Frog_With_Max_Velocity(const int atom_numbers, VECTOR *vel, VECTOR *crd, VECTOR *frc, in MD_Iteration_Leap_Frog_With_Max_Velocity() argument
30 crd[i] = crd[i] + dt * vel_i; in MD_Iteration_Leap_Frog_With_Max_Velocity()
35 void MDIterationLeapFrogWithMaxVelocity(const int atom_numbers, float *vel, float *crd, float *frc,… in MDIterationLeapFrogWithMaxVelocity() argument
39 VECTOR *d_crd = reinterpret_cast<VECTOR *>(crd); in MDIterationLeapFrogWithMaxVelocity()
Dmd_iteration_leap_frog_liujian_gpu_impl.cu23 … float *sqrt_mass_inverse, VECTOR *vel, VECTOR *crd, in MD_Iteration_Leap_Frog_With_LiuJian_kernel() argument
37 output[i].x = crd[i].x + half_dt * vel[i].x; in MD_Iteration_Leap_Frog_With_LiuJian_kernel()
38 output[i].y = crd[i].y + half_dt * vel[i].y; in MD_Iteration_Leap_Frog_With_LiuJian_kernel()
39 output[i].z = crd[i].z + half_dt * vel[i].z; in MD_Iteration_Leap_Frog_With_LiuJian_kernel()
53 … float *sqrt_mass_inverse, float *vel, float *crd, float *frc, float *acc, in MD_Iteration_Leap_Frog_With_LiuJian() argument
59 VECTOR *d_crd = reinterpret_cast<VECTOR *>(crd); in MD_Iteration_Leap_Frog_With_LiuJian()
Dmd_iteration_leap_frog_liujian_with_max_vel_impl.cu23 …const float *sqrt_mass_inverse, VECTOR *vel, VECTOR *crd, VECTOR *frc, VECTOR *acc, VECTOR *random… in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Velocity() argument
45 output[i].x = crd[i].x + half_dt * vel[i].x; in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Velocity()
46 output[i].y = crd[i].y + half_dt * vel[i].y; in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Velocity()
47 output[i].z = crd[i].z + half_dt * vel[i].z; in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Velocity()
64 … float *sqrt_mass_inverse, float *vel, float *crd, float *frc, in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Vel() argument
71 VECTOR *d_crd = reinterpret_cast<VECTOR *>(crd); in MD_Iteration_Leap_Frog_With_LiuJian_With_Max_Vel()
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/crdmcmap/
Dcal_no_wrap_crd_impl.cu22 …_ void Calculate_No_Wrap_Crd(int atom_numbers, INT_VECTOR *box_map_times, VECTOR *box, VECTOR *crd, in Calculate_No_Wrap_Crd() argument
25 nowrap_crd[i].x = static_cast<float>(box_map_times[i].int_x) * box[0].x + crd[i].x; in Calculate_No_Wrap_Crd()
26 nowrap_crd[i].y = static_cast<float>(box_map_times[i].int_y) * box[0].y + crd[i].y; in Calculate_No_Wrap_Crd()
27 nowrap_crd[i].z = static_cast<float>(box_map_times[i].int_z) * box[0].z + crd[i].z; in Calculate_No_Wrap_Crd()
37 VECTOR *crd = reinterpret_cast<VECTOR *>(crd_f); in calculatenowrapcrd() local
40 Calculate_No_Wrap_Crd<<<20, 256, 0, stream>>>(atom_numbers, box_map_times, box, crd, in calculatenowrapcrd()
Drefresh_boxmaptimes_impl.cu21 __global__ void Refresh_BoxMapTimes_CUDA(int atom_numbers, VECTOR *box_length_inverse, VECTOR *crd, in Refresh_BoxMapTimes_CUDA() argument
25 crd_i = crd[i]; in Refresh_BoxMapTimes_CUDA()
38 VECTOR *crd = reinterpret_cast<VECTOR *>(crd_f); in refresh_boxmaptimes() local
41 Refresh_BoxMapTimes_CUDA<<<1, 256, 0, stream>>>(atom_numbers, box_length_inverse, crd, in refresh_boxmaptimes()
/third_party/FreeBSD/sys/dev/usb/
Dusb_dev.c117 struct usb_cdev_refdata *crd, int need_uref) in usb_ref_device() argument
125 (void)memset_s(crd, sizeof(*crd), 0, sizeof(*crd)); in usb_ref_device()
159 crd->do_unlock = usbd_enum_lock(cpd->udev); in usb_ref_device()
166 crd->is_uref = 1; in usb_ref_device()
169 if (crd->do_unlock > 1) { in usb_ref_device()
170 crd->do_unlock = 0; in usb_ref_device()
183 crd->txfifo = f; in usb_ref_device()
184 crd->is_write = 1; /* ref */ in usb_ref_device()
192 crd->is_usbfs = 1; in usb_ref_device()
200 crd->rxfifo = f; in usb_ref_device()
[all …]
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/simple_constrain/
Drefresh_crd_vel_impl.cu21 … float half_exp_gamma_plus_half, VECTOR *test_frc, float *mass_inverse, VECTOR *crd, in Refresh_Crd_Vel() argument
25 VECTOR crd_lin = crd[atom_i]; in Refresh_Crd_Vel()
42 crd[atom_i] = crd_lin; in Refresh_Crd_Vel()
51 VECTOR *crd = reinterpret_cast<VECTOR *>(crd_f); in refreshcrdvel() local
56 …atom_numbers, dt_inverse, dt, exp_gamma, half_exp_gamma_plus_half, test_frc, mass_inverse, crd, ve… in refreshcrdvel()
Drefresh_uint_crd_impl.cu20 __global__ void Refresh_Uint_Crd(int atom_numbers, const VECTOR *crd, const VECTOR *quarter_crd_to_… in Refresh_Uint_Crd() argument
26 VECTOR crd_lin = crd[atom_i]; in Refresh_Uint_Crd()
49 const VECTOR *crd = reinterpret_cast<const VECTOR *>(crd_f); in refreshuintcrd() local
54 Refresh_Uint_Crd<<<block_per_grid, thread_per_block, 0, stream>>>(atom_numbers, crd, in refreshuintcrd()
Dconstrain_force_virial_impl.cu81 __global__ void refresh_uint_crd_update_kernel(int atom_numbers, const VECTOR *crd, in refresh_uint_crd_update_kernel() argument
88 VECTOR crd_lin = crd[atom_i]; in refresh_uint_crd_update_kernel()
159 const VECTOR *crd = reinterpret_cast<const VECTOR *>(crd_f); in refresh_uint_crd_update() local
165 …atom_numbers, crd, quarter_crd_to_uint_crd_cof, uint_crd, test_frc, mass_inverse, half_exp_gamma_p… in refresh_uint_crd_update()
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/sponge/restrain/
Drestrain_energy_impl.cu21 …oid restrain_energy_kernel(const int restrain_numbers, const int *restrain_list, const VECTOR *crd, in restrain_energy_kernel() argument
26 VECTOR dr = Get_Periodic_Displacement(crd_ref[atom_i], crd[atom_i], boxlength); in restrain_energy_kernel()
36 const VECTOR *crd = reinterpret_cast<const VECTOR *>(crd_f); in restrainenergy() local
39 …nel<<<block_per_grid, thread_per_block, 0, stream>>>(restrain_numbers, restrain_list, crd, crd_ref, in restrainenergy()
Drestrain_force_atom_energy_virial_impl.cu21 … const VECTOR *crd, const VECTOR *crd_ref, const float weight, in restrain_force_with_atom_energy_and_virial() argument
27 VECTOR dr = Get_Periodic_Displacement(crd_ref[atom_i], crd[atom_i], boxlength[0]); in restrain_force_with_atom_energy_and_virial()
47 const VECTOR *crd = reinterpret_cast<const VECTOR *>(crd_f); in restrainforcewithatomenergyandvirial() local
52 restrain_numbers, restrain_list, crd, crd_ref, weight, boxlength, atom_ene, atom_virial, frc); in restrainforcewithatomenergyandvirial()
/third_party/mesa3d/src/gallium/drivers/nouveau/codegen/
Dnv50_ir_lowering_gm107.cpp121 Value *crd[3], *arr, *shadow; in handleManualTXD() local
133 crd[c] = bld.getScratch(); in handleManualTXD()
153 bld.mkOp3(OP_SHFL, TYPE_F32, crd[c], i->getSrc(c + array), lane, quad); in handleManualTXD()
159 add = bld.mkOp2(OP_QUADOP, TYPE_F32, crd[c], tmp, crd[c]); in handleManualTXD()
167 add = bld.mkOp2(OP_QUADOP, TYPE_F32, crd[c], tmp, crd[c]); in handleManualTXD()
175 src[c] = bld.mkOp1v(OP_ABS, TYPE_F32, bld.getSSA(), crd[c]); in handleManualTXD()
181 src[c] = bld.mkOp2v(OP_MUL, TYPE_F32, bld.getSSA(), crd[c], val); in handleManualTXD()
184 src[c] = crd[c]; in handleManualTXD()
/third_party/mindspore/mindspore/ops/operations/
Dsponge_update_ops.py543 def infer_shape(self, crd_to_uint_crd_cof, crd): argument
546 validator.check_int(len(crd), 2, Rel.EQ, "crd_dim", cls_name)
549 validator.check_int(crd[0], n, Rel.EQ, "crd[0]", self.name)
550 validator.check_int(crd[1], 3, Rel.EQ, "crd[1]", self.name)
551 return crd
553 def infer_dtype(self, crd_to_uint_crd_cof, crd): argument
555 validator.check_tensor_dtype_valid('crd', crd, [mstype.float32], self.name)
629 … def infer_shape(self, inverse_mass, sqrt_mass_inverse, vel, crd, frc, acc, rand_state, rand_frc): argument
636 validator.check_int(len(crd), 2, Rel.EQ, "crd_dim", self.name)
643 validator.check_int(crd[0], n, Rel.EQ, "crd_shape[0]", self.name)
[all …]
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/sponge/nvtit/
Dmd_iteration_gradient_descent_kernel.h53 auto crd = GetDeviceAddress<float>(inputs, 0); in Launch() local
56 …MDIterationGradientDescent(atom_numbers, crd, frc, learning_rate, reinterpret_cast<cudaStream_t>(s… in Launch()
Dmd_iteration_leap_frog_kernel.h57 auto crd = GetDeviceAddress<float>(inputs, 1); in Launch() local
62 …MDIterationLeapFrog(atom_numbers, vel, crd, frc, acc, inverse_mass, dt, reinterpret_cast<cudaStrea… in Launch()

123