natas_hw
/
jidt
mirror of https://github.com/jlizier/jidt
0
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Separated digamma and reduction into host and device functions.

This commit is contained in:
Pedro Martinez Mediano 2017-05-31 02:36:45 +10:00
parent 7a89dd7908
commit 85dd28d13b
2 changed files with 28 additions and 42 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

64
java/source/infodynamics/measures/continuous/kraskov/cuda/gpuKnnLibrary.c
View File

@ -372,44 +372,50 @@ int computeSumDigammas(float *sumDiGammas, int *nx, int *ny, unsigned int N) {
}
int parallelDigammas(float *digammas, int *nx, int *ny, int signallength) {
int *d_nx, *d_ny;
float *d_digammas;
int d_parallelDigammas(float *digammas, float *d_digammas, int *d_nx,
int *d_ny, int signalLength) {
// Kernel parameters
dim3 threads(1,1,1);
dim3 grid(1,1,1);
threads.x = 512;
grid.x = (signallength-1)/threads.x + 1;
grid.x = (signalLength-1)/threads.x + 1;
// Launch kernel
gpuDigammas<<<grid.x, threads.x>>>(d_digammas, d_nx, d_ny, signalLength);
checkCudaErrors( cudaDeviceSynchronize() );
checkCudaErrors( cudaMemcpy(digammas, d_digammas, signalLength * sizeof(float), cudaMemcpyDeviceToHost) );
checkCudaErrors( cudaDeviceSynchronize() );
return 1;
}
int parallelDigammas(float *digammas, int *nx, int *ny, int signalLength) {
int *d_nx, *d_ny;
float *d_digammas;
{
CPerfTimer pt = startTimer("Digammas_upload");
checkCudaErrors( cudaMalloc((void **) &d_nx, signallength * sizeof(int)) );
checkCudaErrors( cudaMalloc((void **) &d_ny, signallength * sizeof(int)) );
checkCudaErrors( cudaMalloc((void **) &d_digammas, signallength * sizeof(float)) );
checkCudaErrors( cudaMalloc((void **) &d_nx, signalLength * sizeof(int)) );
checkCudaErrors( cudaMalloc((void **) &d_ny, signalLength * sizeof(int)) );
checkCudaErrors( cudaMalloc((void **) &d_digammas, signalLength * sizeof(float)) );
checkCudaErrors( cudaMemcpy(d_nx, nx, signallength*sizeof(int), cudaMemcpyHostToDevice) );
checkCudaErrors( cudaMemcpy(d_ny, ny, signallength*sizeof(int), cudaMemcpyHostToDevice) );
checkCudaErrors( cudaMemcpy(d_nx, nx, signalLength*sizeof(int), cudaMemcpyHostToDevice) );
checkCudaErrors( cudaMemcpy(d_ny, ny, signalLength*sizeof(int), cudaMemcpyHostToDevice) );
stopTimer(pt);
}
{
CPerfTimer pt = startTimer("Digammas_kernel");
// Launch kernel
gpuDigammas<<<grid.x, threads.x>>>(d_digammas, d_nx, d_ny, signallength);
checkCudaErrors( cudaDeviceSynchronize() );
d_parallelDigammas(digammas, d_digammas, d_nx, d_ny, signalLength);
stopTimer(pt);
}
{
CPerfTimer pt = startTimer("Digammas_download");
checkCudaErrors( cudaMemcpy(digammas, d_digammas, signallength * sizeof(float), cudaMemcpyDeviceToHost) );
checkCudaErrors( cudaDeviceSynchronize() );
checkCudaErrors( cudaFree(d_nx) );
checkCudaErrors( cudaFree(d_ny) );
checkCudaErrors( cudaFree(d_digammas) );
@ -419,26 +425,6 @@ int parallelDigammas(float *digammas, int *nx, int *ny, int signallength) {
return 1;
}
int computeSumDigammasChunks(float *sumDiGammas, int *nx, int *ny,
int triallength, int nchunks) {
int signallength = triallength * nchunks;
float digammas[signallength];
int err = parallelDigammas(digammas, nx, ny, signallength);
CPerfTimer pt = startTimer("Digammas_sum");
for (int c = 0; c < nchunks; c++) {
float sum = 0;
for (int i = 0; i < triallength; i++) {
sum += digammas[triallength*c + i];
}
sumDiGammas[c] = sum;
}
stopTimer(pt);
return 1;
}
int d_cudaSumDigammas(float *sumDigammas, int *d_nx, int *d_ny,
float *d_digammas, int trialLength, int nchunks) {

6
java/source/infodynamics/measures/continuous/kraskov/cuda/gpuKnnLibrary.h
View File

@ -31,9 +31,6 @@ int findRadiiAlgorithm2(float *radii, const float *data, const int *indexes,
int computeSumDigammas(float *sumDiGammas, int *nx, int *ny, unsigned int N);
int computeSumDigammasChunks(float *sumDiGammas, int *nx, int *ny,
int trialLength, int nchunks);
int parallelDigammas(float *digammas, int *nx, int *ny, int signallength);
int cudaSumDigammas(float *sumDigammas, int *nx, int *ny, int trialLength,
@ -47,6 +44,9 @@ int d_cudaFindRSAll(int* d_bf_npointsrange, float* d_bf_pointset, float* d_bf_qu
float* d_bf_vecradius, int thelier, int nchunks, int pointdim,
int signallength, int useMaxNorm);
int d_parallelDigammas(float *digammas, float *d_digammas, int *d_nx,
int *d_ny, int signalLength);
int d_cudaSumDigammas(float *sumDigammas, int *d_nx, int *d_ny,
float *d_digammas, int trialLength, int nchunks);