Пример #1
0
bool
OsdCLKernelBundle::Compile(cl_context clContext,
                           int numVertexElements, int numVaryingElements) {

    cl_int ciErrNum;

    _numVertexElements = numVertexElements;
    _numVaryingElements = numVaryingElements;

    char constantDefine[256];
    snprintf(constantDefine, sizeof(constantDefine),
             "#define NUM_VERTEX_ELEMENTS %d\n"
             "#define NUM_VARYING_ELEMENTS %d\n",
             numVertexElements, numVaryingElements);

    const char *sources[] = { constantDefine, clSource };

    _clProgram = clCreateProgramWithSource(clContext, 2, sources, 0, &ciErrNum);
    CL_CHECK_ERROR(ciErrNum, "clCreateProgramWithSource\n");

    ciErrNum = clBuildProgram(_clProgram, 0, NULL, NULL, NULL, NULL);
    if (ciErrNum != CL_SUCCESS) {
        OsdError(OSD_CL_PROGRAM_BUILD_ERROR, "CLerr=%d", ciErrNum);

        cl_int numDevices = 0;
        clGetContextInfo(clContext, CL_CONTEXT_NUM_DEVICES,
                         sizeof(cl_uint), &numDevices, NULL);
        cl_device_id *devices = new cl_device_id[numDevices];
        clGetContextInfo(clContext, CL_CONTEXT_DEVICES,
                         sizeof(cl_device_id)*numDevices, devices, NULL);
        for (int i = 0; i < numDevices; ++i) {
            char cBuildLog[10240];
            clGetProgramBuildInfo(_clProgram, devices[i], CL_PROGRAM_BUILD_LOG,
                                  sizeof(cBuildLog), cBuildLog, NULL);
            OsdError(OSD_CL_PROGRAM_BUILD_ERROR, cBuildLog);
        }
        delete[] devices;
        return false;
    }

    _clBilinearEdge   = buildKernel(_clProgram, "computeBilinearEdge");
    _clBilinearVertex = buildKernel(_clProgram, "computeBilinearVertex");
    _clCatmarkFace    = buildKernel(_clProgram, "computeFace");
    _clCatmarkEdge    = buildKernel(_clProgram, "computeEdge");
    _clCatmarkVertexA = buildKernel(_clProgram, "computeVertexA");
    _clCatmarkVertexB = buildKernel(_clProgram, "computeVertexB");
    _clLoopEdge       = buildKernel(_clProgram, "computeEdge");
    _clLoopVertexA    = buildKernel(_clProgram, "computeVertexA");
    _clLoopVertexB    = buildKernel(_clProgram, "computeLoopVertexB");
    _clVertexEditAdd  = buildKernel(_clProgram, "editVertexAdd");

    return true;
}
Пример #2
0
  kernel device::buildKernelFromString(const std::string &content,
                                       const std::string &kernelName,
                                       const occa::properties &props) const {

    occa::properties allProps = props + kernelProperties();
    allProps["mode"] = mode();

    hash_t kernelHash = (hash()
                         ^ occa::hash(allProps)
                         ^ occa::hash(content));

    io::lock_t lock(kernelHash, "occa-device");
    std::string stringSourceFile = io::hashDir(kernelHash);
    stringSourceFile += "stringSource.okl";

    if (lock.isMine()) {
      if (!sys::fileExists(stringSourceFile)) {
        io::write(stringSourceFile, content);
      }
      lock.release();
    }

    return buildKernel(stringSourceFile,
                       kernelName,
                       props);
  }
Пример #3
0
bool DeconvolutionTool::doDeconvolution(QImage *inputImage, QImage *outputImage, Blur* blur) {
    isProcessingCancelled = false;
    // Create kernel
    buildKernel(kernelMatrix, width, height, blur);
    fftw_execute(realForwardKernelPlan);
    // Fill processingContext
    ProcessingContext* processingContext = new ProcessingContext();
    processingContext->inputImage = inputImage;
    processingContext->outputImage = outputImage;
    processingContext->inputImageMatrix = inputImageMatrix;
    processingContext->outputImageMatrix = outputImageMatrix;
    processingContext->kernelFFT = kernelFFT;
    processingContext->width = width;
    processingContext->height = height;
    processingContext->blur = blur;

    if (blur->mode == PREVIEW_GRAY) {
        doDeconvolutionForChannel(processingContext, GRAY);
    } else {
        QString progressText= "High-Quality";
        if (blur->mode == PREVIEW_COLOR) {
            progressText = "Color Preview";
        }
        setProgressInterval(10,40, progressText);
        doDeconvolutionForChannel(processingContext, RED);
        setProgressInterval(40,70, progressText);
        doDeconvolutionForChannel(processingContext, GREEN);
        setProgressInterval(70,100, progressText);
        doDeconvolutionForChannel(processingContext, BLUE);

    }

    delete(processingContext);
    return !isProcessingCancelled;
}
Пример #4
0
void OpenCL::init(int isGPU)
{
	if (isGPU)
		getDevices(CL_DEVICE_TYPE_GPU);
	else
		getDevices(CL_DEVICE_TYPE_CPU);

	buildKernel();
}
Пример #5
0
PTEX_NAMESPACE_BEGIN

