AngularAlignmentGpu::AngularSphAlignment Class Reference

#include <cuda_angular_sph_alignment.h>

Public Member Functions
void	associateWith (ProgAngularSphAlignmentGpu *prog)
void	setupConstantParameters ()
void	setupChangingParameters ()
void	pretuneKernel ()
void	runKernel ()
void	runKernelTest (Matrix1D< double > &clnm, size_t idxY0, double RmaxF2, double iRmaxF, Matrix2D< double > R, const MultidimArray< double > &mV, Matrix1D< double > &steps_cp, Matrix1D< int > &vL1, Matrix1D< int > &vN, Matrix1D< int > &vL2, Matrix1D< int > &vM, MultidimArray< int > &V_mask, MultidimArray< double > &mP)
void	transferResults ()
KernelOutputs	getOutputs ()
	AngularSphAlignment ()
	~AngularSphAlignment ()

Detailed Description

Definition at line 67 of file cuda_angular_sph_alignment.h.

Constructor & Destructor Documentation

AngularSphAlignment()

AngularAlignmentGpu::AngularSphAlignment::AngularSphAlignment

(

)

Definition at line 77 of file cuda_angular_sph_alignment.cpp.

{
}

~AngularSphAlignment()

AngularAlignmentGpu::AngularSphAlignment::~AngularSphAlignment

(

)

Definition at line 81 of file cuda_angular_sph_alignment.cpp.

{
    if (program->useFakeKernel) {
    } else {
        cudaFree(dVolData);
        cudaFree(dRotation);
        cudaFree(dZshParams);
        cudaFree(dClnm);
        cudaFree(dVolMask);
        cudaFree(dProjectionPlane);
    }
}

Member Function Documentation

associateWith()

void AngularAlignmentGpu::AngularSphAlignment::associateWith

(

ProgAngularSphAlignmentGpu *

prog

)

Definition at line 94 of file cuda_angular_sph_alignment.cpp.

{
    program = prog;
}

getOutputs()

KernelOutputs AngularAlignmentGpu::AngularSphAlignment::getOutputs

(

)

Definition at line 190 of file cuda_angular_sph_alignment.cpp.

{
    return outputs;
}

pretuneKernel()

void AngularAlignmentGpu::AngularSphAlignment::pretuneKernel

(

)

runKernel()

void AngularAlignmentGpu::AngularSphAlignment::runKernel

(

)

Definition at line 277 of file cuda_angular_sph_alignment.cpp.

{
    if (program->useFakeKernel) {
        fakeKernel(
                Rmax2,
                iRmax,
                imageMetaData,
                dVolData,
                dRotation,
                steps,
                dZshParams,
                dClnm,
                dVolMask,
                dProjectionPlane,
                &outputs);
    } else {
        // Define thrust reduction vector
        thrust::device_vector<PrecisionType> thrustVec(totalGridSize * 3, 0.0);

        // TEST make sure everything is ready before kernel starts
        cudaDeviceSynchronize();

        // Run kernel
        projectionKernel<<<grid, block, constantSharedMemSize + changingSharedMemSize>>>(
                Rmax2,
                iRmax,
                imageMetaData,
                dVolData,
                dRotation,
                steps,
                dZshParams,
                dClnm,
                dVolMask,
                dProjectionPlane,
                thrust::raw_pointer_cast(thrustVec.data())
                );

        cudaDeviceSynchronize();

        auto countIt = thrustVec.begin();
        auto sumVDIt = countIt + totalGridSize;
        auto modgIt = sumVDIt + totalGridSize;

        outputs.count = thrust::reduce(countIt, sumVDIt);
        outputs.sumVD = thrust::reduce(sumVDIt, modgIt);
        outputs.modg = thrust::reduce(modgIt, thrustVec.end());
    }
}

runKernelTest()

void AngularAlignmentGpu::AngularSphAlignment::runKernelTest	(	Matrix1D< double > &	clnm,
		size_t	idxY0,
		double	RmaxF2,
		double	iRmaxF,
		Matrix2D< double >	R,
		const MultidimArray< double > &	mV,
		Matrix1D< double > &	steps_cp,
		Matrix1D< int > &	vL1,
		Matrix1D< int > &	vN,
		Matrix1D< int > &	vL2,
		Matrix1D< int > &	vM,
		MultidimArray< int > &	V_mask,
		MultidimArray< double > &	mP
	)

Definition at line 415 of file cuda_angular_sph_alignment.cpp.

