const textureReference* getTextureReference(const void* symbol) {
const textureReference* tex = nullptr;
if (Ctx->isCreated()) {
checkCudaError(cudaGetTextureReference(&tex, symbol));
}
return tex;
}
void CudaImagePyramidHost::unbindTexture()
{
const textureReference* constTexRefPtr=NULL;
cudaGetTextureReference(&constTexRefPtr, _texture_name);
checkCUDAError("Can't get tex ref for unbind TEXTURE_PYRAMID", _name);
cudaUnbindTexture(constTexRefPtr);
checkCUDAError("Unbind error", _name);
}
void CudaImagePyramidHost::bindTexture()
{
const textureReference* constTexRefPtr=NULL;
cudaGetTextureReference(&constTexRefPtr, _texture_name);
checkCUDAError("Can't get tex ref for bind TEXTURE_PYRAMID", _name);
cudaChannelFormatDesc formatDesc = constTexRefPtr->channelDesc;
cudaBindTextureToArray(constTexRefPtr, _storage, &formatDesc);
checkCUDAError("Bind error", _name);
}
template<class T> static inline void bindTexture(const char* name, const DevMem2D_<T>& img/*, bool normalized = false,
enum cudaTextureFilterMode filterMode = cudaFilterModePoint, enum cudaTextureAddressMode addrMode = cudaAddressModeClamp*/)
{
//!!!! const_cast is disabled!
//!!!! Please use constructor of 'class texture' instead.
//textureReference* tex;
//cudaSafeCall( cudaGetTextureReference((const textureReference**)&tex, name) );
//tex->normalized = normalized;
//tex->filterMode = filterMode;
//tex->addressMode[0] = addrMode;
//tex->addressMode[1] = addrMode;
const textureReference* tex;
cudaSafeCall( cudaGetTextureReference(&tex, name) );
cudaChannelFormatDesc desc = cudaCreateChannelDesc<T>();
cudaSafeCall( cudaBindTexture2D(0, tex, img.ptr(), &desc, img.cols, img.rows, img.step) );
}
void CudaImagePyramidHost::initialize(int width, int height, cudaTextureFilterMode filter_mode, int depth)
{
qDebug() << "pyramid host initializing with params: " << width << height << filter_mode << depth;
if (isInitialized() && width == _baseWidth && height == _baseHeight && filter_mode == _filterMode) {
return;
}
clear();
qDebug() << "Clear done.";
_baseWidth = width;
_baseHeight = height;
_filterMode = filter_mode;
_numLayers = depth;
// Get the texture and its channel descriptor to allocate the storage.
const textureReference* constTexRefPtr=NULL;
cudaGetTextureReference(&constTexRefPtr, _texture_name);
qDebug() << "Texture Ref got:" << _name;
if (constTexRefPtr == 0)
{
qDebug() << "constTexRefPtr==0";
}
checkCUDAError("Can't get tex ref for init TEXTURE_PYRAMID", _name);
cudaChannelFormatDesc formatDesc = constTexRefPtr->channelDesc;
if(_textureType == cudaTextureType2DLayered){
cudaDeviceProp prop;
qDebug() << "to get CUDA device prop";
cudaGetDeviceProperties(&prop,0);
qDebug() << "CUDA Device Prop got";
if(prop.maxTexture2DLayered[0] < _baseWidth || prop.maxTexture2DLayered[1] < _baseHeight || prop.maxTexture2DLayered[2] < _numLayers){
qDebug()<< "Max layered texture size:" << prop.maxTexture2DLayered[0] << " x " << prop.maxTexture2DLayered[1] << " x " << prop.maxTexture2DLayered[2];
assert(0);
}
cudaExtent extent = make_cudaExtent(_baseWidth, _baseHeight, _numLayers);
cudaMalloc3DArray(&_storage, &formatDesc, extent, cudaArrayLayered);
}else{
cudaMallocArray(&_storage, &formatDesc, _baseWidth, _baseHeight);
}
checkCUDAError("Failure to allocate", _name);
qDebug() << "allocate done";
// Set texture parameters.
// Evil hack to get around an apparent bug in the cuda api:
// cudaGetTextureReference only returns a const reference, and
// there is no way to set the parameters with a reference other
// than cast it to non-const.
textureReference* texRefPtr=NULL;
texRefPtr = const_cast<textureReference*>( constTexRefPtr );
texRefPtr->addressMode[0] = cudaAddressModeClamp;
texRefPtr->addressMode[1] = cudaAddressModeClamp;
texRefPtr->filterMode = filter_mode;
texRefPtr->normalized = false; // Use unnormalized (pixel) coordinates for addressing. This forbids texture mode wrap.
bindTexture();
qDebug() << "texture binded";
bool found = false;
for (size_t i = 0; i < _instances.size(); i++) {
if (_instances[i] == this)
found = true;
}
if (!found) {
qDebug() << "Not found";
_instances.push_back(this);
}
qDebug() << "paramid host initialized.";
}
cudaError_t WINAPI wine_cudaGetTextureReference( const struct textureReference **texref, const char *symbol ) {
WINE_TRACE("\n");
return cudaGetTextureReference( texref, symbol );
}
static inline void unbindTexture(const char *name)
{
const textureReference* tex;
cudaSafeCall( cudaGetTextureReference(&tex, name) );
cudaSafeCall( cudaUnbindTexture(tex) );
}
