Dear Jason and everybody,
finally I found that this Amber bug comes from the file
$AMBERHOME/src/pmemd/src/cuda/kPMEInterpolation.cu.
There is the following hardcoded value (instead of getting it by asking
to the card) in the function kPMEFillChargeGridBuffer which I replaced
by the settings of my card (Tesla M2090):
+ (line ~400)
int lblocks = min(blocks, 65535);
kPMEFillChargeGridBuffer_kernel<<<lblocks, 64>>>(offset);
LAUNCHERROR("kPMEFillChargeGridBuffer");
offset += 65535;
blocks -= 65535;
-
int lblocks = min(blocks, 65536);
kPMEFillChargeGridBuffer_kernel<<<lblocks, 64>>>(offset);
LAUNCHERROR("kPMEFillChargeGridBuffer");
offset += 65536;
blocks -= 65536;
After this change, Amber continues its execution until the next error:
kPMEReduceChargeGridBuffer. I have also solved this error by modifying
the function kPMEReduceChargeGridBuffer and its cuda kernel function
kPMEReduceChargeGridBuffer_kernel. So my changes are:
+ (line ~166)
kPMEReduceChargeGridBuffer_kernel(int offset)
{
unsigned int pos = blockIdx.x *
blockDim.x + threadIdx.x + offset * blockDim.x;
-
kPMEReduceChargeGridBuffer_kernel()
{
unsigned int pos = blockIdx.x *
blockDim.x + threadIdx.x;
and
+ (line ~209)
long long blocks = (gpu->sim.nfft1 * gpu->sim.nfft2 * gpu->sim.nfft3
+ 127) >> 7;
int offset = 0;
while (blocks > 0)
{
long long lblocks = min(blocks,
65535ll);
kPMEReduceChargeGridBuffer_kernel<<<lblocks, 128>>>(offset);
LAUNCHERROR("kPMEReduceChargeGridBuffer");
offset += 65535;
blocks -= 65535;
}
-
unsigned int blocks = (gpu->sim.nfft1 * gpu->sim.nfft2 *
gpu->sim.nfft3 + 127) >> 7;
kPMEReduceChargeGridBuffer_kernel<<<blocks, 128>>>();
LAUNCHERROR("kPMEReduceChargeGridBuffer");
Now it seems to work and I got 0 errors in the amber cuda tests. But I
cannot ensure that this code will produce the right values for our
calculates, the execution is simply working.
Please, can you check your pmemd.cuda code in order to get it working
for cards with lower grid and block size? Then we will be very glad if
you make a new update with a tested patch solving these issues.
Thank you for your attention.
Best regards,
Pablo.
El 04/09/14 a las #4, Jason Swails escribió:
> On Thu, Sep 4, 2014 at 2:17 AM, Pablo Ródenas <pablo.rodenas.bsc.es> wrote:
>
>> Good morning,
>>
>> could you reproduce the problem with the files provided?
>>
>>
> O
> n my computer (GTX 680, 2 GB of memory), I get a memory allocation error
> because 2 GB is not enough for your system (ca. 700K+ atoms). When I move
> to a K20c (4 GB of memory), it runs fine for over 10 minutes (after which I
> killed it because your input files would have run for 10 hours on the
> K20c). That machine has the nVidia toolkit version 5.0 and the 331.38
> drivers on it.
>
> I'm not sure why you're having problems... Have you tried running the GPU
> validation suite? I know Ross Walker posted a link to it on a previous
> post, but I can't seem to locate it right now...
>
> HTH,
> Jason
>
--
Pablo Ródenas Barquero (pablo.rodenas.bsc.es)
BSC - Centro Nacional de Supercomputación
C/ Jordi Girona, 31 WWW: http://www.bsc.es
08034 Barcelona, Spain Tel: +34-93-405 42 29
e-mail: support.bsc.es Fax: +34-93-413 77 21
-----------------------------------------------
CNAG - Centre Nacional Anàlisi Genòmica
C/ Baldiri Reixac, 4 WWW: http://www.cnag.cat
08028 Barcelona, Spain Tel: +34-93-403 37 54
e-mail: cnag_support.bsc.es
-----------------------------------------------
WARNING / LEGAL TEXT: This message is intended only for the use of the
individual or entity to which it is addressed and may contain
information which is privileged, confidential, proprietary, or exempt
from disclosure under applicable law. If you are not the intended
recipient or the person responsible for delivering the message to the
intended recipient, you are strictly prohibited from disclosing,
distributing, copying, or in any way using this message. If you have
received this communication in error, please notify the sender and
destroy and delete any copies you may have received.
http://www.bsc.es/disclaimer
_______________________________________________
AMBER mailing list
AMBER.ambermd.org
http://lists.ambermd.org/mailman/listinfo/amber
Received on Thu Oct 09 2014 - 07:30:03 PDT
