#ifndef RTS_COLORMAP_H #define RTS_COLORMAP_H #include #include #include #include "rts/cuda/error.h" #define BREWER_CTRL_PTS 11 #ifdef __CUDACC__ texture cudaTexBrewer; static cudaArray* gpuBrewer; #endif namespace rts{ namespace colormap{ enum colormapType {cmBrewer, cmGrayscale}; static void buffer2image(unsigned char* buffer, std::string filename, unsigned int x_size, unsigned int y_size) { //create an image object QImage image(x_size, y_size, QImage::Format_RGB32); int i; unsigned char r, g, b; unsigned int x, y; for(y=0; y __global__ static void applyBrewer(T* gpuSource, unsigned char* gpuDest, unsigned int N, T minVal = 0, T maxVal = 1) { int i = blockIdx.x * blockDim.x + threadIdx.x; if(i >= N) return; //compute the normalized value on [minVal maxVal] float a = (gpuSource[i] - minVal) / (maxVal - minVal); //lookup the color float shift = 1.0/BREWER_CTRL_PTS; float4 color = tex1D(cudaTexBrewer, a+shift); gpuDest[i * 3 + 0] = 255 * color.x; gpuDest[i * 3 + 1] = 255 * color.y; gpuDest[i * 3 + 2] = 255 * color.z; } template __global__ static void applyGrayscale(T* gpuSource, unsigned char* gpuDest, unsigned int N, T minVal = 0, T maxVal = 1) { int i = blockIdx.x * blockDim.x + threadIdx.x; if(i >= N) return; //compute the normalized value on [minVal maxVal] float a = (gpuSource[i] - minVal) / (maxVal - minVal); gpuDest[i * 3 + 0] = 255 * a; gpuDest[i * 3 + 1] = 255 * a; gpuDest[i * 3 + 2] = 255 * a; } template static void gpu2gpu(T* gpuSource, unsigned char* gpuDest, unsigned int nVals, T minVal = 0, T maxVal = 1, colormapType cm = cmGrayscale, int blockDim = 128) { //This function converts a scalar field on the GPU to a color image on the GPU int gridDim = (nVals + blockDim - 1)/blockDim; if(cm == cmGrayscale) applyGrayscale<<>>(gpuSource, gpuDest, nVals, minVal, maxVal); else if(cm == cmBrewer) { initBrewer(); applyBrewer<<>>(gpuSource, gpuDest, nVals, minVal, maxVal); destroyBrewer(); } } template static void gpu2cpu(T* gpuSource, unsigned char* cpuDest, unsigned int nVals, T minVal, T maxVal, colormapType cm = cmGrayscale) { //this function converts a scalar field on the GPU to a color image on the CPU //first create the color image on the GPU //allocate GPU memory for the color image unsigned char* gpuDest; HANDLE_ERROR(cudaMalloc( (void**)&gpuDest, sizeof(unsigned char) * nVals * 3 )); //HANDLE_ERROR(cudaMemset(gpuSource, 0, sizeof(T) * nVals)); //create the image on the gpu gpu2gpu(gpuSource, gpuDest, nVals, minVal, maxVal, cm); //HANDLE_ERROR(cudaMemset(gpuDest, 0, sizeof(unsigned char) * nVals * 3)); //copy the image from the GPU to the CPU HANDLE_ERROR(cudaMemcpy(cpuDest, gpuDest, sizeof(unsigned char) * nVals * 3, cudaMemcpyDeviceToHost)); HANDLE_ERROR(cudaFree( gpuDest )); } template static void gpu2image(T* gpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, T valMin, T valMax, colormapType cm = cmGrayscale) { //allocate a color buffer unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size); //do the mapping gpu2cpu(gpuSource, cpuBuffer, x_size * y_size, valMin, valMax, cm); //copy the buffer to an image buffer2image(cpuBuffer, fileDest, x_size, y_size); free(cpuBuffer); } #endif template static void cpu2cpu(T* cpuSource, unsigned char* cpuDest, unsigned int nVals, T valMin, T valMax, colormapType cm = cmGrayscale) { int i; float a; float range = valMax - valMin; for(i = 0; i static void cpu2image(T* cpuSource, std::string fileDest, unsigned int x_size, unsigned int y_size, T valMin, T valMax, colormapType cm = cmGrayscale) { //allocate a color buffer unsigned char* cpuBuffer = (unsigned char*) malloc(sizeof(unsigned char) * 3 * x_size * y_size); //do the mapping cpu2cpu(cpuSource, cpuBuffer, x_size * y_size, valMin, valMax, cm); //copy the buffer to an image buffer2image(cpuBuffer, fileDest, x_size, y_size); free(cpuBuffer); } }} //end namespace colormap and rts #endif