cudadl/dpstream.cu

/*
 *  cudadl version 0.8: 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_11_atomic_functions.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+4); // 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;
}

__device__ inline void write_dp(unsigned int*, unsigned long, unsigned long);

// 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 void DPstreamWrite(unsigned long a, unsigned long b)
{
    if (!DPbuffer) return;

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

    write_dp(ourbuffer, a, b);
}

static void dpcallback(void *data, unsigned short threadId,
    unsigned short blockId, string x, unsigned long a, unsigned long b);

void DPstreamParse(void *data)
{
    // Get the number of words in the buffer
    unsigned int bufsize;
    cudaMemcpyFromSymbol(&bufsize, DPbuffertail, sizeof(unsigned int));
    cerr << getpid() << " " << bufsize / DPrecordsize << " DPs\n";
    unsigned int *dpbuf = (unsigned int*)calloc(bufsize, sizeof(unsigned int));
    if (dpbuf) {
	cudaMemcpy(dpbuf, DPbuffer_device, bufsize * sizeof(unsigned int),
	    cudaMemcpyDeviceToHost);

	for (unsigned int *dp=dpbuf; dp < dpbuf+bufsize; dp += DPrecordsize) {
	    unsigned short threadId = (unsigned short)(dp[0]);
	    unsigned short blockId = dp[0]>>16;
	    unsigned long a = dp[WORDS+2];
	    a <<= 32;
	    a += dp[WORDS+1];
	    unsigned long b = dp[WORDS+4];
	    b <<= 32;
	    b += dp[WORDS+3];
	    string x((const char *)(dp+1), WORDS*sizeof(unsigned int));

	    dpcallback(data, threadId, blockId, x, a, b);
	}
    }
    free(dpbuf);
    // Reset the buffer
    unsigned int zero = 0;
    cudaMemcpyToSymbol(DPbuffertail, &zero, sizeof(unsigned int));
}
#endif