void PtexTriangleFilter::eval(float* result, int firstChan, int nChannels,
                              int faceid, float u, float v,
                              float uw1, float vw1, float uw2, float vw2,
                              float width, float blur)
{
    // init
    if (!_tx || nChannels <= 0) return;
    if (faceid < 0 || faceid >= _tx->numFaces()) return;
    _ntxchan = _tx->numChannels();
    _dt = _tx->dataType();
    _firstChanOffset = firstChan*DataSize(_dt);
    _nchan = PtexUtils::min(nChannels, _ntxchan-firstChan);

    // get face info
    const FaceInfo& f = _tx->getFaceInfo(faceid);

    // if neighborhood is constant, just return constant value of face
    if (f.isNeighborhoodConstant()) {
        PtexPtr<PtexFaceData> data ( _tx->getData(faceid, 0) );
        if (data) {
            char* d = (char*) data->getData() + _firstChanOffset;
            Ptex::ConvertToFloat(result, d, _dt, _nchan);
        }
        return;
    }

    // clamp u and v
    u = PtexUtils::clamp(u, 0.0f, 1.0f);
    v = PtexUtils::clamp(v, 0.0f, 1.0f);

    // build kernel
    PtexTriangleKernel k;
    buildKernel(k, u, v, uw1, vw1, uw2, vw2, width, blur, f.res);

    // accumulate the weight as we apply
    _weight = 0;

    // allocate temporary result
    _result = (float*) alloca(sizeof(float)*_nchan);
    memset(_result, 0, sizeof(float)*_nchan);

    // apply to faces
    splitAndApply(k, faceid, f);

    // normalize (both for data type and cumulative kernel weight applied)
    // and output result
    float scale = 1.0f / (_weight * OneValue(_dt));
    for (int i = 0; i < _nchan; i++) result[i] = float(_result[i] * scale);

    // clear temp result
    _result = 0;
}
Пример #6
0
  void device::loadKernels(const std::string &library) {
    // TODO 1.1: Load kernels
#if 0
    std::string devHash = hash().toFullString();
    strVector dirs = io::directories("occa://" + library);
    const int dirCount = (int) dirs.size();
    int kernelsLoaded = 0;

    for (int d = 0; d < dirCount; ++d) {
      const std::string buildFile = dirs[d] + kc::buildFile;

      if (!sys::fileExists(buildFile)) {
        continue;
      }

      json info = json::read(buildFile)["info"];
      if ((std::string) info["device/hash"] != devHash) {
        continue;
      }
      ++kernelsLoaded;

      const std::string sourceFilename = dirs[d] + kc::parsedSourceFile;

      json &kInfo = info["kernel"];
      hash_t hash = hash_t::fromString(kInfo["hash"]);
      jsonArray metadataArray = kInfo["metadata"].array();
      occa::properties kernelProps = kInfo["props"];

      // Ignore how the kernel was setup, turn off verbose
      kernelProps["verbose"] = false;

      const int kernels = metadataArray.size();
      for (int k = 0; k < kernels; ++k) {
        buildKernel(sourceFilename,
                    hash,
                    kernelProps,
                    lang::kernelMetadata::fromJson(metadataArray[k]));
      }
    }

    // Print loaded info
    if (properties().get("verbose", false) && kernelsLoaded) {
      std::cout << "Loaded " << kernelsLoaded;
      if (library.size()) {
        std::cout << " ["<< library << "]";
      } else {
        std::cout << " cached";
      }
      std::cout << ((kernelsLoaded == 1)
                    ? " kernel\n"
                    : " kernels\n");
    }
#endif
  }
