public void kernel(HThread thread, byte[] ptr )
{
// map from threadIdx/BlockIdx to pixel position
int x = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
int y = thread.threadIdx.y + thread.blockIdx.y * thread.blockDim.y;
int offset = x + y * thread.blockDim.x * thread.gridDim.x;
float ox = (x - DIM/2);
float oy = (y - DIM/2);
float r=0, g=0, b=0;
float maxz = -INF;
for(int i=0; i<SPHERES; i++)
{
float n = 0;
float t = s[i].hit( ox, oy, ref n );
if (t > maxz)
{
float fscale = n;
r = s[i].r * fscale;
g = s[i].g * fscale;
b = s[i].b * fscale;
maxz = t;
}
}
ptr[offset * 4 + 0] = (byte)(r * 255);
ptr[offset * 4 + 1] = (byte)(g * 255);
ptr[offset * 4 + 2] = (byte)(b * 255);
ptr[offset * 4 + 3] = 255;
}
private void SvcArbitrateLock(AThreadState ThreadState)
{
int OwnerThreadHandle = (int)ThreadState.X0;
long MutexAddress = (long)ThreadState.X1;
int RequestingThreadHandle = (int)ThreadState.X2;
HThread RequestingThread = Ns.Os.Handles.GetData <HThread>(RequestingThreadHandle);
Mutex M = new Mutex(Process, MutexAddress, OwnerThreadHandle);
M = Ns.Os.Mutexes.GetOrAdd(MutexAddress, M);
M.WaitForLock(RequestingThread, RequestingThreadHandle);
ThreadState.X0 = (int)SvcResult.Success;
}
private void SendSyncRequest(AThreadState ThreadState, bool UserBuffer)
{
long CmdPtr = ThreadState.Tpidr;
long Size = 0x100;
int Handle = 0;
if (UserBuffer)
{
CmdPtr = (long)ThreadState.X0;
Size = (long)ThreadState.X1;
Handle = (int)ThreadState.X2;
}
else
{
Handle = (int)ThreadState.X0;
}
HThread CurrThread = Process.GetThread(ThreadState.Tpidr);
Process.Scheduler.Suspend(CurrThread.ProcessorId);
byte[] CmdData = AMemoryHelper.ReadBytes(Memory, CmdPtr, (int)Size);
HSession Session = Ns.Os.Handles.GetData <HSession>(Handle);
IpcMessage Cmd = new IpcMessage(CmdData, CmdPtr, Session is HDomain);
if (Session != null)
{
IpcHandler.IpcCall(Ns, Memory, Session, Cmd, CmdPtr, Handle);
byte[] Response = AMemoryHelper.ReadBytes(Memory, CmdPtr, (int)Size);
ThreadState.X0 = (int)SvcResult.Success;
}
else
{
ThreadState.X0 = (int)SvcResult.ErrBadIpcReq;
}
Thread.Yield();
Process.Scheduler.Resume(CurrThread);
}
public void WaitForLock(HThread RequestingThread, int RequestingThreadHandle)
{
lock (EnterWaitLock)
{
int CurrentThreadHandle = ReadMutexValue() & ~MutexHasListenersMask;
if (CurrentThreadHandle == RequestingThreadHandle ||
CurrentThreadHandle == 0)
{
return;
}
WriteMutexValue(CurrentThreadHandle | MutexHasListenersMask);
WaitingThreads.Enqueue(RequestingThread);
}
Process.Scheduler.WaitForSignal(RequestingThread);
}
public void Unlock()
{
lock (EnterWaitLock)
{
int HasListeners = WaitingThreads.Count > 1 ? MutexHasListenersMask : 0;
WriteMutexValue(HasListeners);
HThread[] UnlockedThreads = new HThread[WaitingThreads.Count];
int Index = 0;
while (WaitingThreads.TryDequeue(out HThread Thread))
{
UnlockedThreads[Index++] = Thread;
}
Process.Scheduler.Signal(UnlockedThreads);
}
}
public void WaitForLock(HThread RequestingThread, int RequestingThreadHandle)
{
AcquireMutexValue();
lock (EnterWaitLock)
{
int CurrentThreadHandle = Process.Memory.ReadInt32(MutexAddress) & ~MutexHasListenersMask;
if (CurrentThreadHandle == RequestingThreadHandle ||
CurrentThreadHandle == 0)
{
return;
}
Process.Memory.WriteInt32(MutexAddress, CurrentThreadHandle | MutexHasListenersMask);
ReleaseMutexValue();
WaitingThreads.Enqueue(RequestingThread);
}
Process.Scheduler.WaitForSignal(RequestingThread);
}
public void WaitForSignal(HThread Thread)
{
int Count = Process.Memory.ReadInt32(CondVarAddress);
if (Count <= 0)
{
lock (WaitingThreads)
{
WaitingThreads.Add(Thread);
}
if (Timeout == -1)
{
Process.Scheduler.WaitForSignal(Thread);
}
else
{
Process.Scheduler.WaitForSignal(Thread, (int)(Timeout / 1000000));
lock (WaitingThreads)
{
WaitingThreads.Remove(Thread);
}
}
}
AcquireCondVarValue();
Count = Process.Memory.ReadInt32(CondVarAddress);
if (Count > 0)
{
Process.Memory.WriteInt32(CondVarAddress, Count - 1);
}
ReleaseCondVarValue();
}
public bool Run()
{
if (Executables.Count == 0)
{
return false;
}
long StackBot = TlsPageAddr - MaxStackSize;
Memory.Manager.MapPhys(StackBot, MaxStackSize, (int)MemoryType.Normal, AMemoryPerm.RW);
int Handle = MakeThread(Executables[0].ImageBase, TlsPageAddr, 0, 0, 0);
if (Handle == -1)
{
return false;
}
MainThread = Ns.Os.Handles.GetData<HThread>(Handle);
Scheduler.StartThread(MainThread);
return true;
}
public void Unlock()
{
AcquireMutexValue();
lock (EnterWaitLock)
{
int HasListeners = WaitingThreads.Count > 1 ? MutexHasListenersMask : 0;
Process.Memory.WriteInt32(MutexAddress, HasListeners);
ReleaseMutexValue();
HThread[] UnlockedThreads = new HThread[WaitingThreads.Count];
int Index = 0;
while (WaitingThreads.TryDequeue(out HThread Thread))
{
UnlockedThreads[Index++] = Thread;
}
Process.Scheduler.Signal(UnlockedThreads);
}
}
public static void doVecAdd(HThread thread, int[][] a, int[][] b, int[][] output, int myNumber, int[] array1d)
{
#warning make this so you specify format and size
int[] cache = thread.block.AllocateShared("cache", new int[thread.block.Dim.TotalSize]);
#warning change vars to be props of thread only
int threadId = thread.threadIdx.x + thread.block.Idx.x * thread.block.Dim.x;
for (int i = 0; i < a[threadId].Length; i++)
output[threadId][i] = a[threadId][i] + b[threadId][i];
cache[threadId] = threadId;
thread.block.SyncThreads();
if (threadId == 0)
{
Console.WriteLine("Thread {0} is done! My number is {1}", threadId, myNumber);
int total = 0;
for (int i = 0; i < thread.block.Dim.x; i++)
total += output[i][1];
// for (int i = 0; i < cache.Length; i++)
// Console.WriteLine("{0} {1}", i, cache[i]);
//Console.WriteLine("Thread {0} is done! Total is {1}", threadId, total);
array1d.ToList().ForEach(v => Console.Write(v + ", "));
Console.WriteLine();
}
}
public bool TryGetThread(long Tpidr, out HThread Thread)
{
return(ThreadsByTpidr.TryGetValue(Tpidr, out Thread));
}