Cuda Solver

[Bird]

see if it get to line UpdateTransformNotify(threadIndex, body);

Yes, I am getting there when I set a break point.

[Bird]

here what it looks like here.

https://youtu.be/nHuFFKJITrw

I'm shutting down for the night though. Be back tomorrow.

[Julio Jerez]

Ok. So it is getting the call.
Later I will add a trace to print the rotation and angular velocity.
Them will will see if those value are as we expect.

But that can be tomorrow.

[Bird]

If I let it run for a while in debug mode but without debugging, I hit this assert

[Julio Jerez]

I think the line is this, but since the file changed, try sync again and run in debug.

Code: Select all

   cudaError_t cudaStatus = cudaMemcpyAsync(cpuInfo, gpuInfo, sizeof(cuSceneInfo), cudaMemcpyDeviceToHost, m_context->m_stream0);
   dAssert(cudaStatus == cudaSuccess);
   if (cudaStatus != cudaSuccess)
   {
      dAssert(0);
   }


if it can't get the memory buffer, that a really bad sign.
so far it usually means is over running memory.

I comment it with most broad phase code commneted out on that file.
so it will only do the transform, let us see if that works.

that the part by the end on teh same file line 749

Code: Select all

   //CudaInitBodyArray << <blocksCount, D_THREADS_PER_BLOCK, 0, stream >> > (CalcuateBodyAabb, *infoGpu);
   //CudaMergeAabb << <1, D_THREADS_PER_BLOCK, 0, stream >> > (ReducedAabb, *infoGpu);
   //CudaCountAabb << <blocksCount, D_THREADS_PER_BLOCK, 0, stream >> > (CountAabb, *infoGpu);
   //CudaPrefixScanSum0 << <blocksCount, D_THREADS_PER_BLOCK, 0, stream >> > (PrefixScanSum0, *infoGpu);
   //CudaPrefixScanSum1 << <1, D_THREADS_PER_BLOCK, 0, stream >> > (PrefixScanSum1, *infoGpu);
   //CudaGenerateGridHash << <blocksCount, D_THREADS_PER_BLOCK, 0, stream >> > (GenerateHashGrids, *infoGpu);
   //CudaEndGridHash << <1, 1, 0, stream >> > (EndGridHash, *infoGpu);
   //CudaSortGridHash (m_context);
}


if it still fail them we know is in the transform kernels.
if it passes, then uncomment those lines one by one from top to bottom, until we get the one cause the problem.

[Bird]

I'm not hitting the assert anymore in this latest version. But the transform is still not working.

My OptiX renderer uses a cuda context too. so just as a sanity check, I disabled OptiX completely and I'm still not seeing the transform in cuNewton

Just to be clear, the Newton Sandbox demo works perfectly. Its only when I try to hook up my app to the Newton SDK that I'm having a problem. I've compiled Newton as a static library.

[Julio Jerez]

ok we can rule out static or dynamics linking. but if the assert goes away and the transform still do no how indicate more than one bug.

let us find the rotation bug first. I add that so that I can get some feedback since it will be some time until we make enough progress to see actual physics. so that was a good call.

let us add some traces form bottom up to see when is going wrong.
I committed the and if you runn it should print something like this.

omega(0.654416 -0.409420 -1.027692) rot(0.467492 0.718702 0.193209 0.477063)
omega(0.654416 -0.409420 -1.027692) rot(0.476125 0.712812 0.191755 0.477947)
omega(0.654416 -0.409420 -1.027692) rot(0.484717 0.706879 0.190225 0.478733)
omega(0.654416 -0.409420 -1.027692) rot(0.476125 0.712812 0.191755 0.477947)
omega(0.654416 -0.409420 -1.027692) rot(0.501776 0.694884 0.186942 0.480010)
omega(0.654416 -0.409420 -1.027692) rot(0.510240 0.688825 0.185190 0.480500)

see if Omega is not zero and rot should be changing, if rot is not changing, them we will move the trace to the shaders only we find where is going wrong.

keep only one body so that the print is not so intrusive

[Bird]

I'm not seeing any debug output when I run the sandbox demo or my own app. I see you're using dTrace and I don't think that works in cuda code, does it?

[Bird]

okay I got it working by changing dTrace to printf and I'm seeing no rotation in the output

Code: Select all

printf("omega(%f %f %f) rot(%f %f %f %f)\n",
            body->m_omega[0], body->m_omega[1], body->m_omega[2],
            rotation[0], rotation[1], rotation[2], rotation[3]);

omega(0.000000 0.000000 0.000000) rot(0.000000 0.000000 0.000000 1.000000)

[Julio Jerez]

Ok, at least is so getting garbage.

I will add few prints to some place to see where the value is going wrong.
Stand by.

[Julio Jerez]