Пример #7
0
bool Engine::start() {
    bool retval {false};

    m_media = std::make_unique<Media>();

    if (m_media->init()) {
        buildKernel();
        if (buildModules()) {
            m_kernel->start();
            retval = true;
        }
    }

    return retval;
}
Пример #8
0
bool
OsdClKernelDispatcher::ClKernel::Compile(cl_context clContext, int numVertexElements, int numVaryingElements) {

    cl_int ciErrNum;

    _numVertexElements = numVertexElements;
    _numVaryingElements = numVaryingElements;

    char constantDefine[256];
    snprintf(constantDefine, 256, "#define NUM_VERTEX_ELEMENTS %d\n"
             "#define NUM_VARYING_ELEMENTS %d\n", numVertexElements, numVaryingElements);

    const char *sources[] = { constantDefine, clSource };

    _clProgram = clCreateProgramWithSource(clContext, 2, sources, 0, &ciErrNum);
    CL_CHECK_ERROR(ciErrNum, "clCreateProgramWithSource\n");

    ciErrNum = clBuildProgram(_clProgram, 0, NULL, NULL, NULL, NULL);
    if (ciErrNum != CL_SUCCESS) {
        OSD_ERROR("ERROR in clBuildProgram %d\n", ciErrNum);
        char cBuildLog[10240];
        clGetProgramBuildInfo(_clProgram, _clDevice, CL_PROGRAM_BUILD_LOG,
                              sizeof(cBuildLog), cBuildLog, NULL);
        OSD_ERROR(cBuildLog);
        return false;
    }

    // -------
    _clBilinearEdge   = buildKernel(_clProgram, "computeBilinearEdge");
    _clBilinearVertex = buildKernel(_clProgram, "computeBilinearVertex");
    _clCatmarkFace    = buildKernel(_clProgram, "computeFace");
    _clCatmarkEdge    = buildKernel(_clProgram, "computeEdge");
    _clCatmarkVertexA = buildKernel(_clProgram, "computeVertexA");
    _clCatmarkVertexB = buildKernel(_clProgram, "computeVertexB");
    _clLoopEdge       = buildKernel(_clProgram, "computeEdge");
    _clLoopVertexA    = buildKernel(_clProgram, "computeVertexA");
    _clLoopVertexB    = buildKernel(_clProgram, "computeLoopVertexB");

    return true;
}
VolumeMaxCLProcessor::VolumeMaxCLProcessor()
    : Processor()
    , ProcessorKernelOwner(this)
    , inport_("inVolume")
    , outport_("outVolume")
    , volumeRegionSize_("region", "Region size", 8, 1, 100)
    , workGroupSize_("wgsize", "Work group size", ivec3(8), ivec3(0), ivec3(256))
    , useGLSharing_("glsharing", "Use OpenGL sharing", true)
    , supportsVolumeWrite_(false)
    , tmpVolume_(nullptr)
    , kernel_(nullptr) {
    addPort(inport_);
    addPort(outport_);

    addProperty(volumeRegionSize_);
    addProperty(workGroupSize_);
    addProperty(useGLSharing_);

    buildKernel();
}
Пример #10
0
///apply filter effect to the PixelBuffer* buffer passed into this function
void FBlur::applyFilter(PixelBuffer* imageBuffer){
  // if kernel is already initialized, do not need initialize it again.
  kernel = buildKernel(std::round(getFloatParameter()));

  // printKernel();

  if(getName() == "FEdgeDetection"){
    imageBuffer -> convertToLuminance();
  }

  int width = imageBuffer -> getWidth();
  int height = imageBuffer -> getHeight();
  //create a new pixel buffer for storing the convolution result.
  PixelBuffer* newImageBuffer
    = new PixelBuffer(width, height, imageBuffer -> getBackgroundColor());
  for(int i = 0; i < width; i++){
    for(int j = 0; j < height; j++){
      float newRed = 0;
      float newGreen = 0;
      float newBlue = 0;
      for(size_t filterI = 0; filterI < kernel.size(); filterI++){
        for(size_t filterJ = 0; filterJ < kernel[filterI].size(); filterJ++){
          //The location imageI and imageJ is calculated so that
          //for the center element of the filter it'll be i, j
          int imageI = (i - kernel.size()/2 + filterI + width) % width;
          int imageJ = (j - kernel[filterI].size()/2 + filterJ + height) % height;
          ColorData currPixel = imageBuffer -> getPixel(imageI, imageJ);
          newRed += currPixel.getRed() * kernel[filterI][filterJ];
          newGreen += currPixel.getGreen()*kernel[filterI][filterJ];
          newBlue += currPixel.getBlue()*kernel[filterI][filterJ];
        }
      }
      ColorData newPixel = ColorData(newRed, newGreen, newBlue);
      newPixel = newPixel.clampedColor();
      newImageBuffer -> setPixel(i, j, newPixel);
    }
  }
  newImageBuffer -> copyPixelBuffer(newImageBuffer, imageBuffer);
  delete newImageBuffer;
}
Пример #11
0
JNIEXPORT jint JNICALL Java_pt_floraon_ecospace_nativeFunctions_computeKernelDensities(JNIEnv *env, jclass obj, jstring filename, jstring outfilename, jintArray variables,jint freqthresh,jfloat sigmapercent,jboolean downweight) {
// NOTE: "variables" must be in increasing order! 0 is latitude, 1 is longitude, 2... are the other climatic variables
	const char *pfilename=(*env)->GetStringUTFChars(env, filename , NULL );
	const char *poutfilename=(*env)->GetStringUTFChars(env, outfilename , NULL );
	
	FILE *varsfile,*densfile,*freqfile;
	int nvars,i,j,k,*freqs,ntaxafiltered=0,*mapIDs,*outIDs,d1,d2,d1from,d1to,d2from,d2to,d3from,d3to,d1kern,d2kern,sidesq;
	int lastID=-1,d2p,kernelhalfside,kernelside,kernelsidesq;
	register int d3p,d3;
	register float *d3kern;
	jsize nrecs;
	jint ntaxa,*pIDs;
	VARIABLE *vararray;
	VARIABLEHEADER *varheader;
	int nvarstouse=(int)(*env)->GetArrayLength(env,variables);		// how many vars will be used for kernel density
	jint *pvariables=(*env)->GetIntArrayElements(env, variables, 0);
	float sigma;
	float *kernel,*tmpdens;
	unsigned long *weight;
	DENSITY *densities;
	bool anythingtosave=false,skiprec;
	size_t dummy;
	
	varsfile=fopen(STANDARDVARIABLEFILE(pfilename),"r");
	dummy=fread(&nrecs,sizeof(jsize),1,varsfile);
	dummy=fread(&ntaxa,sizeof(jint),1,varsfile);
	dummy=fread(&nvars,sizeof(int),1,varsfile);
//	fseek(varsfile,nvars*sizeof(tmp.filename)+sizeof(long)*ntaxa,SEEK_CUR);
	fseek(varsfile,sizeof(long)*ntaxa,SEEK_CUR);	// skip index

	{	// redirect stdout to file
		int fd;
		fpos_t pos;
		fflush(stdout);
		fgetpos(stdout, &pos);
		fd = dup(fileno(stdout));
		FILE *dummy=freopen("logfile.txt", "a", stdout);
	}


	vararray=malloc(nvarstouse*nrecs*sizeof(VARIABLE));
	varheader=malloc(nvarstouse*sizeof(VARIABLEHEADER));
	pIDs=malloc(nrecs*sizeof(jint));
	dummy=fread(pIDs,sizeof(jint),nrecs,varsfile);
	fseek(varsfile,2*sizeof(jfloat)*nrecs,SEEK_CUR);	// skip original coordinates
	weight=malloc(sizeof(long)*nrecs);
	dummy=fread(weight,sizeof(long),nrecs,varsfile);

	for(i=0;i<nvarstouse;i++) {
		fseek(varsfile,(sizeof(VARIABLEHEADER) + sizeof(VARIABLE)*nrecs)*(pvariables[i]-(i==0 ? 0 : (pvariables[i-1]+1))),SEEK_CUR);
		dummy=fread(&varheader[i],sizeof(VARIABLEHEADER),1,varsfile);
		printf("Variable %d: min %f max %f\n",pvariables[i],varheader[i].min,varheader[i].max);
		dummy=fread(&vararray[i*nrecs],sizeof(VARIABLE),nrecs,varsfile);
	}
	
// count the # of records of each taxon
	freqs=calloc(ntaxa,sizeof(int));
	for(i=0;i<nrecs;i++) freqs[pIDs[i]]++;
// write out frequencies in a text file (for java)
	freqfile=fopen(FREQANALYSISFILE(pfilename),"w");
	for(i=0;i<ntaxa;i++) fprintf(freqfile,"%d\n",freqs[i]);
	fclose(freqfile);
// count the # of taxa after filtering out rarest	
	for(i=0;i<ntaxa;i++) if(freqs[i] >= freqthresh) ntaxafiltered++;
// create a mapping of IDs: because some IDs were removed, make them sequential without holes (remember that memory is the limiting factor here!)
	mapIDs=malloc(ntaxa*sizeof(int));
	for(i=0;i<ntaxa;i++) mapIDs[i]=-1;
	for(i=0,j=0;i<nrecs;i++) {
		if(freqs[pIDs[i]] >= freqthresh) {
			if(mapIDs[pIDs[i]]==-1) {
				mapIDs[pIDs[i]]=j;
				j++;
			}
		}
	}
//for(i=0;i<ntaxa;i++) printf("%d ",mapIDs[i]);
	
// compute the resolution of the multidimensional space so that not too much memory is occupied
	side=(int)pow((float)(MAXMEMORYPERBATCH)/(float)ntaxafiltered,(float)1/nvarstouse);
	if(side>MAXSIDE) side=MAXSIDE;
//side=40;
	sidesq=side*side;
	sigma=side*sigmapercent;
	arraysize=(int)pow(side,nvarstouse);

	printf("Using a grid with a side of %d cells, in %d variables (dimensions).\n",side,nvarstouse);
	printf("Reading %d variables of %d records of %d taxa (after filtering out those with less than %d occurrences)...\n",nvars,nrecs,ntaxafiltered,freqthresh);
// build the kernel for this case
	kernel=buildKernel(side,sigma,nvarstouse,&kernelhalfside,&kernelside,&kernelsidesq);

// compute densities
// allocate N multidimensional arrays (multidimensional grids to compute kernel densities in each cell)
	densities=malloc(ntaxafiltered*sizeof(DENSITY));
	outIDs=calloc(ntaxafiltered,sizeof(int));

	for(i=0;i<ntaxafiltered;i++) {
		densities[i].density=malloc(arraysize);
		memset(densities[i].density,0,arraysize);
	}

	printf("Computing kernel densities");fflush(stdout);

// scale the variables to the size of the grid
	for(i=0;i<nrecs;i++) {	
		for(k=0;k<nvarstouse;k++) {
			if(vararray[i+nrecs*k]!=RASTERNODATA) vararray[i+nrecs*k]=(vararray[i+nrecs*k]*side)/10000;
		}
	}
	
	#pragma omp parallel private(i,k,skiprec,tmpdens,anythingtosave,d1from,d1to,d1,d1kern,d2,d2from,d2to,d2p,d2kern,d3,d3from,d3to,d3p,d3kern)
	{
		tmpdens=malloc(arraysize*sizeof(float));
		memset(tmpdens,0,arraysize*sizeof(float));

		#pragma omp for
		for(j=0;j<ntaxa;j++) {		// NOTE: this loop doesn't need the records to be sorted, that's why it takes much longer
// TODO: we might have an index here, i.e. for each taxon, a list of tthe respective records, but it's so fast that maybe it's not a big issue, we're talking about a few thousands of taxa only.
			if(freqs[j]<freqthresh) continue;
			anythingtosave=false;
			for(i=0;i<nrecs;i++) {	// iterate through all records in search for this taxon ACK!! give me an index
				if(pIDs[i]!=j) continue;
				for(k=0,skiprec=false;k<nvarstouse;k++) {	// check if any one of the variables is NA. if it is, skip this record
					if(vararray[i+nrecs*k] == RASTERNODATA) {
						skiprec=true;
						break;
					}
				}
				if(skiprec) continue;		// skip NAs
	// now yeah, create density surface by summing the kernels record by record			
	// since it is not computationally feasible more than 3 dimensions, just make the optimized code for each case...
				switch(nvarstouse) {
					case 1:
						d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
						d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
					
						for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;d1<d1to;d1++,d1kern++) {
							tmpdens[d1]+=kernel[d1kern] * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
						}
						anythingtosave=true;
					break;
				
					case 2:
						d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
						d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
						d2from=(vararray[i+nrecs]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs]-kernelhalfside);
						d2to=(vararray[i+nrecs]+kernelhalfside+1>side ? side : vararray[i+nrecs]+kernelhalfside+1);
					
						for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;d1<d1to;d1++,d1kern++) {
							for(d2=d2from,d2p=d1+d2from*side,d2kern=d1kern+((vararray[i+nrecs]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs])*kernelside : 0);d2<d2to;d2++,d2p+=side,d2kern+=kernelside) {
								tmpdens[d2p]+=kernel[d2kern] * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
							}
						}
						anythingtosave=true;
					break;
				
					case 3:
						d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
						d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
						d2from=(vararray[i+nrecs]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs]-kernelhalfside);
						d2to=(vararray[i+nrecs]+kernelhalfside+1>side ? side : vararray[i+nrecs]+kernelhalfside+1);
						d3from=(vararray[i+nrecs*2]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs*2]-kernelhalfside);
						d3to=(vararray[i+nrecs*2]+kernelhalfside+1>side ? side : vararray[i+nrecs*2]+kernelhalfside+1);
						for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;
							d1<d1to;
							d1++,d1kern++) {
							for(d2=d2from,d2p=d1+d2from*side,d2kern=d1kern+((vararray[i+nrecs]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs])*kernelside : 0)
								;d2<d2to
								;d2++,d2p+=side,d2kern+=kernelside) {
								for(d3=d3from,d3p=d2p+d3from*sidesq
										,d3kern=&kernel[d2kern+((vararray[i+nrecs*2]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs*2])*kernelsidesq : 0)]
									;d3<d3to
									;d3++,d3p+=sidesq,d3kern+=kernelsidesq) {
									tmpdens[d3p]+=*d3kern * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
									//kernel[d3kern];
								}
							}
						}
						anythingtosave=true;
					break;
				}
			}	// end record loop
			if(anythingtosave) {
				saveKernelDensity(tmpdens,freqs[j],&densities[mapIDs[j]]);		// save kernel density of previous taxon
				outIDs[mapIDs[j]]=j;
			} else {
				saveKernelDensity(NULL, freqs[j], &densities[mapIDs[j]]);		// save kernel density of previous taxon
				outIDs[mapIDs[j]]=j;
//				outIDs[mapIDs[j]]=-1;
			}
			anythingtosave = false;
			memset(tmpdens,0,arraysize*sizeof(float));
			printf(".");
			fflush(stdout);
		}	// end taxon loop
	}
