Bit_Compare.cu File Reference

Functions
__global__ void	Bit_gpu (Integer_t *dResPtr, Integer_t const LD, Integer_t const N, Integer_t const dloc, Integer_t const nTot)
int	main (int argc, char **argv)

Function Documentation

Bit_gpu()

__global__ void Bit_gpu	(	Integer_t *	dResPtr,
		Integer_t const	LD,
		Integer_t const	N,
		Integer_t const	dloc,
		Integer_t const	nTot
	)

                                                                                                                               {
    int idx = blockIdx.x*blockDim.x +threadIdx.x;
    int idy = blockIdx.y*blockDim.y +threadIdx.y;
 
    if(idx>=nTot || idy>=N) return;
    dResPtr[idx+LD*idy] = Bit(idx, idy, dloc, N);
}

main()

int main	(	int	argc,
		char **	argv
	)

                                {
    if(argc != 3) {
        std::cerr << "Usage: Bit_Compare(Integer_t const N, Integer_t const dloc)" << std::endl;
        std::exit(EX_USAGE);
    }
    Integer_t N    = (Integer_t)std::atoi(argv[1]);
    Integer_t dloc = (Integer_t)std::atoi(argv[2]);
    Integer_t nTot = (Integer_t)std::pow(dloc, N);
 
    magma_init();
    magma_queue_t queue = NULL;
    magma_int_t     dev = 0;
    magma_getdevice( &dev );
    magma_queue_create( dev, &queue );
 
    constexpr Integer_t GPU_UNIT = 32;
    Integer_t const LDT = magma_roundup(nTot, GPU_UNIT);
    matrix_gpu<Integer_t> dRes(LDT, N);
 
    void (*funcPtr)(Integer_t*, Integer_t const, Integer_t const, Integer_t const, Integer_t const);
    funcPtr = Bit_gpu;
    struct cudaFuncAttributes attr; cudaFuncGetAttributes(&attr, funcPtr);
    Integer_t nThread = (Integer_t)sqrt(attr.maxThreadsPerBlock);
    Integer_t nBlock  = (Integer_t)nTot/nThread;
    if(nTot > nBlock*nThread) nBlock += 1;
    GPUconfig conf(dim3(nBlock,nBlock,1), dim3(nThread,nThread,1),0,queue);
    Bit_gpuMatrixElementsInSector<<<conf.dimGrid(),conf.dimBlock(),conf.shared(),conf.stream()>>>(dRes.ptr(), dRes.LD(), N, dloc, nTot);
    cudaDeviceSynchronize();
 
 
    matrix<Integer_t> Res(nTot, N);
    magma_getmatrix(nTot, N, sizeof(Integer_t), dRes.ptr(), dRes.LD(), &*Res.begin(), nTot, queue);
 
    bool flag;
    for(Integer_t n = 0;n < nTot; ++n) {
        flag = true;
        for(Integer_t pos = 0;pos < N; ++pos) {
            if(Res.at(n,pos) != Bit(n, pos, dloc, N)) { flag = false; break; };
        }
        if(flag == false) std::cerr << "Bit(" << n << ") != Bit_gpu" << std::endl;
        else std::cerr << "(OK, n=" << n << ") Bit = Bit_gpu" <<  std::endl;
    }
 
 
 
 
    cudaDeviceSynchronize();
    return 0;
}