ManyBodySpaceBase.cuh File Reference

Classes
class	ObjectOnGPU< ManyBodySpaceBase< Derived > >

Functions
template<class Derived >
void	copyTransEqClass (ObjectOnGPU< ManyBodySpaceBase< Derived > > &obj, ManyBodySpaceBase< Derived > const &src)
template<typename Derived >
__global__ void	setTransEqDim_kernel (ManyBodySpaceBase< Derived > *obj, int transEqDim)
template<typename Derived >
__global__ void	copyTransEqClass_kernel (ManyBodySpaceBase< Derived > *obj, int *srcTransEqClassRep, int *srcTransPeriod)

Function Documentation

copyTransEqClass()

template<class Derived >

void copyTransEqClass	(	ObjectOnGPU< ManyBodySpaceBase< Derived > > &	obj,
		ManyBodySpaceBase< Derived > const &	src
	)

                                                                      {
    debug_print(__PRETTY_FUNCTION__ << "\n\tsrc.transEqDim() = " << src.transEqDim() << "\n");
    if(src.transEqDim() == -1) return;
 
    int* srcTransEqClassRep = nullptr;
    int* srcTransPeriod     = nullptr;
    cuCHECK(cudaMalloc(&srcTransEqClassRep, src.transEqDim() * sizeof(int)));
    cuCHECK(cudaMalloc(&srcTransPeriod, src.transEqDim() * sizeof(int)));
    cuCHECK(cudaMemcpy(srcTransEqClassRep, src.transEqClassRep().data(),
                       src.transEqDim() * sizeof(int), cudaMemcpyHostToDevice));
    cuCHECK(cudaMemcpy(srcTransPeriod, src.transPeriod().data(), src.transEqDim() * sizeof(int),
                       cudaMemcpyHostToDevice));
    setTransEqDim_kernel<<<1, 1>>>(obj.ptr(), src.transEqDim());
 
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, 0);  // 0-th device
    constexpr int warpSize    = 32;
    int const     nWarps      = (src.transEqDim() % warpSize == 0 ? src.transEqDim() / warpSize
                                                                  : src.transEqDim() / warpSize + 1);
    int const     nWarpsPerSM = (nWarps % deviceProp.multiProcessorCount == 0
                                     ? nWarps / deviceProp.multiProcessorCount
                                     : nWarps / deviceProp.multiProcessorCount + 1);
    int const     nThreads    = min(nWarpsPerSM * warpSize, deviceProp.maxThreadsPerBlock);
    int const     nBlocks     = (src.transEqDim() % nThreads == 0 ? src.transEqDim() / nThreads
                                                                  : src.transEqDim() / nThreads + 1);
    // std::cout << "\tdeviceProp.multiProcessorCount = " << deviceProp.multiProcessorCount
    //           << ", deviceProp.sharedMemPerBlock = " << deviceProp.sharedMemPerBlock << std::endl;
    // std::cout << "\tnBlocks = " << nBlocks << ", nThreads = " << nThreads << std::endl;
    copyTransEqClass_kernel<<<nBlocks, nThreads>>>(obj.ptr(), srcTransEqClassRep, srcTransPeriod);
    cuCHECK(cudaGetLastError());
    cuCHECK(cudaFree(srcTransEqClassRep));
    cuCHECK(cudaFree(srcTransPeriod));
}

copyTransEqClass_kernel()

template<typename Derived >

__global__ void copyTransEqClass_kernel	(	ManyBodySpaceBase< Derived > *	obj,
		int *	srcTransEqClassRep,
		int *	srcTransPeriod
	)

                                                             {
    int const idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx >= obj->transEqDim()) return;
    obj->m_transEqClassRep(idx) = srcTransEqClassRep[idx];
    obj->m_transPeriod(idx)     = srcTransPeriod[idx];
    for(auto trans = 0; trans != srcTransPeriod[idx]; ++trans) {
        auto translated                        = obj->translate(srcTransEqClassRep[idx], trans);
        obj->m_stateToTransEqClass(translated) = idx;
    }
}

setTransEqDim_kernel()

template<typename Derived >

__global__ void setTransEqDim_kernel	(	ManyBodySpaceBase< Derived > *	obj,
		int	transEqDim
	)

                                                                                      {
    debug_printf("%s\n\ttransEqDim=%d\n", __PRETTY_FUNCTION__, transEqDim);
    obj->m_transEqDim = transEqDim;
    obj->m_transEqClassRep.resize(transEqDim);
    obj->m_transPeriod.resize(transEqDim);
    obj->m_stateToTransEqClass.resize(obj->dim());
}