Markov::API::CUDA Namespace Reference

Namespace for objects requiring CUDA libraries. More...

Namespaces
	Random
	Namespace for Random engines operable under device space.

Classes
class	CUDADeviceController
	Controller class for CUDA device. More...
class	CUDAModelMatrix
	Extension of Markov::API::ModelMatrix which is modified to run on GPU devices. More...

Functions
__global__ void	FastRandomWalkCUDAKernel (unsigned long int n, int minLen, int maxLen, char *outputBuffer, char *matrixIndex, long int *totalEdgeWeights, long int *valueMatrix, char *edgeMatrix, int matrixSize, int memoryPerKernelGrid, unsigned long *seed)
	CUDA kernel for the FastRandomWalk operation. More...
__device__ char *	strchr (char *p, char c, int s_len)
	srtchr implementation on device space More...

Detailed Description

Namespace for objects requiring CUDA libraries.

Function Documentation

FastRandomWalkCUDAKernel()

__global__ void Markov::API::CUDA::FastRandomWalkCUDAKernel	(	unsigned long int	n,
		int	minLen,
		int	maxLen,
		char *	outputBuffer,
		char *	matrixIndex,
		long int *	totalEdgeWeights,
		long int *	valueMatrix,
		char *	edgeMatrix,
		int	matrixSize,
		int	memoryPerKernelGrid,
		unsigned long *	seed
	)

CUDA kernel for the FastRandomWalk operation.

Will be initiated by CPU and continued by GPU (global tag)

Parameters

n	- Number of passwords to generate.
minlen	- minimum string length for a single generation
maxLen	- maximum string length for a single generation
outputBuffer	- VRAM ptr to the output buffer
matrixIndex	- VRAM ptr to the matrix indices
totalEdgeWeights	- VRAM ptr to the totalEdgeWeights array
valueMatrix	- VRAM ptr to the edge weights array
edgeMatrix	- VRAM ptr to the edge representations array
matrixSize	- Size of the matrix dimensions
memoryPerKernelGrid	- Maximum memory usage per kernel grid
seed	- seed chunk to generate the random from (generated & used by Marsaglia)

Definition at line 194 of file cudaModelMatrix.cu.

                                                                                                                                                         {
        
        int kernelWorkerIndex = threadIdx.x + blockIdx.x * blockDim.x;
 
        if(n==0) return;
 
        char* e;
        int index = 0;
        char next;
        int len=0;
        long int selection;
        char cur;
        long int bufferctr = 0;
        unsigned long int *x,*y,*z,t;
        char* res = &outputBuffer[kernelWorkerIndex*memoryPerKernelGrid];
        x=&seed[kernelWorkerIndex*3];
        y=&seed[kernelWorkerIndex*3+1];
        z=&seed[kernelWorkerIndex*3+2];
        for (int i = 0; i < n; i++) {
            cur=199;
            len=0;
            while (true) {
                e = strchr(matrixIndex, cur, matrixSize);
                index = e - matrixIndex;
                /*selection = Markov::API::CUDA::Random::devrandom(
                    seed[kernelWorkerIndex*3],
                    seed[kernelWorkerIndex*3+1],
                    seed[kernelWorkerIndex*3+2]) % totalEdgeWeights[index];*/
                *x ^= *x << 16;
                *x ^= *x >> 5;
                *x ^= *x << 1;
 
                t = *x;
                *x = *y;
                *y = *z;
                *z = t ^ *x ^ *y;
                selection = *z % totalEdgeWeights[index];
            for(int j=0;j<matrixSize-1;j++){
                    selection -= valueMatrix[index*matrixSize + j];
                    if (selection < 0){
                        next = edgeMatrix[index*sizeof(char)*matrixSize + j];
                        break;
                    }
                }
 
                if (len >= maxLen)  break;
                else if ((next < 0) && (len < minLen)) continue;
                else if (next < 0) break;  
                cur = next;
                res[bufferctr + len++] = cur;
            }
            res[bufferctr + len++] = '\n';
            bufferctr+=len;
        }
        res[bufferctr] = '\0';
    }

strchr()

__device__ char * Markov::API::CUDA::strchr	(	char *	p,
		char	c,
		int	s_len
	)

srtchr implementation on device space

Fint the first matching index of a string

Parameters

p	- string to check
c	- character to match
s_len	- maximum string length

Returns

pointer to the match

Definition at line 252 of file cudaModelMatrix.cu.

                                                       {
       for (;; ++p, s_len--) {
            if (*p ==  c)
                return((char *)p);
            if (!*p)
                return((char *)NULL);
        }
    }