/*	THIS is the working code. Above still developing.
	for(i=0;i<nrecs;i++) {	// iterate through all records IMPORTANT: records must be sorted by taxon ID
		if(freqs[pIDs[i]]>=freqthresh) {
			if(pIDs[i]!=lastID) {	// this record is already a different species
				if(anythingtosave) {
					saveKernelDensity(tmpdens,freqs[lastID],&densities[mapIDs[lastID]]);		// save kernel density of previous taxon
					outIDs[mapIDs[lastID]]=lastID;
				} else outIDs[mapIDs[lastID]]=-1;
				anythingtosave=false;
				memset(tmpdens,0,arraysize*sizeof(float));
				lastID=pIDs[i];
				printf(".");
				fflush(stdout);
			}
// scale the variables to the size of the grid
			for(j=0,skiprec=false;j<nvarstouse;j++) {
				if(vararray[i+nrecs*j]==RASTERNODATA) {
					skiprec=true;
					continue;
				} else vararray[i+nrecs*j]=(vararray[i+nrecs*j]*side)/10000;
			}
			if(skiprec) continue;		// skip NAs
// now yeah, create density surface by summing the kernels record by record			
// since it is not computationally feasible more than 3 dimensions, just make the optimized code for each case...
			switch(nvarstouse) {
				case 1:
					d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
					d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
					
					for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;d1<d1to;d1++,d1kern++) {
						tmpdens[d1]+=kernel[d1kern] * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
					}
					anythingtosave=true;
				break;
				
				case 2:
					d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
					d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
					d2from=(vararray[i+nrecs]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs]-kernelhalfside);
					d2to=(vararray[i+nrecs]+kernelhalfside+1>side ? side : vararray[i+nrecs]+kernelhalfside+1);
					
					for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;d1<d1to;d1++,d1kern++) {
						for(d2=d2from,d2p=d1+d2from*side,d2kern=d1kern+((vararray[i+nrecs]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs])*kernelside : 0);d2<d2to;d2++,d2p+=side,d2kern+=kernelside) {
							tmpdens[d2p]+=kernel[d2kern] * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
						}
					}
					anythingtosave=true;
				break;
				
				case 3:
					d1from=(vararray[i]-kernelhalfside)<0 ? 0 : (vararray[i]-kernelhalfside);
					d1to=(vararray[i]+kernelhalfside+1>side ? side : vararray[i]+kernelhalfside+1);
					d2from=(vararray[i+nrecs]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs]-kernelhalfside);
					d2to=(vararray[i+nrecs]+kernelhalfside+1>side ? side : vararray[i+nrecs]+kernelhalfside+1);
					d3from=(vararray[i+nrecs*2]-kernelhalfside)<0 ? 0 : (vararray[i+nrecs*2]-kernelhalfside);
					d3to=(vararray[i+nrecs*2]+kernelhalfside+1>side ? side : vararray[i+nrecs*2]+kernelhalfside+1);
					for(d1=d1from,d1kern=vararray[i]-kernelhalfside<0 ? kernelhalfside-vararray[i] : 0;
						d1<d1to;
						d1++,d1kern++) {
						for(d2=d2from,d2p=d1+d2from*side,d2kern=d1kern+((vararray[i+nrecs]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs])*kernelside : 0)
							;d2<d2to
							;d2++,d2p+=side,d2kern+=kernelside) {
							for(d3=d3from,d3p=d2p+d3from*sidesq
									,d3kern=&kernel[d2kern+((vararray[i+nrecs*2]-kernelhalfside)<0 ? (kernelhalfside-vararray[i+nrecs*2])*kernelsidesq : 0)]
								;d3<d3to
								;d3++,d3p+=sidesq,d3kern+=kernelsidesq) {
								tmpdens[d3p]+=*d3kern * (downweight ? ((float)weight[i]/MULTIPLIER) : 1);
								//kernel[d3kern];
							}
						}
					}
					anythingtosave=true;
				break;
			}
		}
	}
	saveKernelDensity(tmpdens,freqs[lastID],&densities[mapIDs[lastID]]);		// save kernel density of the last taxon
	outIDs[mapIDs[lastID]]=lastID;
*/

