GPUOcelot
Public Types | Public Member Functions | Public Attributes | Protected Member Functions

executive::CooperativeThreadArray Class Reference

#include <CooperativeThreadArray.h>

Collaboration diagram for executive::CooperativeThreadArray:
Collaboration graph
[legend]

List of all members.

Public Types

typedef std::vector< CTABarrierBarrierVector
typedef std::vector< ir::PTXU64RegisterFile

Public Member Functions

 CooperativeThreadArray (EmulatedKernel *kernel, const ir::Dim3 &gridDim, bool trace)
 CooperativeThreadArray ()
 ~CooperativeThreadArray ()
void reset ()
void execute (const ir::Dim3 &block)
void jumpToPC (int PC)
RegisterFile getCurrentRegisterFile () const
CTAContextgetActiveContext ()
ir::PTXU8 getRegAsU8 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXU16 getRegAsU16 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXU32 getRegAsU32 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXU64 getRegAsU64 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXS8 getRegAsS8 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXS16 getRegAsS16 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXS32 getRegAsS32 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXS64 getRegAsS64 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXF32 getRegAsF32 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXF64 getRegAsF64 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXB8 getRegAsB8 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXB16 getRegAsB16 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXB32 getRegAsB32 (int threadID, ir::PTXOperand::RegisterType reg)
ir::PTXB64 getRegAsB64 (int threadID, ir::PTXOperand::RegisterType reg)
bool getRegAsPredicate (int threadID, ir::PTXOperand::RegisterType reg)
void setRegAsU8 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXU8 value)
void setRegAsU16 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXU16 value)
void setRegAsU32 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXU32 value)
void setRegAsU64 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXU64 value)
void setRegAsS8 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXS8 value)
void setRegAsS16 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXS16 value)
void setRegAsS32 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXS32 value)
void setRegAsS64 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXS64 value)
void setRegAsF32 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXF32 value)
void setRegAsF64 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXF64 value)
void setRegAsB8 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXB8 value)
void setRegAsB16 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXB16 value)
void setRegAsB32 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXB32 value)
void setRegAsB64 (int threadID, ir::PTXOperand::RegisterType reg, ir::PTXB64 value)
void setRegAsPredicate (int threadID, ir::PTXOperand::RegisterType reg, bool value)
ir::PTXU8 operandAsU8 (int, const ir::PTXOperand &)
ir::PTXU16 operandAsU16 (int, const ir::PTXOperand &)
ir::PTXU32 operandAsU32 (int, const ir::PTXOperand &)
ir::PTXU64 operandAsU64 (int, const ir::PTXOperand &)
ir::PTXS8 operandAsS8 (int, const ir::PTXOperand &)
ir::PTXS16 operandAsS16 (int, const ir::PTXOperand &)
ir::PTXS32 operandAsS32 (int, const ir::PTXOperand &)
ir::PTXS64 operandAsS64 (int, const ir::PTXOperand &)
ir::PTXF32 operandAsF32 (int, const ir::PTXOperand &)
ir::PTXF64 operandAsF64 (int, const ir::PTXOperand &)
ir::PTXB8 operandAsB8 (int, const ir::PTXOperand &)
ir::PTXB16 operandAsB16 (int, const ir::PTXOperand &)
ir::PTXB32 operandAsB32 (int, const ir::PTXOperand &)
ir::PTXB64 operandAsB64 (int, const ir::PTXOperand &)
bool operandAsPredicate (int, const ir::PTXOperand &)
void eval_Abs (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Add (CTAContext &context, const ir::PTXInstruction &instr)
void eval_AddC (CTAContext &context, const ir::PTXInstruction &instr)
void eval_And (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Atom (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Bar (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Bfi (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Bfind (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Bfe (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Bra (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Brev (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Brkpt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Call (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Clz (CTAContext &context, const ir::PTXInstruction &instr)
void eval_CNot (CTAContext &context, const ir::PTXInstruction &instr)
void eval_CopySign (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Cos (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Cvt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Cvta (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Div (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Ex2 (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Exit (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Fma (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Isspacep (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Ld (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Ldu (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Lg2 (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mad24 (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mad (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Max (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Membar (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Min (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mul24 (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mul (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Neg (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Not (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Or (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Pmevent (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Popc (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Prmt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Rcp (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Red (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Rem (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Ret (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Rsqrt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Sad (CTAContext &context, const ir::PTXInstruction &instr)
void eval_SelP (CTAContext &context, const ir::PTXInstruction &instr)
void eval_SetP (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Set (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Shl (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Shr (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Sin (CTAContext &context, const ir::PTXInstruction &instr)
void eval_SlCt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Sqrt (CTAContext &context, const ir::PTXInstruction &instr)
void eval_St (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Sub (CTAContext &context, const ir::PTXInstruction &instr)
void eval_SubC (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Suld (CTAContext &context, const ir::PTXInstruction &instr)
 load from surface memory
void eval_Sured (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Sust (CTAContext &context, const ir::PTXInstruction &instr)
 store to surface memory
void eval_Suq (CTAContext &context, const ir::PTXInstruction &instr)
void eval_TestP (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Tex (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Trap (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Txq (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Vote (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Xor (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Reconverge (CTAContext &context, const ir::PTXInstruction &instr)

Public Attributes

ir::Dim3 blockDim
ir::Dim3 gridDim
int threadCount
EmulatedKernelkernel
ir::Dim3 blockId
EmulatorCallStack functionCallStack
BarrierVector barriers
ReconvergenceMechanismreconvergenceMechanism
 abstraction for reconvergence mechanism
ir::PTXU64 clock
bool traceEvents
int counter
trace::TraceEvent currentEvent

Protected Member Functions

void initialize (const ir::Dim3 &block)
void finalize ()
const ir::PTXInstructioncurrentInstruction (CTAContext &context)
ir::PTXU32 getSpecialValue (const int threadId, const ir::PTXOperand::SpecialRegister, const ir::PTXOperand::VectorIndex) const
ir::PTXF32 sat (int modifier, ir::PTXF32 f)
void trace ()
void postTrace ()
 invokes TraceGenerator::postEvent() on all trace generators
void eval_Mov_reg (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov_sreg (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov_imm (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov_indirect (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov_addr (CTAContext &context, const ir::PTXInstruction &instr)
void eval_Mov_func (CTAContext &context, const ir::PTXInstruction &instr)
void copyArgument (const ir::PTXOperand &s, CTAContext &context)

Detailed Description

Defines state of cooperative thread array


Member Typedef Documentation


Constructor & Destructor Documentation

executive::CooperativeThreadArray::CooperativeThreadArray ( EmulatedKernel k,
const ir::Dim3 grid,
bool  trace 
)

Constructs a cooperative thread array from an EmulatedKernel instance

Parameters:
kernelpointer to EmulatedKernel to which this CTA belongs
gridDimThe dimensions of the kernel
traceEnable trace generation

Constructs a cooperative thread array from an EmulatedKernel instance

Parameters:
kernelpointer to EmulatedKernel to which this CTA belongs
executive::CooperativeThreadArray::CooperativeThreadArray ( )
executive::CooperativeThreadArray::~CooperativeThreadArray ( )

Destroys state associated with CTA


Member Function Documentation

void executive::CooperativeThreadArray::copyArgument ( const ir::PTXOperand s,
CTAContext context 
) [protected]
const ir::PTXInstruction & executive::CooperativeThreadArray::currentInstruction ( CTAContext context) [protected]

Gets current instruction

void executive::CooperativeThreadArray::eval_Abs ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Add ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_AddC ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_And ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Atom ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Bar ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Bfe ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Bfi ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Bfind ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Bra ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Brev ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Brkpt ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Call ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Clz ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_CNot ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_CopySign ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Cos ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Cvt ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Cvta ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Div ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Ex2 ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Exit ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Fma ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Isspacep ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Ld ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Ldu ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Lg2 ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Mad ( CTAContext context,
const ir::PTXInstruction instr 
)

Mad:

mad[.hi,.lo,.wide][.sat].itype d, a, b, c; mad[.rnd][.sat].ftype d, a, b, c;

t = a * b; n = bitwidth of type; d = t + c; // for floating-point and .wide d = t<2n-1..n> + c; // for .hi variant d = t<n-1..0> + c; // for .lo variant

void executive::CooperativeThreadArray::eval_Mad24 ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Max ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Membar ( CTAContext context,
const ir::PTXInstruction instr 
)

No need to do anything here.

void executive::CooperativeThreadArray::eval_Min ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Mov ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Mov_addr ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mov_func ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mov_imm ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mov_indirect ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mov_reg ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mov_sreg ( CTAContext context,
const ir::PTXInstruction instr 
) [protected]
void executive::CooperativeThreadArray::eval_Mul ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Mul24 ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Neg ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Not ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Or ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Pmevent ( CTAContext context,
const ir::PTXInstruction instr 
)

No need to do anything here.

void executive::CooperativeThreadArray::eval_Popc ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Prmt ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Rcp ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Reconverge ( CTAContext context,
const ir::PTXInstruction instr 
)

Reconverge instruction is inserted into the PTX during analysis and construction of the emulated kernel.

void executive::CooperativeThreadArray::eval_Red ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Rem ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Ret ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Rsqrt ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Sad ( CTAContext context,
const ir::PTXInstruction instr 
)

sad.type d, a, b, c;

d = c + ((a<b) ? b-a : a-b);

.type = { .u16, .u32, .u64, .s16, .s32, .s64 };

void executive::CooperativeThreadArray::eval_SelP ( CTAContext context,
const ir::PTXInstruction instr 
)

Select between source operands according to the predicate value selp.type d, a, b, c;

.type = { .b16, .b32, .b64, .u16, .u32, .u64, .s16, .s32, .s64, .f32, .f64 };

void executive::CooperativeThreadArray::eval_Set ( CTAContext context,
const ir::PTXInstruction instr 
)

PTX set instruction

t = (a CmpOp b) ? 1 : 0; if (isFloat(dtype)) d = BoolOp(t, c) ? 1.0f : 0x00000000; else d = BoolOp(t, c) ? 0xFFFFFFFF : 0x00000000;

void executive::CooperativeThreadArray::eval_SetP ( CTAContext context,
const ir::PTXInstruction instr 
)

Compare two numeric values with a relational operator and [optionally] combine this result with a predicate value by applying a boolean operator

setp.CmpOp.type p[|q], a, b; setp.CmpOp.BoolOp.type p[|q], a, b, [!]c;

.type = { .b16, .b32, .b64, .u16, .u32, .u64, .s16, .s32, .s64, .f32, .f64 };

void executive::CooperativeThreadArray::eval_Shl ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Shr ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Sin ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_SlCt ( CTAContext context,
const ir::PTXInstruction instr 
)

slct.dtype.ctype d, a, b, c;

.dtype = { .b16, .b32, .b64, .u16, .u32, .u64, .s16, .s32, .s64, .f32, .f64 };

.ctype = { .s32, .f32 };

d = (c >= 0) ? a : b;

void executive::CooperativeThreadArray::eval_Sqrt ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_St ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Sub ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_SubC ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Suld ( CTAContext context,
const ir::PTXInstruction instr 
)

load from surface memory

void executive::CooperativeThreadArray::eval_Suq ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Sured ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Sust ( CTAContext context,
const ir::PTXInstruction instr 
)

store to surface memory

void executive::CooperativeThreadArray::eval_TestP ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Tex ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Trap ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Txq ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Vote ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::eval_Xor ( CTAContext context,
const ir::PTXInstruction instr 
)
void executive::CooperativeThreadArray::execute ( const ir::Dim3 block)

Initializes the CTA and executes the kernel for a given block

Called by the worker thread to evaluate a block

void executive::CooperativeThreadArray::finalize ( ) [protected]

finishes execution of the CTA

executive::CTAContext & executive::CooperativeThreadArray::getActiveContext ( )

gets the active context of the cooperative thread array

executive::CooperativeThreadArray::RegisterFile executive::CooperativeThreadArray::getCurrentRegisterFile ( ) const
ir::PTXB16 executive::CooperativeThreadArray::getRegAsB16 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXB32 executive::CooperativeThreadArray::getRegAsB32 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXB64 executive::CooperativeThreadArray::getRegAsB64 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXB8 executive::CooperativeThreadArray::getRegAsB8 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXF32 executive::CooperativeThreadArray::getRegAsF32 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXF64 executive::CooperativeThreadArray::getRegAsF64 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
bool executive::CooperativeThreadArray::getRegAsPredicate ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread
regindex of register
ir::PTXS16 executive::CooperativeThreadArray::getRegAsS16 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXS32 executive::CooperativeThreadArray::getRegAsS32 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXS64 executive::CooperativeThreadArray::getRegAsS64 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXS8 executive::CooperativeThreadArray::getRegAsS8 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXU16 executive::CooperativeThreadArray::getRegAsU16 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXU32 executive::CooperativeThreadArray::getRegAsU32 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXU64 executive::CooperativeThreadArray::getRegAsU64 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXU8 executive::CooperativeThreadArray::getRegAsU8 ( int  threadID,
ir::PTXOperand::RegisterType  reg 
)

Gets a register value

Parameters:
threadIDID of the active thread register index
ir::PTXU32 executive::CooperativeThreadArray::getSpecialValue ( const int  threadId,
const ir::PTXOperand::SpecialRegister  reg,
const ir::PTXOperand::VectorIndex  index 
) const [protected]

Gets special value

void executive::CooperativeThreadArray::initialize ( const ir::Dim3 block) [protected]

initializes elements of the CTA

void executive::CooperativeThreadArray::jumpToPC ( int  PC)

Jump to a specific PC for the current context

ir::PTXB16 executive::CooperativeThreadArray::operandAsB16 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXB32 executive::CooperativeThreadArray::operandAsB32 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXB64 executive::CooperativeThreadArray::operandAsB64 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXB8 executive::CooperativeThreadArray::operandAsB8 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXF32 executive::CooperativeThreadArray::operandAsF32 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXF64 executive::CooperativeThreadArray::operandAsF64 ( int  threadID,
const ir::PTXOperand op 
)
bool executive::CooperativeThreadArray::operandAsPredicate ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXS16 executive::CooperativeThreadArray::operandAsS16 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXS32 executive::CooperativeThreadArray::operandAsS32 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXS64 executive::CooperativeThreadArray::operandAsS64 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXS8 executive::CooperativeThreadArray::operandAsS8 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXU16 executive::CooperativeThreadArray::operandAsU16 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXU32 executive::CooperativeThreadArray::operandAsU32 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXU64 executive::CooperativeThreadArray::operandAsU64 ( int  threadID,
const ir::PTXOperand op 
)
ir::PTXU8 executive::CooperativeThreadArray::operandAsU8 ( int  threadID,
const ir::PTXOperand op 
)
void executive::CooperativeThreadArray::postTrace ( ) [protected]

invokes TraceGenerator::postEvent() on all trace generators

void executive::CooperativeThreadArray::reset ( )
ir::PTXF32 executive::CooperativeThreadArray::sat ( int  modifier,
ir::PTXF32  f 
) [protected]
void executive::CooperativeThreadArray::setRegAsB16 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXB16  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsB32 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXB32  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsB64 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXB64  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsB8 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXB8  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsF32 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXF32  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsF64 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXF64  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsPredicate ( int  threadID,
ir::PTXOperand::RegisterType  reg,
bool  value 
)

Sets a register value

Parameters:
threadIDID of the active thread
regindex of register
valuevalue of predicate register
void executive::CooperativeThreadArray::setRegAsS16 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXS16  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsS32 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXS32  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsS64 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXS64  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsS8 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXS8  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsU16 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXU16  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsU32 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXU32  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsU64 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXU64  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::setRegAsU8 ( int  threadID,
ir::PTXOperand::RegisterType  reg,
ir::PTXU8  value 
)

Sets a register value

Parameters:
threadIDID of the active thread register index
void executive::CooperativeThreadArray::trace ( ) [protected]

Member Data Documentation

Vector of named barriers

Dimensions of the cooperative thread array

ID of block implemented by this CooperativeThreadArray instance

Counter incremented 4 times per instruction

Number of dynamic instructions executed

An object used to trace execution of the CooperativeThreadArray

Function call stack

Dimensions of the kernel

Pointer to EmulatedKernel instance that this CTA is executing

abstraction for reconvergence mechanism

Number of threads in CTA (equal to blockDim.x * blockDim.y * blockDim.z)

Flag to enable or disable tracing of events


The documentation for this class was generated from the following files:
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Defines