"csrc/gfx93/prefill/sparse/phase1.cuh" did not exist on "7a8722d7cde64ef222f39dc71e3a106ee47cb886"
  • Anton Gorenko's avatar
    Fix computeNonbonded hang on the HIP platform (#4959) · a4b43a04
    Anton Gorenko authored
    * Add a workaround for infinite loop in computeNonbonded (HIP)
    
    computeNonbonded hangs in some tests (without neighbor list).
    Reproducible on ROCm 6.4 and 6.4.1 (maybe on older versions too) on various architectures (both CDNA and RDNA).
    Affected tests: TestHipATMForce, TestHipMonteCarloBarostat, TestHipNonbondedForce, TestHipVirtualSites.
    
    Disassembly shows that the compiler splits branches of `if (skipBase+tgx < NUM_TILES_WITH_EXCLUSIONS)` and does
    `SHFL(skipTiles, TILE_SIZE-1) < pos` checks in them separately, even though `__builtin_amdgcn_ds_bpermute`
    is a convergent function. Apparently in this case not all lanes participate in each call.
    
    * Simplify includeTile check using ballot
    a4b43a04
nonbonded.hip 20.4 KB