// scale all densities to the absolute maximum
/*	unsigned long maxmax=0;
	float factor;
	for(i=0;i<ntaxafiltered;i++) if(maxmax<densities[i].max) maxmax=densities[i].max;
	for(i=0;i<ntaxafiltered;i++) {
		factor=(float)densities[i].max/maxmax;
		for(j=0;j<arraysize;j++) densities[i].density[j]=(unsigned char)((float)densities[i].density[j]*factor);
	}
	*/
// now write to output file
	printf("\nWriting file...\n");
	densfile=fopen(DENSITYFILE(pfilename,poutfilename),"w");
	fwrite(&ntaxafiltered,sizeof(int),1,densfile);		// how many densities in file
	fwrite(outIDs,sizeof(int),ntaxafiltered,densfile);	// the real taxon IDs of each density
	fwrite(&side,sizeof(int),1,densfile);				// the size of the grid
	fwrite(&nvarstouse,sizeof(int),1,densfile);			// the number of variables
	for(k=0;k<ntaxafiltered;k++) {						// now the densities!
		fwrite(&densities[k],sizeof(DENSITY),1,densfile);	// of course, the pointer will be meaningless
		fwrite(densities[k].density,arraysize,1,densfile);
	}
	fclose(densfile);
#ifdef VERBOSE	
	for(k=3;k<4;k++) {	//ntaxafiltered
		printf("************* Taxon Nrecs: %d ************\n",densities[k].nrecords);
		switch(nvarstouse) {
			case 2:
				for(d1=0;d1<side;d1++) {
					for(d2=0;d2<side;d2++) {
							printf("%3d",densities[k].density[d1+d2*side]);
					}
					printf("\n");
				}
			break;
			
			case 3:
				for(d1=0;d1<5;d1++) {
					for(d2=0;d2<side;d2++) {
						for(d3=0;d3<side;d3++) {
							printf("%3d",densities[k].density[d1+d2*side+d3*sidesq]);
						}
						printf("\n");
					}
					printf("*****************\n");
				}
			break;
		}
		printf("********SUM: %f NRECS: %d*********\n",sum,densities[k].nrecords);
	}
	for(k=0;k<ntaxafiltered;k++) {
		for(d1=0,sum=0;d1<arraysize;d1++) sum+=(float)densities[k].density[d1]*densities[k].max/255.0;
		printf("SUM: %f NRECS: %d\n",sum,densities[k].nrecords);
	}
