You missed this part:
unsigned int blocks = (gpu->sim.nfft1 * gpu->sim.nfft2 * gpu->sim.nfft3 +
255) >> 8;
On Thu, Oct 9, 2014 at 8:54 AM, Pablo Ródenas <pablo.rodenas.bsc.es> wrote:
> Dear Scott,
>
> I tried with both changes:
> kPMEReduceChargeGridBuffer_kernel<<<blocks, 255>>>();
> and
> kPMEReduceChargeGridBuffer_kernel<<<blocks, 256>>>();
>
> and the error arises again:
> BLOCKS: 66150
> NFFT: 360 196 120
> Error: invalid configuration argument launching kernel
> kPMEReduceChargeGridBuffer
>
> Is it there where I have to try with the values 255 and 256 for
> ReduceChargeGridBuffer?
>
> Thanks,
> Pablo.
>
> El 09/10/14 a las #4, Pablo Ródenas escribió:
> > Hi Scott,
> >
> > the nfft values are respectively the followings:
> > 360 196 120
> >
> > I just undo all my changes in ReduceChargeGridBuffer and I will try with:
> > kPMEReduceChargeGridBuffer_kernel<<<blocks, 255>>>();
> > or
> > kPMEReduceChargeGridBuffer_kernel<<<blocks, 256>>>();
> >
> > I will let you know any advance.
> >
> > Thank you very much.
> >
> > Best regards,
> > Pablo.
> >
> >
> > El 09/10/14 a las #4, Scott Le Grand escribió:
> >> Also try 255 and 256 for ReduceChargeGridBuffer...
> >>
> >>
> >> On Thu, Oct 9, 2014 at 8:23 AM, Scott Le Grand <varelse2005.gmail.com>
> >> wrote:
> >>
> >>> What is nfft1, nfft2, and nfft3 for this system?
> >>>
> >>> On Thu, Oct 9, 2014 at 8:20 AM, Pablo Ródenas <pablo.rodenas.bsc.es>
> >>> wrote:
> >>>
> >>>> Dear Scott,
> >>>>
> >>>> you are right, this is a messy change.
> >>>> We only saw clear to change from 65536 to 65535.
> >>>> But after this change on the kPMEFillChargeGridBuffer function, we got
> >>>> the error on the kPMEReduceChargeGridBuffer one. We saw that the value
> >>>> of block variable on the kPMEReduceChargeGridBuffer function is 66150,
> >>>> greater than the one of our card, so we try to solve it with a quick
> >>>> change in order to adapt it to our card maximum values. Please, could
> >>>> you try to fix this function too?
> >>>>
> >>>> FYI, the maximum values of our card are:
> >>>> Maximum number of threads per block: 1024
> >>>> Maximum sizes of each dimension of a block: 1024 x 1024 x 64
> >>>> Maximum sizes of each dimension of a grid: 65535 x 65535 x
> 65535
> >>>>
> >>>> Thank you very much for your answers!
> >>>>
> >>>> Best regards,
> >>>> Pablo.
> >>>>
> >>>>
> >>>> El 09/10/14 a las #4, Scott Le Grand escribió:
> >>>>> Also do *not* *mess* *with* LOADSIZE. You will wake the CUDAThuhlu
> with
> >>>>> such actions and no one wants that...
> >>>>>
> >>>>> The only change you should make here is changing 65536 to 65535.
> >>>>>
> >>>>> That said, thanks for hitting this corner case! Fix shortly...
> >>>>>
> >>>>> On Thu, Oct 9, 2014 at 8:06 AM, Scott Le Grand <
> varelse2005.gmail.com>
> >>>>> wrote:
> >>>>>
> >>>>>> Looking closer at this thread, could you try using 65,535 instead of
> >>>>>> 65,536?
> >>>>>>
> >>>>>> http://en.wikipedia.org/wiki/CUDA, specifically maximum x dimension
> >>>> on SM
> >>>>>> 2.0 GPUs is 65,535...
> >>>>>>
> >>>>>> Missed it by one... Works fine on any Kepler or better class GPU
> >>>> because
> >>>>>> this limit was raised to 2^31 - 1
> >>>>>>
> >>>>>> Ironically, y and z are still limited to 65,535. I'll check in a
> fix
> >>>>>> shortly...
> >>>>>>
> >>>>>>
> >>>>>>
> >>>>>> On Thu, Oct 9, 2014 at 7:58 AM, Scott Le Grand <
> varelse2005.gmail.com>
> >>>>>> wrote:
> >>>>>>
> >>>>>>> Broken.
> >>>>>>>
> >>>>>>> Do not use this.
> >>>>>>>
> >>>>>>> The threadblocks *have* to be 64 for this to work (4 x 4 x 4
> >>>>>>> interpolation)...
> >>>>>>>
> >>>>>>>
> >>>>>>>
> >>>>>>>
> >>>>>>> On Thu, Oct 9, 2014 at 7:22 AM, Pablo Ródenas <
> pablo.rodenas.bsc.es>
> >>>>>>> wrote:
> >>>>>>>
> >>>>>>>> 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
> >>>>>>>>
> >>>>> _______________________________________________
> >>>>> AMBER mailing list
> >>>>> AMBER.ambermd.org
> >>>>> http://lists.ambermd.org/mailman/listinfo/amber
> >>>> --
> >>>> 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
> >>>>
> >> _______________________________________________
> >> AMBER mailing list
> >> AMBER.ambermd.org
> >> http://lists.ambermd.org/mailman/listinfo/amber
>
> --
> 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
>
_______________________________________________
AMBER mailing list
AMBER.ambermd.org
http://lists.ambermd.org/mailman/listinfo/amber
Received on Thu Oct 09 2014 - 09:30:03 PDT