ok I now added printf to the shader that calculate omega and rotation.
the print should look like this

IntegrateVelocity: id(1) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateExternalForce: id(0) w(0.736507 0.335142 -0.895999) r(0.978725 0.096457 -0.110242 0.143666)
IntegrateExternalForce: id(1) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.736507 0.335142 -0.895999) r(0.979360 0.093341 -0.111848 0.140115)
IntegrateVelocity: id(1) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
GetTransform: id(0) w(0.736507 0.335142 -0.895999) r(0.979360 0.093341 -0.111848 0.140115)
IntegrateExternalForce: id(0) w(0.738657 0.334356 -0.894186) r(0.979360 0.093341 -0.111848 0.140115)
IntegrateExternalForce: id(1) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.738657 0.334356 -0.894186) r(0.979971 0.090230 -0.113445 0.136552)
IntegrateVelocity: id(1) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)

I see a problem because is shoe two bodies when is should just show one.
the last body is always the sentinel.
so I wonder if it is getting the sentinel.

I will fix that, but you try it to see if you get one of the two bodies changing rotations and non zero omega.

[Julio Jerez]

ok if sync before, please do it again, I fixed that is was operating on the sentinel.
the print out should looks like this.

GetTransform: id(0) w(0.482468 0.196393 -1.128870) r(0.763875 0.466181 0.091196 0.436868)
IntegrateExternalForce: id(0) w(0.484116 0.199775 -1.127159) r(0.763875 0.466181 0.091196 0.436868)
IntegrateVelocity: id(0) w(0.484116 0.199775 -1.127159) r(0.767011 0.462767 0.089447 0.435361)
IntegrateExternalForce: id(0) w(0.485790 0.203121 -1.125429) r(0.767011 0.462767 0.089447 0.435361)
IntegrateVelocity: id(0) w(0.485790 0.203121 -1.125429) r(0.770128 0.459352 0.087692 0.433831)
GetTransform: id(0) w(0.485790 0.203121 -1.125429) r(0.770128 0.459352 0.087692 0.433831)
IntegrateExternalForce: id(0) w(0.487488 0.206433 -1.123679) r(0.770128 0.459352 0.087692 0.433831)
IntegrateVelocity: id(0) w(0.487488 0.206433 -1.123679) r(0.773225 0.455935 0.085930 0.432276)
IntegrateExternalForce: id(0) w(0.489212 0.209709 -1.121911) r(0.773225 0.455935 0.085930 0.432276)
IntegrateVelocity: id(0) w(0.489212 0.209709 -1.121911) r(0.776302 0.452517 0.084162 0.430698)
GetTransform: id(0) w(0.489212 0.209709 -1.121911) r(0.776302 0.452517 0.084162 0.430698)
IntegrateExternalForce: id(0) w(0.490959 0.212949 -1.120124) r(0.776302 0.452517 0.084162 0.430698)

[Bird]

Here's my output

25.036558 DEBUG [9144 NewtonEngine.cpp->tick:36] TICK 93
IntegrateExternalForce: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
GetTransform: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateExternalForce: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateExternalForce: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
GetTransform: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateExternalForce: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateExternalForce: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)
IntegrateVelocity: id(0) w(0.000000 0.000000 0.000000) r(0.000000 0.000000 0.000000 1.000000)

[Julio Jerez]

ah ok, so it seems it is failling at initialize the buffer in the first time.
that happens on this function file: sdk\dExtensions\dCuda\cuNatives\cuDeviceBuffer.h

Code: Select all

template<class T>
void cuDeviceBuffer<T>::ReadData(const T* const src, ndInt32 elements)
{
   dAssert(elements <= m_size);
   cudaMemcpy(m_array, src, sizeof (T) * elements, cudaMemcpyHostToDevice);
}


can you edit to this:

Code: Select all

template<class T>
void cuDeviceBuffer<T>::ReadData(const T* const src, ndInt32 elements)
{
   dAssert(elements <= m_size);
   cudaMemcpy(m_array, src, sizeof (T) * elements, cudaMemcpyHostToDevice);

   printf("ReadData: id(%d) w(%f %f %f) r(%f %f %f %f)\n", 0,
      src[0].m_omega.x, src[0].m_omega.y, src[0].m_omega.z,
      src[0].m_rotation.x, src[0].m_rotation.y, src[0].m_rotation.z, src[0].m_rotation.w);
}


then se a break point to check that the values are no zero.
that whet the proxies are initialized. If the value is zero, them is most be that for some reason a rotation is failing.
play try doing that, and we will know where to check next.

it also seems that the rotation is the problem because unless the position is also zero, it seems the body is placed in the correct location.

[Julio Jerez]

I added the printf, just sync and place a break point to see if the value are correct comming in.