#endif
	
	
	fclose(varsfile);
	free(pIDs);
	free(freqs);
	free(vararray);
	free(varheader);
	free(kernel);
	free(mapIDs);
	free(tmpdens);
	for(i=0;i<ntaxafiltered;i++) free(densities[i].density);
	free(densities);
	free(outIDs);
	(*env)->ReleaseStringUTFChars(env,filename,pfilename);
	(*env)->ReleaseStringUTFChars(env,outfilename,poutfilename);
	return 1;
}
void PtexSeparableFilter::eval(float* result, int firstChan, int nChannels,
			       int faceid, float u, float v, 
			       float uw1, float vw1, float uw2, float vw2,
			       float width, float blur)
{
    // init
    if (!_tx || nChannels <= 0) return;
    if (faceid < 0 || faceid >= _tx->numFaces()) return;
    _ntxchan = _tx->numChannels();
    _dt = _tx->dataType();
    _firstChanOffset = firstChan*DataSize(_dt);
    _nchan = PtexUtils::min(nChannels, _ntxchan-firstChan);

    // get face info
    const FaceInfo& f = _tx->getFaceInfo(faceid);

    // if neighborhood is constant, just return constant value of face
    if (f.isNeighborhoodConstant()) {
	PtexPtr<PtexFaceData> data ( _tx->getData(faceid, 0) );
	if (data) {
	    char* d = (char*) data->getData() + _firstChanOffset;
	    Ptex::ConvertToFloat(result, d, _dt, _nchan);
	}
	return;
    }

    // find filter width as bounding box of vectors w1 and w2
    float uw = fabs(uw1) + fabs(uw2), vw = fabs(vw1) + fabs(vw2);

    // handle border modes
    switch (_uMode) {
    case m_clamp: u = PtexUtils::clamp(u, 0.0f, 1.0f); break;
    case m_periodic: u = u-floor(u); break;
    case m_black: break; // do nothing
    }

    switch (_vMode) {
    case m_clamp: v = PtexUtils::clamp(v, 0.0f, 1.0f); break;
    case m_periodic: v = v-floor(v);
    case m_black: break; // do nothing
    }

    // build kernel
    PtexSeparableKernel k;
    if (f.isSubface()) {
	// for a subface, build the kernel as if it were on a main face and then downres
	uw = uw * width + blur * 2;
	vw = vw * width + blur * 2;
	buildKernel(k, u*.5, v*.5, uw*.5, vw*.5, f.res);
	if (k.res.ulog2 == 0) k.upresU();
	if (k.res.vlog2 == 0) k.upresV();
	k.res.ulog2--; k.res.vlog2--;
    }
    else {
	uw = uw * width + blur;
	vw = vw * width + blur;
	buildKernel(k, u, v, uw, vw, f.res);
    }
    k.stripZeros();

    // check kernel (debug only)
    assert(k.uw > 0 && k.vw > 0);
    assert(k.uw <= PtexSeparableKernel::kmax && k.vw <= PtexSeparableKernel::kmax);
    _weight = k.weight();

    // allocate temporary double-precision result
    _result = (double*) alloca(sizeof(double)*_nchan);
    memset(_result, 0, sizeof(double)*_nchan);

    // apply to faces
    splitAndApply(k, faceid, f);

    // normalize (both for data type and cumulative kernel weight applied)
    // and output result
    double scale = 1.0 / (_weight * OneValue(_dt));
    for (int i = 0; i < _nchan; i++) result[i] = float(_result[i] * scale);

    // clear temp result
    _result = 0;
}
Пример #13
0
CLKernel *EasyCL::buildKernel(string kernelfilepath, string kernelname) {
    return buildKernel(kernelfilepath, kernelname, "");
}
Пример #14
0
void VolumeMaxCLProcessor::initialize() {
    Processor::initialize();
    buildKernel();
}
Пример #15
0
// Serial ray casting
unsigned char* raycast_serial(unsigned char* data, unsigned char* region){
    unsigned char* image = (unsigned char*)malloc(sizeof(unsigned char)*IMAGE_DIM*IMAGE_DIM);
    
    // Camera/eye position, and direction of viewing. These can be changed to look
    // at the volume from different angles.
    float3 camera = {.x=1000,.y=1000,.z=1000};
    float3 forward = {.x=-1, .y=-1, .z=-1};
    float3 z_axis = {.x=0, .y=0, .z = 1};
    
    // Finding vectors aligned with the axis of the image
    float3 right = cross(forward, z_axis);
    float3 up = cross(right, forward);
    
    // Creating unity lenght vectors
    forward = normalize(forward);
    right = normalize(right);
    up = normalize(up);
    
    float fov = 3.14/4;
    float pixel_width = tan(fov/2.0)/(IMAGE_DIM/2);
    float step_size = 0.5;
    
    // For each pixel
    for(int y = -(IMAGE_DIM/2); y < (IMAGE_DIM/2); y++){
        for(int x = -(IMAGE_DIM/2); x < (IMAGE_DIM/2); x++){
            
            // Find the ray for this pixel
            float3 screen_center = add(camera, forward);
            float3 ray = add(add(screen_center, scale(right, x*pixel_width)), scale(up, y*pixel_width));
            ray = add(ray, scale(camera, -1));
            ray = normalize(ray);
            float3 pos = camera;
            
            // Move along the ray, we stop if the color becomes completely white,
            // or we've done 5000 iterations (5000 is a bit arbitrary, it needs 
            // to be big enough to let rays pass through the entire volume)
            int i = 0;
            float color = 0;
            while(color < 255 && i < 5000){
                i++;
                pos = add(pos, scale(ray, step_size));          // Update position
                int r = value_at(pos, region);                  // Check if we're in the region
                color += value_at(pos, data)*(0.01 + r) ;       // Update the color based on data value, and if we're in the region
            }
            
            // Write final color to image
            image[(y+(IMAGE_DIM/2)) * IMAGE_DIM + (x+(IMAGE_DIM/2))] = color > 255 ? 255 : color;
        }
    }
    
    return image;
}


