Re: [AMBER] about "...gti_controlVariable.i for GaMD-PPI ..." from Christian Seitz via AMBER on 2025-10-20 (Amber Archive Oct 2025)

From: Christian Seitz via AMBER <amber.ambermd.org>
Date: Mon, 20 Oct 2025 22:22:08 +0000

Hi Patricio,

I tried one of your suggestions about removing the shared memory. Inside $AMBERHOME/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu I removed the line and while loop that you mentioned. I'm not sure I understood the final part you recommended changing. Inside the same file I changed

for (unsigned i = 0; i < listSize; i++) {
        PMEFloat dx = coordTI[i][0] - x0;
        PMEFloat dy = coordTI[i][1] - y0;
        PMEFloat dz = coordTI[i][2] - z0;

To this:

      for (unsigned i = 0; i < listSize; i++) {
        CSim.pImage(X,Y,Z)[iatom];

After doing that and trying to recompile, I find errors with this specific change:

CMake Error at pmemd_cuda_SPFP_generated_gti_NBList_kernels.cu.o.RELEASE.cmake:278 (message):
  Error generating file
  /lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/build/src/pmemd/src/cuda/CMakeFiles/pmemd_cuda_SPFP.dir//./pmemd_cuda_SPFP_generated_gti_NBList_kernels.cu.o

And later in the error logs:

/lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu(140): error: class "gti_simulationConst" has no member "pImage"
   cSim.pImage(X,Y,Z)[iatom]
        ^
/lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu(140): error: identifier "X" is undefined
   cSim.pImage(X,Y,Z)[iatom]
               ^
/lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu(140): error: identifier "Y" is undefined
   cSim.pImage(X,Y,Z)[iatom]
                 ^
/lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu(140): error: identifier "Z" is undefined
   cSim.pImage(X,Y,Z)[iatom]
                   ^
/lus/eagle/projects/FoundEpidem/cseitz/amber/pmemd24_src/src/pmemd/src/cuda/gti_NBList_kernels.cu(140): error: identifier "iatom" is undefined
   cSim.pImage(X,Y,Z)[iatom]

Should I be adding CSim.pImage(X,Y,Z)[iatom] a different way? Many thanks for your help!

Best,
Christian

Christian Seitz
PhD, Dept. of Computer Science | UChicago
cgseitz.uchicago.edu<mailto:cgseitz.uchicago.edu>
[www.linkedin.com/in/christianseitz21]<http://www.linkedin.com/in/christianseitz21>
________________________________
From: Christian Seitz via AMBER <amber.ambermd.org>
Sent: Wednesday, October 8, 2025 11:52
To: Miao, Yinglong <Yinglong_Miao.med.unc.edu>; taisung.gmail.com <taisung.gmail.com>; Patil Pranita Uttamrao via AMBER <amber.ambermd.org>
Cc: Adrian Roitberg <roitberg.ufl.edu>
Subject: Re: [AMBER] about "...gti_controlVariable.i for GaMD-PPI ..."

Hi Patricio, Adrian, Taisung, and Yinglong,

Thanks for the suggestions! Yes, as Yinglong said, I am trying to place a protein within the TI mask region so I can use GaMD-PPI to explore its interaction with another protein. I am certainly willing to provide test cases for Yinglong, Jinan or other GaMD developers interested in expanding GaMD-PPI, but rewriting the shared memory kernels and NB calculations are out of the scope of this project. Would it make sense to follow Patricio's suggestion and do away with the shared memory kernels and only use global memory kernels? I understand that speed would decrease significantly; this is probably ok as long as the results are still accurate. If not, what other MD simulation programs for protein-protein interactions would you recommend where one can see reversible association and dissociation for larger systems?

Best,
Christian

Christian Seitz
PhD, Dept. of Computer Science | UChicago
cgseitz.uchicago.edu<mailto:cgseitz.uchicago.edu>
[https://urldefense.com/v3/__http://www.linkedin.com/in/christianseitz21__;!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1QxSxpWiA$ ]<https://urldefense.com/v3/__http://www.linkedin.com/in/christianseitz21__;!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1QxSxpWiA$ >
________________________________
From: Miao, Yinglong <Yinglong_Miao.med.unc.edu>
Sent: Wednesday, October 8, 2025 05:46
To: taisung.gmail.com <taisung.gmail.com>
Cc: Adrian Roitberg <roitberg.ufl.edu>; Christian Seitz <cgseitz.uchicago.edu>
Subject: Re: [AMBER] about "...gti_controlVariable.i for GaMD-PPI ..."

Hi All,

I assume this is related to PPI-GaMD. A previous postdoc in my lab Jinan implemented the code in Amber. As we need to extract non-bonded interaction potential energies between two proteins for applying the GaMD boost, Jinan may use timask1 and scmask1 to select one of the proteins for the energy calculations. This may answer Taisung’s question why we want to include a fairly large number of atoms in the TI region. But there should be no alchemical change (as with TI) for the current PPI-GaMD simulations.

Hope this helps,
Yinglong

On Oct 8, 2025, at 8:33 AM, taisung.gmail.com wrote:

You don't often get email from taisung@gmail.com<mailto:taisung@gmail.com>. Learn why this is important<https://urldefense.com/v3/__https://aka.ms/LearnAboutSenderIdentification__;!!BpyFHLRN4TMTrA!-uJRThE51mzQJmcYrpvFoMt1KKBMRx6LU7n8LRqc-9mxyvS2BY3O4YbudqcHy9VpC06EjOV-IBoDgMJUjtv_moPxiHw_qyA$>
Hi Adrian and Christian,

               I think Patricio already answered the questions. The original intended usage for this part was only for TI simulations, where 5-100 atoms are typically defined in the special region. Hence, these kernels, which utilize shared memory for optimal performance, were never designed to handle thousands of atoms. To handle a large number of atoms, one may need to rewrite these kernels. The code piece Patricio described is just a small portion of the relevant kernels. You might need to go through all kernels to make sure everything is not only compliable but also producing the right results.

               The real question: why do you need so many atoms in the special region (TI or GaMD)? Depending on your specific problem, the current GPU-TI framework might not suit your needs. Not only the special region sizes, but the algorithm of NB calculations would also need to be rewritten. I expect the performance will be significantly impacted when you have more than 250 atoms in the special regions, even if you can use it.

Best,

Taisung

From: Adrian Roitberg <roitberg.ufl.edu>
Sent: Tuesday, October 7, 2025 12:37 PM
To: Christian Seitz <cgseitz.uchicago.edu>; Yinglong_Miao.med.unc.edu; Taisung Lee <taisung.gmail.com>
Subject: Fwd: [AMBER] about "...gti_controlVariable.i for GaMD-PPI ..."

Yinglong and Taisung, can one of you or your people help Christian please ?
Thanks !

-------- Forwarded Message --------
Subject:
Re: [AMBER] about "...gti_controlVariable.i for GaMD-PPI ..."
Date:
Tue, 7 Oct 2025 01:33:13 +0000
From:
Patricio Barletta via AMBER <amber.ambermd.org><mailto:amber.ambermd.org>
Reply-To:
Patricio Barletta <pb777.iqb.rutgers.edu><mailto:pb777.iqb.rutgers.edu>, AMBER Mailing List <amber.ambermd.org><mailto:amber.ambermd.org>
To:
Christian Seitz <cgseitz.uchicago.edu><mailto:cgseitz.uchicago.edu>, Carmen Al Masri Said via AMBER <amber.ambermd.org><mailto:amber.ambermd.org>

[External Email]

Yeah, those limits are too high then.

I glanced over the kernels that, I assume, are the ones giving you issues:

kgBuildSpecial2RestNBPreList_kernel(specialType type)
kgBuildSpecial2RestNBList_kernel(specialType type)
kgBuildTI2TINBList_kernel() {
kgBuildTI2TINBList_gamd2_kernel()

If you really need such a high atom count, then the easiest solution is to do away with the shared memory usage and use global memory:

Eg: at kgBuildSpecial2RestNBPreList_kernel()
```
__shared__ volatile float coordTI[gti_simulationConst::MaxNumberTIAtom][3]; # REMOVE THIS
...
while (th < listSize) { # REMOVE THIS WHOLE LOOP
unsigned iatom = cSim.pImageAtomLookup[pList[th]];
coordTI[th][0] = cSim.pImageX[iatom];
coordTI[th][1] = cSim.pImageY[iatom];
coordTI[th][2] = cSim.pImageZ[iatom];
th += blockDim.x;
}
...
PMEFloat dx = coordTI[i][0] - x0; # REPLACE WITH cSim.pImage(X,Y,Z)[iatom]
PMEFloat dy = coordTI[i][1] - y0;
PMEFloat dz = coordTI[i][2] - z0;
```

The performance hit on each kernel will be massive, but I have no idea how the overall performance will be impacted. I've never run these calculations.

The proper fix would be to determine the amount of shared memory, given the architecture, and issue these kernels serially until all atoms are considered. Are you in contact with the gamd developers? Perhaps this is something they'd be willing to collaborate on. I don't know them personally, but my PI (Darrin York) surely does.

Sorry I can't be of much help. I work on the alchemical side of things, so I try not to touch other features.

Best of luck.

________________________________
From: Christian Seitz <cgseitz.uchicago.edu><mailto:cgseitz.uchicago.edu>
Sent: Monday, October 6, 2025 6:41 PM
To: Patricio Barletta <pb777.iqb.rutgers.edu><mailto:pb777.iqb.rutgers.edu>; Carmen Al Masri Said via AMBER <amber.ambermd.org><mailto:amber.ambermd.org>
Subject: Re: about "...gti_controlVariable.i for GaMD-PPI ..."

Hi Patricio,

Many thanks for the suggestions! I am using NVIDIA A100s and have cuda11.8 - cuda 12.6. I used a fresh tarball of amber24, and edited the cmake/CudaConfig.cmake from

elseif((${CUDA_VERSION} VERSION_GREATER_EQUAL 11.8) AND (${CUDA_VERSION} VERSION_LESS 12.7)) message(STATUS "Configuring for SM5.0, SM5.2, SM5.3, SM6.0, SM6.1, SM7.0, SM7.5, SM8.0, SM8.6, and SM9.0") list(APPEND CUDA_NVCC_FLAGS ${SM50FLAGS} ${SM52FLAGS} ${SM53FLAGS} ${SM60FLAGS} ${SM61FLAGS} ${SM70FLAGS} ${SM75FLAGS} ${SM80FLAGS} ${SM86FLAGS} ${SM90FLAGS} -Wno-deprecated-gpu-targets -Wno-deprecated-declarations)

to

elseif((${CUDA_VERSION} VERSION_GREATER_EQUAL 11.8) AND (${CUDA_VERSION} VERSION_LESS 12.7))
message(STATUS "Configuring for SM5.0, SM5.2, SM5.3, SM6.0, SM6.1, SM7.0, SM7.5, SM8.0, SM8.6, and SM9.0")
list(APPEND CUDA_NVCC_FLAGS ${SM80FLAGS} -Wno-deprecated-gpu-targets -Wno-deprecated-declarations)

After doing "make install" I see the same errors as before. I can send the full terminal output if necessary, but I do not see any errors related to global constant data, thankfully - only shared data. Let me know if I'm misunderstanding your suggestions, or any other ideas you have. Thank you!

Best,
Christian

Christian Seitz
PhD, Dept. of Computer Science | UChicago
cgseitz.uchicago.edu<mailto:cgseitz.uchicago.edu><mailto:cgseitz.uchicago.edu><mailto:cgseitz.uchicago.edu>
[https://urldefense.com/v3/__https://nam10.safelinks.protection.outlook.com/?url=http*3A*2F*2Fwww.linkedin.com*2Fin*2Fchristianseitz21&data=05*7C02*7Croitberg*40ufl.edu*7Cac13fd7f1a3c40432f1f08de0541996f*7C0d4da0f84a314d76ace60a62331e1b84*7C0*7C0*7C638953976505975383*7CUnknown*7CTWFpbGZsb3d8eyJFbXB0eU1hcGkiOnRydWUsIlYiOiIwLjAuMDAwMCIsIlAiOiJXaW4zMiIsIkFOIjoiTWFpbCIsIldUIjoyfQ*3D*3D*7C0*7C*7C*7C&sdata=CEU2pnSYkT6n7Rxoy51aFL2PMTSRIw5NTjm2VsY5*2Bes*3D&reserved=0__;JSUlJSUlJSUlJSUlJSUlJSUlJSUlJSU!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1SsYDZIzg$ <https://urldefense.com/v3/__http://www.linkedin.com/in/christianseitz21__;!!BpyFHLRN4TMTrA!-uJRThE51mzQJmcYrpvFoMt1KKBMRx6LU7n8LRqc-9mxyvS2BY3O4YbudqcHy9VpC06EjOV-IBoDgMJUjtv_moPxkyyRa2c$>]<https://urldefense.com/v3/__https://nam10.safelinks.protection.outlook.com/?url=http*3A*2F*2Fwww.linkedin.com*2Fin*2Fchristianseitz21&data=05*7C02*7Croitberg*40ufl.edu*7Cac13fd7f1a3c40432f1f08de0541996f*7C0d4da0f84a314d76ace60a62331e1b84*7C0*7C0*7C638953976505998354*7CUnknown*7CTWFpbGZsb3d8eyJFbXB0eU1hcGkiOnRydWUsIlYiOiIwLjAuMDAwMCIsIlAiOiJXaW4zMiIsIkFOIjoiTWFpbCIsIldUIjoyfQ*3D*3D*7C0*7C*7C*7C&sdata=Y0n3BgDD16ubU*2B7VgJgJ7p61L8dieLVKj0sXvOb1H9g*3D&reserved=0__;JSUlJSUlJSUlJSUlJSUlJSUlJSUlJSU!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1SXOxHl_Q$ ><https://urldefense.com/v3/__http://www.linkedin.com/in/christianseitz21__;!!BpyFHLRN4TMTrA!-uJRThE51mzQJmcYrpvFoMt1KKBMRx6LU7n8LRqc-9mxyvS2BY3O4YbudqcHy9VpC06EjOV-IBoDgMJUjtv_moPxkyyRa2c$>
________________________________
From: Patricio Barletta <pb777.iqb.rutgers.edu><mailto:pb777.iqb.rutgers.edu>
Sent: Wednesday, October 1, 2025 12:25
To: Christian Seitz <cgseitz.uchicago.edu><mailto:cgseitz.uchicago.edu>
Subject: about "...gti_controlVariable.i for GaMD-PPI ..."

I only see errors in shared memory allocations, so that's good. You can also run into issues with the constant memory when doing these edits.

There's no way around the limit on constant memory, so if you also got errors like "... uses too much global constant data...", you'll have to lower those atom counts, or maybe zero out the sizes of some variables for other features (eg: MaxNumberRMSDAtom).

About the shared memory, maybe your GPU can deal with the increased usage, but since amber is being built for multiple targets, your compilation is breaking unnecessarily.

Check your cuda version and your GPU capability,<https://urldefense.com/v3/__https://developer.nvidia.com/cuda-gpus__;!!BpyFHLRN4TMTrA!-OoutJvvOt75yETJdP_OLYYwHYM6hQuCvcYF1JhJYF_tT0CAVw2gjRiF0O1mqqPehoD4Czx3kKDpvJZgXF0MAFOK$><https://urldefense.com/v3/__https:/developer.nvidia.com/cuda-gpus__;!!BpyFHLRN4TMTrA!-OoutJvvOt75yETJdP_OLYYwHYM6hQuCvcYF1JhJYF_tT0CAVw2gjRiF0O1mqqPehoD4Czx3kKDpvJZgXF0MAFOK$<https://urldefense.com/v3/__https://developer.nvidia.com/cuda-gpus__;!!BpyFHLRN4TMTrA!-OoutJvvOt75yETJdP_OLYYwHYM6hQuCvcYF1JhJYF_tT0CAVw2gjRiF0O1mqqPehoD4Czx3kKDpvJZgXF0MAFOK$%3E%3Chttps://urldefense.com/v3/__https:/developer.nvidia.com/cuda-gpus__;!!BpyFHLRN4TMTrA!-OoutJvvOt75yETJdP_OLYYwHYM6hQuCvcYF1JhJYF_tT0CAVw2gjRiF0O1mqqPehoD4Czx3kKDpvJZgXF0MAFOK$>> and edit cmake/CudaConfig.cmake

Say you have CUDA12 and and RTXA4500, you would go to:

```
elseif((${CUDA_VERSION} VERSION_GREATER_EQUAL 11.8) AND (${CUDA_VERSION} VERSION_LESS 12.7))
message(STATUS "Configuring for SM5.0, SM5.2, SM5.3, SM6.0, SM6.1, SM7.0, SM7.5, SM8.0, SM8.6, and SM9.0")
list(APPEND CUDA_NVCC_FLAGS ${SM50FLAGS} ${SM52FLAGS} ${SM53FLAGS} ${SM60FLAGS} ${SM61FLAGS} ${SM70FLAGS} ${SM75FLAGS} ${SM80FLAGS} ${SM86FLAGS} ${SM90FLAGS} -Wno-deprecated-gpu-targets -Wno-deprecated-declarations)

```
And delete all those ${SM*FLAGS} and leave only ${SM86FLAGS}.

Let me know how it goes.
_______________________________________________
AMBER mailing list
AMBER.ambermd.org<mailto:AMBER.ambermd.org>
https://urldefense.com/v3/__https://nam10.safelinks.protection.outlook.com/?url=http*3A*2F*2Flists.ambermd.org*2Fmailman*2Flistinfo*2Famber&data=05*7C02*7Croitberg*40ufl.edu*7Cac13fd7f1a3c40432f1f08de0541996f*7C0d4da0f84a314d76ace60a62331e1b84*7C0*7C0*7C638953976506017278*7CUnknown*7CTWFpbGZsb3d8eyJFbXB0eU1hcGkiOnRydWUsIlYiOiIwLjAuMDAwMCIsIlAiOiJXaW4zMiIsIkFOIjoiTWFpbCIsIldUIjoyfQ*3D*3D*7C0*7C*7C*7C&sdata=SdbxtnFf6nSvc2*2FXzJSHmk*2Fp*2BnOnUIdXTqSMdP3RX78*3D&reserved=0__;JSUlJSUlJSUlJSUlJSUlJSUlJSUlJSUlJSU!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1QazEmieg$ <https://urldefense.com/v3/__http://lists.ambermd.org/mailman/listinfo/amber__;!!BpyFHLRN4TMTrA!-uJRThE51mzQJmcYrpvFoMt1KKBMRx6LU7n8LRqc-9mxyvS2BY3O4YbudqcHy9VpC06EjOV-IBoDgMJUjtv_moPxWCxNPKg$>

_______________________________________________
AMBER mailing list
AMBER.ambermd.org
https://urldefense.com/v3/__http://lists.ambermd.org/mailman/listinfo/amber__;!!BpyFHLRN4TMTrA!8uZWQyiHSxQSsjigF1bwnbmQH45m4j4WdwbWVfc2_z3n2df5BAkFyGk_lVCNNMysn26nqexIG1Q2b48mJQ$
_______________________________________________
AMBER mailing list
AMBER.ambermd.org
http://lists.ambermd.org/mailman/listinfo/amber
Received on Mon Oct 20 2025 - 15:30:02 PDT