#ifndef STIM_CUDA_TEXTURE_H #define STIM_CUDA_TEXTURE_H #include #include #include #include #include #include //#include #include #include #include #include #include #include ///A container for the texture based methods used by the spider class. namespace stim { namespace cuda { class cuda_texture { public: cudaArray* srcArray; cudaGraphicsResource_t resource; struct cudaResourceDesc resDesc; struct cudaTextureDesc texDesc; cudaTextureObject_t tObj; float *result; ///basic constructor that creates the texture with default parameters. cuda_texture() { memset(&texDesc, 0, sizeof(texDesc)); texDesc.addressMode[0] = cudaAddressModeWrap; texDesc.addressMode[1] = cudaAddressModeWrap; texDesc.filterMode = cudaFilterModePoint; texDesc.readMode = cudaReadModeElementType; texDesc.normalizedCoords = 0; } ///basic destructor ~cuda_texture() { UnmapCudaTexture(); if(result != NULL) cudaFree(result); } ///Enable the nromalized texture coordinates. ///@param bool, 1 for on, 0 for off void SetTextureCoordinates(bool val) { if(val) texDesc.normalizedCoords = 1; else texDesc.normalizedCoords = 0; } ///sets the dimension dim to used the mode at the borders of the texture. ///@param dim : 0-x, 1-y, 2-z ///@param mode: cudaAddressModeWrap = 0, /// cudaAddressModeClamp = 1, /// cudaAddressMNodeMirror = 2, /// cudaAddressModeBorder = 3, void SetAddressMode(int dim, int mode) { switch(mode) { case 0: texDesc.addressMode[dim] = cudaAddressModeWrap; break; case 1: texDesc.addressMode[dim] = cudaAddressModeClamp; break; case 2: texDesc.addressMode[dim] = cudaAddressModeMirror; break; case 3: texDesc.addressMode[dim] = cudaAddressModeBorder; break; default: break; } } //-------------------------------------------------------------------------// //-------------------------------CUDA_MAPPING------------------------------// //-------------------------------------------------------------------------// //Methods for creating the cuda texture. ///@param GLuint tex -- GLtexture (must be contained in a frame buffer object) /// that holds that data that will be handed to cuda. ///@param GLenum target -- either GL_TEXTURE_1D, GL_TEXTURE_2D or GL_TEXTURE_3D /// map work with other gl texture types but untested. ///Maps the gl texture in cuda memory, binds that data to a cuda array, and binds the cuda ///array to a cuda texture. void MapCudaTexture(GLuint tex, GLenum target) { HANDLE_ERROR( cudaGraphicsGLRegisterImage( &resource, tex, target, cudaGraphicsMapFlagsReadOnly // cudaGraphicsRegisterFlagsNone ) ); HANDLE_ERROR( cudaGraphicsMapResources(1, &resource) ); HANDLE_ERROR( cudaGraphicsSubResourceGetMappedArray(&srcArray, resource, 0, 0) ); memset(&resDesc, 0, sizeof(resDesc)); resDesc.resType = cudaResourceTypeArray; resDesc.res.array.array = srcArray; HANDLE_ERROR( cudaCreateTextureObject(&tObj, &resDesc, &texDesc, NULL) ); } ///Unmaps the gl texture, binds that data to a cuda array, and binds the cuda ///array to a cuda texture. void UnmapCudaTexture() { // HANDLE_ERROR( // cudaGraphicsUnmapResources(1, &resource) // ); // HANDLE_ERROR( // cudaGraphicsUnregisterResource(resource) // ); // HANDLE_ERROR( // cudaDestroyTextureObject(tObj) // ); // HANDLE_ERROR( // cudaFreeArray(srcArray) // ); } ///Allocate the auxiliary internal 1D float array void Alloc(int x) { cudaMalloc( (void**) &result, x*sizeof(float)); } //-------------------------------------------------------------------------// //------------------------------GET/SET METHODS----------------------------// //-------------------------------------------------------------------------// ///Returns the bound texture object. cudaTextureObject_t getTexture() { return tObj; } cudaArray* getArray() { return srcArray; } float* getAuxArray() { return result; } }; } } #endif