// Check if two values are similar, threshold can be changed.
int similar(unsigned char* data, int3 a, int3 b){
    unsigned char va = data[a.z * DATA_DIM*DATA_DIM + a.y*DATA_DIM + a.x];
    unsigned char vb = data[b.z * DATA_DIM*DATA_DIM + b.y*DATA_DIM + b.x];
    
    int i = abs(va-vb) < 1;
    return i;
}


// Serial region growing, same algorithm as in assignment 2
unsigned char* grow_region_serial(unsigned char* data){
    unsigned char* region = (unsigned char*)calloc(sizeof(unsigned char), DATA_DIM*DATA_DIM*DATA_DIM);
    
    stack_t* stack = new_stack();
    
    int3 seed = {.x=50, .y=300, .z=300};
    push(stack, seed);
    region[seed.z *DATA_DIM*DATA_DIM + seed.y*DATA_DIM + seed.x] = 1;
    
    int dx[6] = {-1,1,0,0,0,0};
    int dy[6] = {0,0,-1,1,0,0};
    int dz[6] = {0,0,0,0,-1,1};
    
    while(stack->size > 0){
        int3 pixel = pop(stack);
        for(int n = 0; n < 6; n++){
            int3 candidate = pixel;
            candidate.x += dx[n];
            candidate.y += dy[n];
            candidate.z += dz[n];
            
            if(!inside_int(candidate)){
                continue;
            }
            
            if(region[candidate.z * DATA_DIM*DATA_DIM + candidate.y*DATA_DIM + candidate.x]){
                continue;
            }
            
            if(similar(data, pixel, candidate)){
                push(stack, candidate);
                region[candidate.z * DATA_DIM*DATA_DIM + candidate.y*DATA_DIM + candidate.x] = 1;
            }
        }
    }
    
    return region;
}

