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
Received on Thu Oct 09 2014 - 09:00:02 PDT
