cudadl/dpstream.cu

/*
 *  cudadl version 0.9: Compute discrete logs in smooth group orders
 *                      using CUDA
 *  Copyright (C) 2012 by Ryan Henry and Ian Goldberg
 *                        {rhenry,iang}@cs.uwaterloo.ca
 *
 *  This program is free software: you can redistribute it and/or modify
 *  it under the terms of version 3 of the GNU General Public License as
 *  published by the Free Software Foundation.
 *
 *  This program is distributed in the hope that it will be useful,
 *  but WITHOUT ANY WARRANTY; without even the implied warranty of
 *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 *  GNU General Public License for more details.
 *
 *  You should have received a copy of the GNU General Public License
 *  along with this program.  If not, see <http://www.gnu.org/licenses/>.
 */

#ifndef __DPSTREAM_H__
#define __DPSTREAM_H__

#include <sm_20_atomic_functions.h>
#include <unistd.h>
#include <stdlib.h>

/* Routines to allow CUDA threads to write distinguished points into
   a buffer that will be read by the host */

__constant__ static unsigned int *DPbuffer = NULL;
__constant__ static unsigned int DPbuffersize = 0;  // in words
static unsigned int *DPbuffer_device = NULL;
__device__ volatile static unsigned int DPbuffertail = 0;  // in words

static const unsigned int DPrecordsize = (1+WORDS+6); // in words
static unsigned int DPbufsize = 0;

cudaError_t DPstreamInit(size_t numrecords)
{
    // Allocate the buffer
    DPbufsize = DPrecordsize * numrecords;

    if (cudaMalloc((void **)&DPbuffer_device, DPbufsize * sizeof(unsigned int))
	    != cudaSuccess) {
	return cudaErrorInitializationError;
    }
    // Zero it
    cudaMemset(DPbuffer_device, 0, DPbufsize * sizeof(unsigned int));
    // Initialize the device-visible pointers
    unsigned int zero = 0;
    cudaMemcpyToSymbol(DPbuffer, &DPbuffer_device, sizeof(unsigned int*));
    cudaMemcpyToSymbol(DPbuffersize, &DPbufsize, sizeof(unsigned int));
    cudaMemcpyToSymbol(DPbuffertail, &zero, sizeof(unsigned int));

    return cudaSuccess;
}

void DPstreamEnd(void)
{
    if (DPbuffer_device) {
	cudaFree(DPbuffer_device);
    }
    DPbuffer_device = NULL;
}

// Call this from the device to write the following into the buffer:
//   block and thread indices (1 word total)
//   x (WORDS words)
//   a (2 words)
//   b (2 words)
__device__ inline unsigned int *DPstreamAlloc()
{
    if (!DPbuffer) return NULL;

    // Atomically increment the DPbuffertail
    if (DPbuffertail >= DPbuffersize) return NULL;
    unsigned int ouroffset = atomicAdd((unsigned int *)&DPbuffertail,
		(unsigned int) (DPrecordsize));
    if (ouroffset + DPrecordsize > DPbuffersize) {
	atomicSub((unsigned int *)&DPbuffertail, (unsigned int)DPrecordsize);
	return NULL;
    }
    unsigned int *ourbuffer = DPbuffer + ouroffset;
    return ourbuffer;
}

#ifdef DERANDOMIZE
int DPrecordComparator(const void* p1, const void* p2)
{
    const unsigned int* p1_uint = (const unsigned int*)p1;
    const unsigned int* p2_uint = (const unsigned int*)p2;

    for (int i=0; i<DPrecordsize; i++) {
	// compare each word
	if (p1_uint[i] < p2_uint[i]) {
	    return 1;
	} else if (p1_uint[i] > p2_uint[i]) {
	    return -1;
	}
    }

    return 0;
}
#endif

// dp points to an array of WORDS+7 unsigned ints:
// - 1 word of threadID/blockID
// - WORDS words of the dp value
// - 3 words of a
// - 3 words of b
extern bool dpcallback(void *data, unsigned int *dpwords);

// Return true if we should stop computing
bool DPstreamParse(void *data, unsigned int *num_dps, bool *filled_dp_buffer)
{
    // Get the number of words in the buffer
    unsigned int bufsize;
    cudaMemcpyFromSymbol(&bufsize, DPbuffertail, sizeof(unsigned int));
    unsigned int max_bufsize;
    cudaMemcpyFromSymbol(&max_bufsize, DPbuffersize, sizeof(unsigned int));

    *num_dps = bufsize/DPrecordsize;
    *filled_dp_buffer = (bufsize == max_bufsize);

    unsigned int *dpbuf = (unsigned int*)calloc(bufsize, sizeof(unsigned int));
    bool stop_computing = false;
    if (dpbuf) {
	cudaMemcpy(dpbuf, DPbuffer_device, bufsize * sizeof(unsigned int),
	    cudaMemcpyDeviceToHost);

#ifdef DERANDOMIZE
	qsort(dpbuf, *num_dps, DPrecordsize*sizeof(unsigned int), DPrecordComparator);
	// make the order of the DPs deterministic after each kernel launch
#endif

	for (unsigned int *dp=dpbuf; dp < dpbuf+bufsize; dp += DPrecordsize) {
	    if (dpcallback(data, dp)) {
		stop_computing = true;
	    }
	}
    }
    free(dpbuf);
    // Reset the buffer
    unsigned int zero = 0;
    cudaMemcpyToSymbol(DPbuffertail, &zero, sizeof(unsigned int));

    return stop_computing;
}
#endif