8 #include "IdentDelivery.hpp" 9 #include "columns/HyPerCol.hpp" 10 #include "utils/MapLookupByType.hpp" 15 IdentDelivery::IdentDelivery(
char const *name, HyPerCol *hc) { initialize(name, hc); }
17 int IdentDelivery::initialize(
char const *name, HyPerCol *hc) {
18 return BaseDelivery::initialize(name, hc);
21 void IdentDelivery::setObjectType() { mObjectType =
"IdentDelivery"; }
26 if (ioFlag == PARAMS_IO_READ) {
27 parent->parameters()->handleUnnecessaryParameter(name,
"receiveGpu",
false );
32 IdentDelivery::communicateInitInfo(std::shared_ptr<CommunicateInitInfoMessage const> message) {
33 auto status = BaseDelivery::communicateInitInfo(message);
34 if (status != Response::SUCCESS) {
38 mSingleArbor = mapLookupByType<SingleArbor>(message->mHierarchy, getDescription());
39 FatalIf(!mSingleArbor,
"%s requires a SingleArbor component.\n", getDescription_c());
41 checkPreAndPostDimensions();
42 return Response::SUCCESS;
45 void IdentDelivery::checkPreAndPostDimensions() {
46 int status = PV_SUCCESS;
47 pvAssert(mPreLayer and mPostLayer);
48 PVLayerLoc const *preLoc = mPreLayer->getLayerLoc();
49 PVLayerLoc const *postLoc = mPostLayer->getLayerLoc();
50 if (preLoc->nx != postLoc->nx) {
52 "%s requires pre and post nx be equal (%d versus %d).\n",
58 if (preLoc->ny != postLoc->ny) {
60 "%s requires pre and post ny be equal (%d versus %d).\n",
66 if (preLoc->nf != postLoc->nf) {
68 "%s requires pre and post nf be equal (%d versus %d).\n",
74 if (preLoc->nbatch != postLoc->nbatch) {
76 "%s requires pre and post nbatch be equal (%d versus %d).\n",
84 "IdentDelivery \"%s\" Error: %s and %s do not have the same dimensions.\n Dims: " 85 "%dx%dx%d vs. %dx%dx%d\n",
88 mPostLayer->getName(),
97 void IdentDelivery::deliver() {
98 if (mChannelCode == CHANNEL_NOUPDATE) {
102 int delay = mSingleArbor->getDelay(0);
104 PVLayerLoc const &preLoc = preActivityCube.loc;
105 PVLayerLoc const &postLoc = *mPostLayer->getLayerLoc();
107 int const nx = preLoc.nx;
108 int const ny = preLoc.ny;
109 int const nf = preLoc.nf;
110 int nxPreExtended = nx + preLoc.halo.lt + preLoc.halo.rt;
111 int nyPreExtended = ny + preLoc.halo.dn + preLoc.halo.up;
112 int numPreExtended = nxPreExtended * nyPreExtended * nf;
113 pvAssert(numPreExtended * preLoc.nbatch == preActivityCube.numItems);
114 int numPostRestricted = nx * ny * nf;
116 float *postChannel = mPostLayer->getChannel(mChannelCode);
117 for (
int b = 0; b < parent->getNBatch(); b++) {
118 float const *preActivityBuffer = preActivityCube.data + b * numPreExtended;
119 float *postGSynBuffer = postChannel + b * numPostRestricted;
120 if (preActivityCube.isSparse) {
123 int numActive = preActivityCube.numActive[b];
124 #ifdef PV_USE_OPENMP_THREADS
125 #pragma omp parallel
for 127 for (
int loopIndex = 0; loopIndex < numActive; loopIndex++) {
128 int kPre = activeIndices[loopIndex].index;
129 int kx = kxPos(kPre, nxPreExtended, nyPreExtended, nf) - preLoc.halo.lt;
130 int ky = kyPos(kPre, nxPreExtended, nyPreExtended, nf) - preLoc.halo.up;
131 if (kx < 0 or kx >= nx or ky < 0 or ky >= ny) {
134 int kf = featureIndex(kPre, nxPreExtended, nyPreExtended, nf);
135 int kPost = kIndex(kx, ky, kf, nx, ny, nf);
136 pvAssert(kPost >= 0 and kPost < numPostRestricted);
137 float a = activeIndices[loopIndex].value;
138 postGSynBuffer[kPost] += a;
142 int const nk = postLoc.nx * postLoc.nf;
143 #ifdef PV_USE_OPENMP_THREADS 144 #pragma omp parallel for 146 for (
int y = 0; y < ny; y++) {
148 kIndex(preLoc.halo.lt, y + preLoc.halo.up, 0, nxPreExtended, nyPreExtended, nf);
150 float const *preActivityLine = &preActivityBuffer[preLineIndex];
151 int postLineIndex = kIndex(0, y, 0, postLoc.nx, ny, postLoc.nf);
152 float *postGSynLine = &postGSynBuffer[postLineIndex];
153 for (
int k = 0; k < nk; k++) {
154 postGSynLine[k] += preActivityLine[k];
160 mPostLayer->setUpdatedDeviceGSynFlag(!mReceiveGpu);
161 #endif // PV_USE_CUDA 164 void IdentDelivery::deliverUnitInput(
float *recvBuffer) {
165 const int numNeuronsPost = mPostLayer->getNumNeuronsAllBatches();
166 #ifdef PV_USE_OPENMP_THREADS 167 #pragma omp parallel for 169 for (
int k = 0; k < numNeuronsPost; k++) {
170 recvBuffer[k] += 1.0f;
176 if (getChannelCode() == CHANNEL_NOUPDATE) {
virtual void ioParam_receiveGpu(enum ParamsIOFlag ioFlag) override
IdentDeliver does not use the GPU. It is an error to set receiveGpu to true.
PVLayerCube createCube(int delay=0)
bool isExchangeFinished(int delay=0)
virtual bool isAllInputReady() override