void cudaCloverField::createTexObject(cudaTextureObject_t &tex, cudaTextureObject_t &texNorm,
void *field, void *norm) {
if (order == QUDA_FLOAT2_CLOVER_ORDER || order == QUDA_FLOAT4_CLOVER_ORDER) {
// create the texture for the field components
cudaChannelFormatDesc desc;
memset(&desc, 0, sizeof(cudaChannelFormatDesc));
if (precision == QUDA_SINGLE_PRECISION) desc.f = cudaChannelFormatKindFloat;
else desc.f = cudaChannelFormatKindSigned; // half is short, double is int2
// always four components regardless of precision
desc.x = (precision == QUDA_DOUBLE_PRECISION) ? 8*sizeof(int) : 8*precision;
desc.y = (precision == QUDA_DOUBLE_PRECISION) ? 8*sizeof(int) : 8*precision;
desc.z = (precision == QUDA_DOUBLE_PRECISION) ? 8*sizeof(int) : 8*precision;
desc.w = (precision == QUDA_DOUBLE_PRECISION) ? 8*sizeof(int) : 8*precision;
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = field;
resDesc.res.linear.desc = desc;
resDesc.res.linear.sizeInBytes = bytes/2;
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
if (precision == QUDA_HALF_PRECISION) texDesc.readMode = cudaReadModeNormalizedFloat;
else texDesc.readMode = cudaReadModeElementType;
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);
checkCudaError();
// create the texture for the norm components
if (precision == QUDA_HALF_PRECISION) {
cudaChannelFormatDesc desc;
memset(&desc, 0, sizeof(cudaChannelFormatDesc));
desc.f = cudaChannelFormatKindFloat;
desc.x = 8*QUDA_SINGLE_PRECISION; desc.y = 0; desc.z = 0; desc.w = 0;
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeLinear;
resDesc.res.linear.devPtr = norm;
resDesc.res.linear.desc = desc;
resDesc.res.linear.sizeInBytes = norm_bytes/2;
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.readMode = cudaReadModeElementType;
cudaCreateTextureObject(&texNorm, &resDesc, &texDesc, NULL);
checkCudaError();
}
}
}
__host__ explicit Texture(const GlobPtrSz<T>& mat,
bool normalizedCoords = false,
cudaTextureFilterMode filterMode = cudaFilterModePoint,
cudaTextureAddressMode addressMode = cudaAddressModeClamp)
{
CV_Assert( deviceSupports(FEATURE_SET_COMPUTE_30) );
rows = mat.rows;
cols = mat.cols;
cudaResourceDesc texRes;
std::memset(&texRes, 0, sizeof(texRes));
texRes.resType = cudaResourceTypePitch2D;
texRes.res.pitch2D.devPtr = mat.data;
texRes.res.pitch2D.height = mat.rows;
texRes.res.pitch2D.width = mat.cols;
texRes.res.pitch2D.pitchInBytes = mat.step;
texRes.res.pitch2D.desc = cudaCreateChannelDesc<T>();
cudaTextureDesc texDescr;
std::memset(&texDescr, 0, sizeof(texDescr));
texDescr.addressMode[0] = addressMode;
texDescr.addressMode[1] = addressMode;
texDescr.addressMode[2] = addressMode;
texDescr.filterMode = filterMode;
texDescr.readMode = cudaReadModeElementType;
texDescr.normalizedCoords = normalizedCoords;
CV_CUDEV_SAFE_CALL( cudaCreateTextureObject(&this->texObj, &texRes, &texDescr, 0) );
}
void TransferFunction::onColorTFChanged()
{
//std::cout<<"Color changed"<<std::endl;
if(compositeTex)
{
CudaSafeCall(cudaDestroyTextureObject(compositeTex));
compositeTex = 0;
}
colorTF->GetTable(0.0, 1.0, TABLE_SIZE, colorTable);
size_t j = 0, k = 0;
for(size_t i = 0; i < TABLE_SIZE; ++i)
{
compositeTable[j++] = colorTable[k++];
compositeTable[j++] = colorTable[k++];
compositeTable[j++] = colorTable[k++];
j++;
}
//CompositeTable();
CudaSafeCall(cudaMemcpyToArray(array, 0, 0, compositeTable, sizeof(float) * TABLE_SIZE * 4, cudaMemcpyHostToDevice));
CudaSafeCall(cudaCreateTextureObject(&compositeTex, &resourceDesc, &texDesc, NULL));
Changed();
}
TransferFunction::TransferFunction(vtkSmartPointer<vtkPiecewiseFunction> otf, vtkSmartPointer<vtkColorTransferFunction> ctf, QObject *parent) : QObject(parent)
{
opacityTF = otf;
colorTF = ctf;
this->otf = QSharedPointer<ctkTransferFunction>(new ctkVTKPiecewiseFunction(opacityTF));
this->ctf = QSharedPointer<ctkTransferFunction>(new ctkVTKColorTransferFunction(colorTF));
connect(this->otf.data(), SIGNAL(changed()), this, SLOT(onOpacityTFChanged()));
connect(this->ctf.data(), SIGNAL(changed()), this, SLOT(onColorTFChanged()));
compositeTex = 0;
// initialize each table
opacityTF->GetTable(0.0, 1.0, TABLE_SIZE, opacityTable);
colorTF->GetTable(0.0, 1.0, TABLE_SIZE, colorTable);
CompositeTable();
channelDesc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
CudaSafeCall(cudaMallocArray(&array, &channelDesc, TABLE_SIZE));
CudaSafeCall(cudaMemcpyToArray(array, 0, 0, compositeTable, sizeof(float) * TABLE_SIZE * 4, cudaMemcpyHostToDevice));
memset(&resourceDesc, 0, sizeof(resourceDesc));
resourceDesc.resType = cudaResourceTypeArray;
resourceDesc.res.array.array = array;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.normalizedCoords = true;
texDesc.readMode = cudaReadModeElementType;
CudaSafeCall(cudaCreateTextureObject(&compositeTex, &resourceDesc, &texDesc, NULL));
}
cuda_texture::cuda_texture(
cuda_linear_buffer_device::const_ptr dev_smart_ptr,
int vector_size)
: tex(0)
, dev_smart_ptr(dev_smart_ptr)
{
struct cudaResourceDesc res_desc;
memset(&res_desc, 0, sizeof(res_desc));
res_desc.resType = cudaResourceTypeLinear;
res_desc.res.linear.devPtr = const_cast<void *>((const void *)(*dev_smart_ptr));
switch (vector_size)
{
case 1:
res_desc.res.linear.desc = cudaCreateChannelDesc<float>();
break;
case 2:
res_desc.res.linear.desc = cudaCreateChannelDesc<float2>();
break;
case 4:
res_desc.res.linear.desc = cudaCreateChannelDesc<float4>();
break;
default:
throw neural_network_exception((boost::format("Invalid vetor_size %1% for cuda_texture") % vector_size).str());
}
res_desc.res.linear.sizeInBytes = dev_smart_ptr->get_size();
struct cudaTextureDesc tex_desc;
memset(&tex_desc, 0, sizeof(tex_desc));
tex_desc.addressMode[0] = cudaAddressModeBorder;
tex_desc.readMode = cudaReadModeElementType;
tex_desc.normalizedCoords = 0;
cuda_safe_call(cudaCreateTextureObject(&tex, &res_desc, &tex_desc, 0));
}
::cudaTextureObject_t
SharedAllocationRecord< Kokkos::CudaSpace , void >::
attach_texture_object( const unsigned sizeof_alias
, void * const alloc_ptr
, size_t const alloc_size )
{
// Only valid for 300 <= __CUDA_ARCH__
// otherwise return zero.
::cudaTextureObject_t tex_obj ;
struct cudaResourceDesc resDesc ;
struct cudaTextureDesc texDesc ;
memset( & resDesc , 0 , sizeof(resDesc) );
memset( & texDesc , 0 , sizeof(texDesc) );
resDesc.resType = cudaResourceTypeLinear ;
resDesc.res.linear.desc = ( sizeof_alias == 4 ? cudaCreateChannelDesc< int >() :
( sizeof_alias == 8 ? cudaCreateChannelDesc< ::int2 >() :
/* sizeof_alias == 16 */ cudaCreateChannelDesc< ::int4 >() ) );
resDesc.res.linear.sizeInBytes = alloc_size ;
resDesc.res.linear.devPtr = alloc_ptr ;
CUDA_SAFE_CALL( cudaCreateTextureObject( & tex_obj , & resDesc, & texDesc, NULL ) );
return tex_obj ;
}
void SingleParticle2dx::Methods::CUDAProjectionMethod::prepareForProjections(SingleParticle2dx::DataStructures::ParticleContainer& cont)
{
cudaSetDevice(getMyGPU());
cudaStreamCreate(&m_stream);
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
cudaExtent VS = make_cudaExtent(m_size, m_size, m_size);
if( m_alloc_done == false )
{
cudaMalloc3DArray(&m_cuArray, &channelDesc, VS);
}
SingleParticle2dx::real_array3d_type real_data( boost::extents[m_size][m_size][m_size] );
m_context->getRealSpaceData(real_data);
unsigned int size = m_size*m_size*m_size*sizeof(float);
if( m_alloc_done == false )
{
res_data_h = (float*)malloc(m_size*m_size*sizeof(float));
cudaMalloc((void**)&res_data_d, m_size*m_size*sizeof(float));
m_alloc_done = true;
}
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = make_cudaPitchedPtr((void*)real_data.origin(), VS.width*sizeof(float), VS.width, VS.height);
copyParams.dstArray = m_cuArray;
copyParams.extent = VS;
copyParams.kind = cudaMemcpyHostToDevice;
// cudaMemcpy3D(©Params);
cudaMemcpy3DAsync(©Params, m_stream);
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = m_cuArray;
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.addressMode[2] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = 0;
if(m_alloc_done == true)
{
cudaDestroyTextureObject(m_texObj);
}
m_texObj = 0;
cudaCreateTextureObject(&m_texObj, &resDesc, &texDesc, NULL);
}
static void addImageToTextureUint (vector<Mat_<uint8_t> > &imgs, cudaTextureObject_t texs[])
{
for (unsigned int i=0; i<imgs.size(); i++)
{
int rows = imgs[i].rows;
int cols = imgs[i].cols;
// Create channel with uint8_t point type
cudaChannelFormatDesc channelDesc =
//cudaCreateChannelDesc (8,
//0,
//0,
//0,
//cudaChannelFormatKindUnsigned);
cudaCreateChannelDesc<char>();
// Allocate array with correct size and number of channels
cudaArray *cuArray;
checkCudaErrors(cudaMallocArray(&cuArray,
&channelDesc,
cols,
rows));
checkCudaErrors (cudaMemcpy2DToArray (cuArray,
0,
0,
imgs[i].ptr<uint8_t>(),
imgs[i].step[0],
cols*sizeof(uint8_t),
rows,
cudaMemcpyHostToDevice));
// Specify texture
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;
// Specify texture object parameters
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeWrap;
texDesc.addressMode[1] = cudaAddressModeWrap;
texDesc.filterMode = cudaFilterModePoint;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = 0;
// Create texture object
//cudaTextureObject_t &texObj = texs[i];
checkCudaErrors(cudaCreateTextureObject(&(texs[i]), &resDesc, &texDesc, NULL));
//texs[i] = texObj;
}
return;
}
Texture::Texture(const float *data, int numElements)
{
cudaResourceDesc desc;
memset(&desc, 0, sizeof(desc));
desc.resType = cudaResourceTypeLinear;
desc.res.linear.devPtr = const_cast<float *>(data);
desc.res.linear.desc = cudaCreateChannelDesc<float>();
desc.res.linear.sizeInBytes = sizeof(float)*numElements;
cudaTextureDesc tdesc;
memset(&tdesc, 0, sizeof(tdesc));
CUDA_CHECK(cudaCreateTextureObject(&_texture, &desc, &tdesc, NULL));
}
Texture::Texture(const float4 *data, int numElements)
{
cudaResourceDesc desc;
memset(&desc, 0, sizeof(desc));
desc.resType = cudaResourceTypeLinear;
desc.res.linear.devPtr = const_cast<float4 *>(data);
desc.res.linear.desc = cudaCreateChannelDesc(32, 32, 32, 32, cudaChannelFormatKindFloat);
desc.res.linear.sizeInBytes = sizeof(float4)*numElements;
cudaTextureDesc tdesc;
memset(&tdesc, 0, sizeof(tdesc));
tdesc.addressMode[0] = cudaAddressModeBorder;
tdesc.addressMode[1] = cudaAddressModeBorder;
tdesc.addressMode[2] = cudaAddressModeBorder;
CUDA_CHECK(cudaCreateTextureObject(&_texture, &desc, &tdesc, NULL));
}
cudaTextureObject_t create_environment_light_texture(const std::string& filename)
{
int w = 0, h = 0, n = 0;
float* data = stbi_loadf(filename.c_str(), &w, &h, &n, 0);
if(!data)
{
std::cerr<<"Unable to load environment map: "<<filename<<std::endl;
exit(0);
}
//create channel desc
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
//create cudaArray
cudaArray* array;
checkCudaErrors(cudaMallocArray(&array, &channelDesc, w, h));
if(n == 3)
{
uint32_t count = w * h;
std::vector<float4> ext_data;
ext_data.reserve(count);
for(auto i = 0; i < count; ++i)
ext_data.push_back(make_float4(data[i * 3], data[i * 3 + 1], data[i * 3 + 2], 0.f));
checkCudaErrors(cudaMemcpyToArray(array, 0, 0, ext_data.data(), sizeof(float4) * w * h, cudaMemcpyHostToDevice));
}
else
checkCudaErrors(cudaMemcpyToArray(array, 0, 0, data, sizeof(float4) * w * h, cudaMemcpyHostToDevice));
//create resource desc
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = array;
//create texture desc
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeWrap;
texDesc.addressMode[1] = cudaAddressModeWrap;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = true;
//create cudaTextureObject
cudaTextureObject_t tex;
checkCudaErrors(cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL));
return tex;
}
static void addImageToTextureFloatColor (vector<Mat > &imgs, cudaTextureObject_t texs[])
{
for (size_t i=0; i<imgs.size(); i++)
{
int rows = imgs[i].rows;
int cols = imgs[i].cols;
// Create channel with floating point type
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float4>();
// Allocate array with correct size and number of channels
cudaArray *cuArray;
checkCudaErrors(cudaMallocArray(&cuArray,
&channelDesc,
cols,
rows));
checkCudaErrors (cudaMemcpy2DToArray (cuArray,
0,
0,
imgs[i].ptr<float>(),
imgs[i].step[0],
cols*sizeof(float)*4,
rows,
cudaMemcpyHostToDevice));
// Specify texture
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;
// Specify texture object parameters
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeWrap;
texDesc.addressMode[1] = cudaAddressModeWrap;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeElementType;
texDesc.normalizedCoords = 0;
// Create texture object
checkCudaErrors(cudaCreateTextureObject(&(texs[i]), &resDesc, &texDesc, NULL));
}
return;
}
void costVol_chamo::updataCV(Mat refImg, Mat projMat, float weightPerImg){
cudaArray* cuArray;
cudaTextureObject_t texObj;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
cudaSafeCall(cudaMallocArray(&cuArray, &channelDesc, width, height));
cudaMemcpyToArray(cuArray, 0, 0, refImg.data, width*height*sizeof(float), cudaMemcpyHostToDevice);
cudaSafeCall(cudaGetLastError());
struct cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));
resDesc.resType = cudaResourceTypeArray;
resDesc.res.array.array = cuArray;
struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModeLinear;
texDesc.readMode = cudaReadModeNormalizedFloat;
texDesc.normalizedCoords = 0;
cudaSafeCall(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL));
Mat finalTran = projMat*baseImgProjMat.inv();
cvInput input;
input.baseImg = (float *)baseImg.data;
input.cvData = (float*)cvData.data;
input.nearZ = nearZ;
input.farZ = farZ;
input.height = height;
input.width = width;
input.lowInd = (float*)lowInd.data;
input.lowValue = (float*)bottonVal.data;
for (int i = 0; i < 12; i++){
input.transMat[i] = finalTran.at<float>(i);
}
input.refImg = texObj;
input.zStep = (nearZ - farZ) / layers;
input.stepCount = layers;
updataCount++;
input.weightPerImg = 1.0 / updataCount;
updataCVCaller(input);
}
::cudaTextureObject_t
SharedAllocationRecord< Kokkos::CudaSpace , void >::
attach_texture_object( const unsigned sizeof_alias
, void * const alloc_ptr
, size_t const alloc_size )
{
enum { TEXTURE_BOUND_1D = 1u << 27 };
if ( ( alloc_ptr == 0 ) || ( sizeof_alias * TEXTURE_BOUND_1D <= alloc_size ) ) {
std::ostringstream msg ;
msg << "Kokkos::CudaSpace ERROR: Cannot attach texture object to"
<< " alloc_ptr(" << alloc_ptr << ")"
<< " alloc_size(" << alloc_size << ")"
<< " max_size(" << ( sizeof_alias * TEXTURE_BOUND_1D ) << ")" ;
std::cerr << msg.str() << std::endl ;
std::cerr.flush();
Kokkos::Impl::throw_runtime_exception( msg.str() );
}
::cudaTextureObject_t tex_obj ;
struct cudaResourceDesc resDesc ;
struct cudaTextureDesc texDesc ;
memset( & resDesc , 0 , sizeof(resDesc) );
memset( & texDesc , 0 , sizeof(texDesc) );
resDesc.resType = cudaResourceTypeLinear ;
resDesc.res.linear.desc = ( sizeof_alias == 4 ? cudaCreateChannelDesc< int >() :
( sizeof_alias == 8 ? cudaCreateChannelDesc< ::int2 >() :
/* sizeof_alias == 16 */ cudaCreateChannelDesc< ::int4 >() ) );
resDesc.res.linear.sizeInBytes = alloc_size ;
resDesc.res.linear.devPtr = alloc_ptr ;
CUDA_SAFE_CALL( cudaCreateTextureObject( & tex_obj , & resDesc, & texDesc, NULL ) );
return tex_obj ;
}
