//! @brief Copy the ND-array device array to host array.
//! You must allocate the array with the appropriate size.
__host__
inline void copy_to_host(T *host_data) const
{
if (N == 2)
{
SHAKTI_SAFE_CUDA_CALL(cudaMemcpy2D(
host_data, _sizes[0] * sizeof(T), _data, _pitch, _sizes[0] * sizeof(T), _sizes[1],
cudaMemcpyDeviceToHost));
}
else if (N == 3)
{
cudaMemcpy3DParms params = { 0 };
params.srcPtr = make_cudaPitchedPtr(reinterpret_cast<void *>(_data),
_pitch, _sizes[0], _sizes[1]);
params.dstPtr = make_cudaPitchedPtr(host_data, _sizes[0] * sizeof(T),
_sizes[0], _sizes[1]);
params.extent = make_cudaExtent(_sizes[0]*sizeof(T), _sizes[1],
_sizes[2]);
params.kind = cudaMemcpyDeviceToHost;
SHAKTI_SAFE_CUDA_CALL(cudaMemcpy3D(¶ms));
}
else
SHAKTI_SAFE_CUDA_CALL(cudaMemcpy(host_data, _data, sizeof(T) * size(),
cudaMemcpyDeviceToHost));
}
bool ControlCubeCache::_readElement(NodeLinkedList<index_node_t> * element)
{
#ifndef NDEBUG
if ((int)element->element > _maxNumCubes)
{
std::cerr<<"Control Cube CPU Cache, try to write outside reserved memory"<<std::endl;
throw;
}
#endif
index_node_t idCube = element->id;
float * cube = (_memory + element->element*_sizeElement);
if (!checkCubeInside(element->id))
{
if (cudaSuccess != cudaMemset((void*)cube, 0, _sizeElement*sizeof(float)))
{
std::cout<<"---> "<<idCube<<" "<<_minValue<<" "<<_maxValue<<std::endl;
LBERROR<<"Control Cube Cache: error copying to a device: "<<cudaGetErrorString(cudaGetLastError()) <<" "<<cube<<" "<<_sizeElement<<std::endl;
throw;
}
return true;
}
index_node_t idCubeCPU = idCube >> 3*(_levelCube - _cpuCache->getCubeLevel());
float * pCube = _cpuCache->getAndBlockElement(idCubeCPU);
if (pCube != 0)
{
vmml::vector<3, int> coord = getMinBoxIndex2(idCube, _levelCube, _nLevels);
vmml::vector<3, int> coordC = getMinBoxIndex2(idCubeCPU, _cpuCache->getCubeLevel(), _nLevels);
coord -= coordC;
vmml::vector<3, int> realDimCPU = _cpuCache->getRealCubeDim();
cudaMemcpy3DParms myParms = {0};
myParms.srcPtr = make_cudaPitchedPtr((void*)pCube, realDimCPU.z()*sizeof(float), realDimCPU.x(), realDimCPU.y());
//myParms.dstPtr = make_cudaPitchedPtr((void*)cube, _realcubeDim.z()*sizeof(float), _realcubeDim.x(), _realcubeDim.y());
myParms.dstPtr = make_cudaPitchedPtr((void*)cube, _dimCube*sizeof(float), _dimCube, _dimCube);
myParms.extent = make_cudaExtent(_dimCube*sizeof(float), _dimCube, _dimCube);
myParms.dstPos = make_cudaPos(0,0,0);
myParms.srcPos = make_cudaPos(coord.z()*sizeof(float), coord.y(), coord.x());
myParms.kind = cudaMemcpyHostToDevice;
if (cudaSuccess != cudaMemcpy3DAsync(&myParms, _stream) || cudaSuccess != cudaStreamSynchronize(_stream))
{
std::cout<<"---> "<<idCube<<" "<<_minValue<<" "<<_maxValue<<std::endl;
LBERROR<<"Control Cube Cache: error copying to a device: "<<cudaGetErrorString(cudaGetLastError()) <<" "<<cube<<" "<<pCube<<" "<<_sizeElement<<std::endl;
throw;
}
_cpuCache->unlockElement(idCubeCPU);
return true;
}
else
{
return false;
}
}
void cpyD2H() {
// D2H“]‘—
cudaMemcpy3DParms parms = { 0 };
parms.srcPos = make_cudaPos(0, 0, 0);
parms.srcPtr = make_cudaPitchedPtr(vdm.dPtr, sizeof(TYPE) * buf.pitch.x, buf.pitch.x, buf.pitch.y);
parms.dstPos = make_cudaPos(sizeof(TYPE) * buf.lower.x, buf.lower.y, buf.lower.z);
parms.dstPtr = make_cudaPitchedPtr(vdm.hPtr, sizeof(TYPE) * vdm.pitch.x, vdm.pitch.x, vdm.pitch.y);
parms.extent = make_cudaExtent(sizeof(TYPE) * buf.pitch.x, buf.pitch.y, buf.pitch.z);
parms.kind = cudaMemcpyDeviceToHost;
cudaMemcpy3D(&parms);
}
void MemManager::cpyD2H(Task task) {
cudaStream_t stream = task.stream;
for (auto vdm : vdms) {
VDMBuffer buf = vdm.bufs[task.n];
// D2H転送
cudaMemcpy3DParms parms = { 0 };
parms.srcPos = make_cudaPos(0, 0, 0);
parms.srcPtr = make_cudaPitchedPtr(vdm.dPtr, sizeof(TYPE) * buf.pitchx(), buf.pitchx(), buf.pitchy());
parms.dstPos = make_cudaPos(sizeof(TYPE) * buf.lower.x, buf.lower.y, buf.lower.z);
parms.dstPtr = make_cudaPitchedPtr(vdm.hPtr, sizeof(TYPE) * vdm.pitch.x, vdm.pitch.x, vdm.pitch.y);
parms.extent = make_cudaExtent(sizeof(TYPE) * buf.pitchx(), buf.pitchy(), buf.pitchz());
parms.kind = cudaMemcpyDeviceToHost;
cudaMemcpy3DAsync(&parms, stream);
}
}
//! Hidden copy constructor
VolumeGPU( const VolumeGPU& src ) : dims(make_uint3(0,0,0)),
d_data(make_cudaPitchedPtr(NULL,0,0,0)) {
std::cerr << __FUNCTION__
<< ": Please don't use copy constructor"
<< std::endl;
exit( EXIT_FAILURE );
}
void MemManager::cpyH2D(Task task) {
cudaStream_t stream = task.stream;
for (auto vdm : vdms) {
VDMBuffer buf = vdm.bufs[task.n];
// H2D転送
cudaMemcpy3DParms parms = { 0 };
parms.dstPos = make_cudaPos(0, 0, 0);
parms.dstPtr = make_cudaPitchedPtr(vdm.dPtr, sizeof(TYPE) * buf.pitchx(), buf.pitchx(), buf.pitchy());
parms.srcPos = make_cudaPos(sizeof(TYPE) * buf.lower.x, buf.lower.y, buf.lower.z);
parms.srcPtr = make_cudaPitchedPtr(vdm.hPtr, sizeof(TYPE) * vdm.pitch.x, vdm.pitch.x, vdm.pitch.y);
parms.extent = make_cudaExtent(sizeof(TYPE) * buf.pitchx(), buf.pitchy(), buf.pitchz());
parms.kind = cudaMemcpyHostToDevice;
cudaMemcpy3DAsync(&parms, stream);
}
cudaMemcpyAsync(confsetDevPtr, confsetHosPtrs[task.n], sizeof(MKConf)*nVdm, cudaMemcpyHostToDevice, stream);
}
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);
}
void CudaImagePyramidHost::copyFromHost(int width, int height, const void* source, int layer)
{
assert(isInitialized());
assert(_textureType == cudaTextureType2DLayered);
cudaMemcpy3DParms myParms = {0};
myParms.srcPtr = make_cudaPitchedPtr((void*)source,width*_typeSize,width,height);
myParms.srcPos = make_cudaPos(0,0,0);
myParms.dstArray = _storage;
myParms.dstPos = make_cudaPos(0,0,layer);
myParms.extent = make_cudaExtent(width,height,1);
myParms.kind = cudaMemcpyHostToDevice;
cudaMemcpy3D(&myParms);
checkCUDAError("Memcpy error", _name);
}
void Buffer3D::setData( void const* srcData, dp::math::Vec3ui const& srcOffset, dp::math::Vec3ui const& srcStride, dp::math::Vec3ui const& srcExtent, dp::math::Vec3ui const& dstOffset )
{
#if !defined(NDEBUG)
for ( int i=0 ; i<3 ; i++ )
{
DP_ASSERT( srcOffset[i] + ( srcExtent[i] ? srcExtent[i] : m_extent[i] ) <= m_extent[i] );
DP_ASSERT( dstOffset[i] + ( srcExtent[i] ? srcExtent[i] : m_extent[i] ) <= m_extent[i] );
}
#endif
cudaMemcpy3DParms parms = { 0 };
parms.srcPos = make_cudaPos( m_elementSize * srcOffset[0], srcOffset[1], srcOffset[2] );
parms.srcPtr = make_cudaPitchedPtr( const_cast<void *>(srcData), m_elementSize * ( srcStride[0] ? srcStride[0] : m_extent[0] ), srcStride[0] ? srcStride[0] : m_extent[0], srcStride[1] ? srcStride[1] : m_extent[1] );
parms.dstPos = make_cudaPos( m_elementSize * dstOffset[0], dstOffset[1], dstOffset[2] );
parms.dstPtr = m_pitchedPtr;
parms.extent = make_cudaExtent( m_elementSize * ( srcExtent[0] ? srcExtent[0] : m_extent[0] ), srcExtent[1] ? srcExtent[1] : m_extent[1], srcExtent[2] ? srcExtent[2] : m_extent[2] );
parms.kind = cudaMemcpyHostToDevice;
CUDA_VERIFY( cudaMemcpy3D( &parms ) );
}
const cudaPitchedPtr getCudaPitched() const
{
__startOperation(ITask::TASK_CUDA);
TYPE* dPointer;
cudaHostGetDevicePointer(&dPointer, pointer, 0);
/* on 1D memory we have no size for y, therefore we set y to 1 to
* get a valid cudaPitchedPtr
*/
int size_y=1;
if(DIM>DIM1)
size_y= this->data_space[1];
return make_cudaPitchedPtr(dPointer,
this->data_space.x() * sizeof (TYPE),
this->data_space.x(),
size_y
);
}
void Buffer3D::getData( dp::cuda::BufferHostSharedPtr const& dstBuffer, dp::math::Vec3ui const& dstOffset, dp::math::Vec3ui const& dstStride, dp::math::Vec3ui const& dstExtent
, dp::math::Vec3ui const& srcOffset, dp::cuda::StreamSharedPtr const& stream )
{
#if !defined(NDEBUG)
for ( int i=0 ; i<3 ; i++ )
{
DP_ASSERT( dstOffset[i] + ( dstExtent[i] ? dstExtent[i] : m_extent[i] ) <= m_extent[i] );
DP_ASSERT( srcOffset[i] + ( dstExtent[i] ? dstExtent[i] : m_extent[i] ) <= m_extent[i] );
}
#endif
cudaMemcpy3DParms parms = { 0 };
parms.srcPos = make_cudaPos( m_elementSize * srcOffset[0], srcOffset[1], srcOffset[2] );
parms.srcPtr = m_pitchedPtr;
parms.dstPos = make_cudaPos( m_elementSize * dstOffset[0], dstOffset[1], dstOffset[2] );
parms.dstPtr = make_cudaPitchedPtr( dstBuffer->getPointer<void>(), m_elementSize * ( dstStride[0] ? dstStride[0] : m_extent[0] ), dstStride[0] ? dstStride[0] : m_extent[0], dstStride[1] ? dstStride[1] : m_extent[1] );
parms.extent = make_cudaExtent( m_elementSize * ( dstExtent[0] ? dstExtent[0] : m_extent[0] ), dstExtent[1] ? dstExtent[1] : m_extent[1], dstExtent[2] ? dstExtent[2] : m_extent[2] );
parms.kind = cudaMemcpyDeviceToHost;
CUDA_VERIFY( cudaMemcpy3DAsync( &parms ) );
}
void VolSkin::init( int width, int height, TetMesh *tm )
{
this->width = width;
this->height = height;
tetMesh = tm;
// TEMP initialize volume data
cudaExtent volumeSize = make_cudaExtent(128, 128, 128);
//cudaExtent volumeSize = make_cudaExtent(256, 256, 256);
// generate raw volume data
float *h_densityData = (float*)malloc( sizeof(float)*volumeSize.width*volumeSize.height*volumeSize.depth );
math::PerlinNoise pn;
pn.setDepth( 4 );
pn.setFrequency(3.0f);
//pn.setInflection(true);
for( int k=0;k<volumeSize.depth;++k )
for( int j=0;j<volumeSize.height;++j )
for( int i=0;i<volumeSize.width;++i )
{
int index = k*volumeSize.width*volumeSize.height + j*volumeSize.width + i;
math::Vec3f uvw( (float)(i)/(float)(volumeSize.width),
(float)(j)/(float)(volumeSize.height),
(float)(k)/(float)(volumeSize.depth));
float t = (float)(j)/(float)(volumeSize.height);
//h_densityData[index] = 0.5f;
//h_densityData[index] = (1.0f-t)*1.0f;
h_densityData[index] = std::max( 0.0f, pn.perlinNoise_3D( uvw.x, uvw.y*2.0, uvw.z ) )*1.0f; // cylinder
//h_densityData[index] = std::max( 0.0f, pn.perlinNoise_3D( uvw.x*2.0f, uvw.y*2.0f, uvw.z*2.0f ))*1.0f; // tetraeder
//h_densityData[index] = (uvw.getLength() < 0.2f ? 1.0f : 0.0f)*2.0f;
}
// create 3D array
d_densityArray = 0;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaMalloc3DArray(&d_densityArray, &channelDesc, volumeSize);
// copy data to 3D array
cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr = make_cudaPitchedPtr((void*)h_densityData, volumeSize.width*sizeof(float), volumeSize.width, volumeSize.height);
copyParams.dstArray = d_densityArray;
copyParams.extent = volumeSize;
copyParams.kind = cudaMemcpyHostToDevice;
cudaMemcpy3D(©Params);
// TMP
/*
h_debugVec.resize( 1000.0f );
d_debugVec = h_debugVec;
h_debugInfo.samples = convertToKernel(d_debugVec);
h_debugInfo.numSamples = 0;
cudaMemcpyToSymbol( d_debugInfo, &h_debugInfo, sizeof(DebugInfo), 0, cudaMemcpyHostToDevice );
*/
// setup lighting
m_light0.cam = base::CameraPtr( new base::Camera() );
m_light0.cam->m_aspectRatio = 1.0;
//m_light0.cam->m_transform = math::createLookAtMatrix( math::Vec3f( -2.0f, -2.0f, 2.0f ), math::Vec3f( 0.0f, 0.0f, 0.0f ), math::Vec3f( 0.0f, 1.0f, 0.0f ), false );
//m_light0.cam->m_transform = math::Matrix44f::TranslationMatrix( 0.3f, 0.15f, 2.0f );
//m_light0.cam->m_transform = math::Matrix44f::TranslationMatrix( -3.0f, 0.0f, 0.0f );
m_light0.cam->m_transform = math::createLookAtMatrix( math::Vec3f( 4.0f, 0.0f, 0.0f ), math::Vec3f( 0.0f, 0.0f, 0.0f ), math::Vec3f( 0.0f, 1.0f, 0.0f ), false );
m_light0.cam->update();
cudaMalloc( &m_light0.d_dctCoefficients, width*height*sizeof(float)*8 );// 8 floats /6 coefficients
// set defaults
setTotalCrossSection( 10.0f );
setAlbedo( 1.0f );
setAbsorptionColor( math::Vec3f(0.5f,0.5f, 0.5f) );
setScatteringColor(math::Vec3f(0.5f, 0.5f, 0.5f));
setLight(0, math::Vec3f(1.0f, 1.0f, 1.0f), 0.0f);
setTime( 0.0f );
setStepSize( 0.01f );
// get tetmesh onto gpu
gpuUploadTetMesh();
}
//! Default constructor
VecVolGPU( void ) : dims(make_uint3(0,0,0)),
d_x(make_cudaPitchedPtr(NULL,0,0,0)),
d_y(make_cudaPitchedPtr(NULL,0,0,0)),
d_z(make_cudaPitchedPtr(NULL,0,0,0)) {};
//! Default constructor
VolumeGPU( void ) : dims(make_uint3(0,0,0)),
d_data(make_cudaPitchedPtr(NULL,0,0,0)) {};
