//! @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)) {};