{
    size_t idxZ0=2*idxY0;
    outputs.sumVD = 0.0;
    outputs.modg = 0.0;
    outputs.count = 0.0;

    Matrix1D<double> pos;
    pos.initZeros(3);

    /*
    std::cout 
        << "clnm: " << clnm[0] << "," << clnm[1] << "," << clnm[2] << "\n"
        << "Rmax2: " << RmaxF2 << "\n"
        << "iRmax: " << iRmaxF << "\n"
        << "Rotation: " << R(0, 0) << "," << R(0, 1) << "," << R(1, 2) << "\n"
        << "Volume: " << mV(0, 0, 0) << "," << mV(0, 0, 1) << "," << mV(0, 1, 2) << "\n"
        << "" << std::endl;
    */

    for (int k=STARTINGZ(mV); k<=FINISHINGZ(mV); k++) {
        for (int i=STARTINGY(mV); i<=FINISHINGY(mV); i++) {
            for (int j=STARTINGX(mV); j<=FINISHINGX(mV); j++) {
                ZZ(pos) = k; YY(pos) = i; XX(pos) = j;
                pos = R * pos;
                //if (k == 10 && i == 10 && j == 10)
                //    std::cout << "pos("<<pos[0]<<","<<pos[1]<<","<<pos[2]<<")" << std::endl;
                double gx=0.0, gy=0.0, gz=0.0;
                double k2=ZZ(pos)*ZZ(pos);
                double kr=ZZ(pos)*iRmaxF;
                double k2i2=k2+YY(pos)*YY(pos);
                double ir=YY(pos)*iRmaxF;
                double r2=k2i2+XX(pos)*XX(pos);
                double jr=XX(pos)*iRmaxF;
                double rr=sqrt(r2)*iRmaxF;
                if (r2<RmaxF2) {
                    for (size_t idx=0; idx<idxY0; idx++) {
                        if (VEC_ELEM(steps_cp,idx) == 1) {
                            double zsph=0.0;
                            int l1 = VEC_ELEM(vL1,idx);
                            int n = VEC_ELEM(vN,idx);
                            int l2 = VEC_ELEM(vL2,idx);
                            int m = VEC_ELEM(vM,idx);
                            zsph=::ZernikeSphericalHarmonics(l1,n,l2,m,jr,ir,kr,rr);
                            if (rr>0 || l2==0) {
                                gx += VEC_ELEM(clnm,idx)        *(zsph);
                                gy += VEC_ELEM(clnm,idx+idxY0)  *(zsph);
                                gz += VEC_ELEM(clnm,idx+idxZ0)  *(zsph);
                            }
                        }
                    }

                    int k_mask, i_mask, j_mask;
                    int voxelI_mask;
                    k_mask = (int)(ZZ(pos)+gz);
                    i_mask = (int)(YY(pos)+gy);
                    j_mask = (int)(XX(pos)+gx);

                    if (V_mask.outside(k_mask, i_mask, j_mask)) {
                        voxelI_mask = 0;
                    } else {
                        voxelI_mask = A3D_ELEM(V_mask, k_mask, i_mask, j_mask);
                    }

                    if (voxelI_mask == 1) {
                        double voxelI=mV.interpolatedElement3D(XX(pos)+gx,YY(pos)+gy,ZZ(pos)+gz);
                        A2D_ELEM(mP,i,j) += voxelI;
                        outputs.sumVD += voxelI;
                        outputs.modg += gx*gx+gy*gy+gz*gz;
                        outputs.count++;
                    }
                }
            }
        }
    }
}

setupChangingParameters()

void AngularAlignmentGpu::AngularSphAlignment::setupChangingParameters

(

)

Definition at line 136 of file cuda_angular_sph_alignment.cpp.

{
    if (program == nullptr)
        throw(std::runtime_error("AngularSphAlignment not associated with the program!"));

    if (program->useFakeKernel) {
        setupClnmCpu();
        setupRotationCpu();
        setupProjectionPlaneCpu();
    } else {
        setupClnm();
        setupRotation();
        setupProjectionPlane();
    }

    steps = program->onesInSteps;

    changingSharedMemSize = 0;
    changingSharedMemSize += sizeof(int4) * steps;
    changingSharedMemSize += sizeof(PrecisionType3) * steps;
}

setupConstantParameters()

void AngularAlignmentGpu::AngularSphAlignment::setupConstantParameters

(

)

Definition at line 102 of file cuda_angular_sph_alignment.cpp.

{
    if (program == nullptr)
        throw(std::runtime_error("AngularSphAlignment not associated with the program!"));

    // kernel arguments
    this->Rmax2 = program->RmaxDef * program->RmaxDef;
    this->iRmax = 1.0 / program->RmaxDef;
    setupImageMetaData(program->V);

    if (program->useFakeKernel) {
        setupVolumeDataCpu();
        setupVolumeMaskCpu();
        setupZSHparamsCpu();
    } else {
        setupVolumeData();
        setupVolumeMask();
        setupZSHparams();
    }

    // kernel dimension
    block.x = BLOCK_X_DIM;
    block.y = BLOCK_Y_DIM;
    block.z = BLOCK_Z_DIM;
    grid.x = ((imageMetaData.xDim + block.x - 1) / block.x);
    grid.y = ((imageMetaData.yDim + block.y - 1) / block.y);
    grid.z = ((imageMetaData.zDim + block.z - 1) / block.z);

    totalGridSize = grid.x * grid.y * grid.z;

    // Dynamic shared memory
    constantSharedMemSize = 0;
}

transferResults()

void AngularAlignmentGpu::AngularSphAlignment::transferResults

(

)

Definition at line 343 of file cuda_angular_sph_alignment.cpp.

{
    if (program->useFakeKernel) {
        transferProjectionPlaneCpu();
    } else {
        transferProjectionPlane();
    }
}

---

The documentation for this class was generated from the following files:

• xmipp/libraries/reconstruction_cuda/cuda_angular_sph_alignment.h
• xmipp/libraries/reconstruction_cuda/cuda_angular_sph_alignment.cpp