unsigned char* grow_region_gpu(unsigned char* data){
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_kernel kernel;
    cl_int err;
    char *source;
    int i;
    
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);
    
    printPlatformInfo(platform);

    queue = clCreateCommandQueue(context, device, 0, &err);
    kernel = buildKernel("region.cl", "region", NULL, context, device);

    //Host variables
    unsigned char* host_region = (unsigned char*)calloc(sizeof(unsigned char), DATA_SIZE);
    int            host_unfinished;

    cl_mem device_region     = clCreateBuffer(context, CL_MEM_READ_WRITE, DATA_SIZE * sizeof(cl_uchar) ,NULL,&err);
    cl_mem device_data       = clCreateBuffer(context, CL_MEM_READ_ONLY, DATA_SIZE * sizeof(cl_uchar), NULL,&err);
    cl_mem device_unfinished = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL,&err);
    clError("Error allocating memory", err);

    //plant seed
    int3 seed = {.x=50, .y=300, .z=300};
    host_region[index(seed.z, seed.y, seed.x)] = 2;

    //Copy data to the device
    clEnqueueWriteBuffer(queue, device_data  , CL_FALSE, 0, DATA_SIZE * sizeof(cl_uchar), data       , 0, NULL, NULL);
    clEnqueueWriteBuffer(queue, device_region, CL_FALSE, 0, DATA_SIZE * sizeof(cl_uchar), host_region, 0, NULL, NULL);

    //Calculate block and grid sizes
    size_t global[] = { 512, 512, 512 }; 
    size_t local[] = { 8, 8, 8 };


    //Run kernel untill completion
    do{
        host_unfinished = 0;
        clEnqueueWriteBuffer(queue, device_unfinished, CL_FALSE, 0, sizeof(cl_int), &host_unfinished  , 0, NULL, NULL);
        clFinish(queue);

        err = clSetKernelArg(kernel, 0, sizeof(device_data), (void*)&device_data);
        err = clSetKernelArg(kernel, 1, sizeof(device_region), (void*)&device_region);
        err = clSetKernelArg(kernel, 2, sizeof(device_unfinished), (void*)&device_unfinished);
        clError("Error setting arguments", err);

        //Run the kernel
        clEnqueueNDRangeKernel(queue, kernel, 3, NULL, &global, &local, 0, NULL, NULL);
        clFinish(queue);
        clError("Error running kernel", err);

        err = clEnqueueReadBuffer(queue, device_unfinished, CL_TRUE, 0, sizeof(cl_int), &host_unfinished, 0, NULL, NULL);
        clFinish(queue);
        clError("Error reading buffer 1", err);

    }while(host_unfinished);

    //Copy result to host
    err = clEnqueueReadBuffer(queue, device_region, CL_TRUE, 0, DATA_SIZE * sizeof(cl_uchar), host_region, 0, NULL, NULL);
    clFinish(queue);
    clError("Error reading buffer 2", err);

    return host_region;
}

unsigned char* raycast_gpu(unsigned char* data, unsigned char* region){
    cl_platform_id platform;
    cl_device_id device;
    cl_context context;
    cl_command_queue queue;
    cl_kernel kernel;
    cl_int err;
    char *source;
    int i;
    
    clGetPlatformIDs(1, &platform, NULL);
    clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
    context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);

    printPlatformInfo(platform);
    printDeviceInfo(device);

    queue = clCreateCommandQueue(context, device, 0, &err);
    kernel = buildKernel("raycast.cl", "raycast", NULL, context, device);


    cl_mem device_region = clCreateBuffer(context, CL_MEM_READ_ONLY, DATA_SIZE * sizeof(cl_uchar) ,NULL,&err);
    cl_mem device_data   = clCreateBuffer(context, CL_MEM_READ_ONLY, DATA_SIZE * sizeof(cl_uchar), NULL,&err);
    cl_mem device_image  = clCreateBuffer(context, CL_MEM_READ_WRITE, IMAGE_SIZE * sizeof(cl_uchar),NULL,&err);
    clError("Error allocating memory", err);
    //Copy data to the device
    clEnqueueWriteBuffer(queue, device_data  , CL_FALSE, 0, DATA_SIZE * sizeof(cl_uchar), data  , 0, NULL, NULL);
    clEnqueueWriteBuffer(queue, device_region, CL_FALSE, 0, DATA_SIZE * sizeof(cl_uchar), region, 0, NULL, NULL);

    int grid_size = IMAGE_DIM;
    int block_size = IMAGE_DIM;


    //Set up kernel arguments
    err = clSetKernelArg(kernel, 0, sizeof(device_data), (void*)&device_data);
    err = clSetKernelArg(kernel, 1, sizeof(device_region), (void*)&device_region);
    err = clSetKernelArg(kernel, 2, sizeof(device_image), (void*)&device_image);
    clError("Error setting arguments", err);

    //Run the kernel
    const size_t globalws[2] = {IMAGE_DIM, IMAGE_DIM};
    const size_t localws[2] = {8, 8};
    clEnqueueNDRangeKernel(queue, kernel, 2, NULL, &globalws, &localws, 0, NULL, NULL);
    
    clFinish(queue);


    //Allocate memory for the result
    unsigned char* host_image = (unsigned char*)malloc(IMAGE_SIZE_BYTES);

    //Copy result from device
    err = clEnqueueReadBuffer(queue, device_image, CL_TRUE, 0, IMAGE_SIZE * sizeof(cl_uchar), host_image, 0, NULL, NULL);
    clFinish(queue);


    //Free device memory
    return host_image;
}


int main(int argc, char** argv){
    
    unsigned char* data = create_data();
    
    unsigned char* region = grow_region_gpu(data);
    
    unsigned char* image = raycast_gpu(data, region);
    
    write_bmp(image, IMAGE_DIM, IMAGE_DIM);
}