void CUDAResourceManager::deallocUSG(GPUUsg *usg)
{
CUDA_SAFE_CALL(cudaFree(usg->getElemList()));
CUDA_SAFE_CALL(cudaFree(usg->getTypeList()));
CUDA_SAFE_CALL(cudaFree(usg->getConnList()));
CUDA_SAFE_CALL(cudaFree(usg->getVertices()));
}
void
LiGL2D::setVbo(int spaceVect)
{
GLuint oldVbo = 0;
GLuint newVbo = 0;
if(vbo != 0){
oldVbo = vbo;
vbo = 0;
}
if(iw != 0 && ih !=0){
GLint bsize;
// create buffer object
unsigned int size = ((int)iw/(spaceVect+1))*((int)ih/(spaceVect+1)) * 6 * sizeof(float2);
glGenBuffers( 1, &newVbo);
glBindBuffer( GL_ARRAY_BUFFER, newVbo);
// initialize buffer object
glBufferData( GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glGetBufferParameterivARB(GL_ARRAY_BUFFER_ARB, GL_BUFFER_SIZE_ARB, &bsize);
glBindBuffer( GL_ARRAY_BUFFER, 0);
// register buffer object with CUDA
CUDA_SAFE_CALL(cudaGLRegisterBufferObject(newVbo));
sVbo = ((int)iw/(spaceVect+1))*((int)ih/(spaceVect+1))*6;
vbo = newVbo;
emit sendVbo(vbo);
}
if(oldVbo != 0){
CUDA_SAFE_CALL(cudaGLUnregisterBufferObject(oldVbo));
glDeleteBuffers(1, &oldVbo);
}
}
void CpuSNN::printSimSummary(FILE *fp)
{
DBG(2, fpLog, AT, "printSimSummary()");
float etime;
if(currentMode == GPU_MODE) {
stopGPUTiming();
etime = gpuExecutionTime;
CUDA_SAFE_CALL( cudaMemcpyFromSymbol( &spikeCountD2, "secD2fireCnt", sizeof(int), 0, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL( cudaMemcpyFromSymbol( &spikeCountD1, "secD1fireCnt", sizeof(int), 0, cudaMemcpyDeviceToHost));
spikeCountAll1sec = spikeCountD1 + spikeCountD2;
CUDA_SAFE_CALL( cudaMemcpyFromSymbol( &spikeCountD2, "spikeCountD2", sizeof(int), 0, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL( cudaMemcpyFromSymbol( &spikeCountD1, "spikeCountD1", sizeof(int), 0, cudaMemcpyDeviceToHost));
spikeCountAll = spikeCountD1 + spikeCountD2;
}
else {
stopCPUTiming();
etime = cpuExecutionTime;
}
fprintf(fp, "\n*** Network configuration dumped in %s.dot file...\n\
Use graphViz to see the network connectivity...\n\n", networkName.c_str());
fprintf(fp, "*********** %s Simulation Summary **********\n", (currentMode == GPU_MODE)?("GPU"):"CPU");
fprintf(fp, "Network Parameters: \n\tN = %d (numNExcReg:numNInhReg=%2.1f:%2.1f), numPostSynapses = %d, D = %d\n", numN, 100.0*numNExcReg/numN, 100.0*numNInhReg/numN, numPostSynapses, D);
fprintf(fp, "Random Seed: %d\n", randSeed);
fprintf(fp, "Timing: \n\tModel Simulation Time = %lld sec \n\tActual Execution Time = %4.2f sec\n", (unsigned long long)simTimeSec, etime/1000.0);
fprintf(fp, "Average Firing Rate \n\t2+ms delay = %3.3f Hz \n\t1ms delay = %3.3f Hz \n\tOverall = %3.3f Hz\n",
spikeCountD2/(1.0*simTimeSec*numNExcReg), spikeCountD1/(1.0*simTimeSec*numNInhReg), spikeCountAll/(1.0*simTimeSec*numN));
fprintf(fp, "Overall Firing Count: \n\t2+ms delay = %d \n\t1ms delay = %d \n\tTotal = %d\n",
spikeCountD2, spikeCountD1, spikeCountAll );
fprintf(fp, "**************************************\n\n");
fflush(fp);
}
bool MultivalueHashTable::Initialize(const unsigned max_table_entries,
const float space_usage,
const unsigned num_hash_functions)
{
bool success = HashTable::Initialize(max_table_entries, space_usage,
num_hash_functions);
target_space_usage_ = space_usage;
// + 2N 32-bit entries
CUDA_SAFE_CALL(cudaMalloc( (void**)&d_scratch_offsets_,
sizeof(unsigned) * max_table_entries ));
CUDA_SAFE_CALL(cudaMalloc( (void**)&d_scratch_is_unique_,
sizeof(unsigned) * max_table_entries ));
success &= (d_scratch_offsets_ != NULL);
success &= (d_scratch_is_unique_ != NULL);
// Allocate memory for the scan.
// + Unknown memory usage
CUDPPConfiguration config;
config.op = CUDPP_ADD;
config.datatype = CUDPP_UINT;
config.algorithm = CUDPP_SCAN;
config.options = CUDPP_OPTION_FORWARD | CUDPP_OPTION_INCLUSIVE;
CUDPPResult result = cudppPlan(theCudpp, &scanplan_, config,
max_table_entries, 1, 0);
if (CUDPP_SUCCESS != result) {
fprintf(stderr, "Failed to create plan.");
return false;
}
return success;
}
void
LiGL2D::setPbo(int image_width, int image_height)
{
makeCurrent();
iw = image_width;
ih = image_height;
GLuint oldPbo = 0;
GLuint newPbo = 0;
GLuint oldTex = 0;
if(pbo != 0){
oldPbo = pbo;
pbo = 0;
oldTex = tex;
}
if(iw != 0 && ih !=0){
glGenBuffers(1, &newPbo);
glBindBuffer(GL_ARRAY_BUFFER, newPbo);
glBufferData(GL_ARRAY_BUFFER, image_height*image_width* 4*sizeof(GLubyte),NULL, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
CUDA_SAFE_CALL(cudaGLRegisterBufferObject(newPbo));
createTexture(&tex, iw, ih);
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
pbo = newPbo;
emit sendPbo(pbo);
}
if(oldPbo != 0){
CUDA_SAFE_CALL(cudaGLUnregisterBufferObject(oldPbo));
glDeleteBuffers(1, &oldPbo);
}
if(oldTex != 0){
glDeleteTextures(1, &oldTex);
}
}
void
FiringBuffer::sync(cudaStream_t stream)
{
memcpyFromDeviceAsync(mh_buffer.get(), md_buffer.get(),
m_mapper.partitionCount() * m_pitch, stream);
CUDA_SAFE_CALL(cudaEventRecord(m_copyDone, stream));
CUDA_SAFE_CALL(cudaEventSynchronize(m_copyDone));
populateSparse(mh_buffer.get());
}
void runbench_warmup(double *cd, long size){
const long reduced_grid_size = size/(UNROLLED_MEMORY_ACCESSES)/32;
const int BLOCK_SIZE = 256;
const int TOTAL_REDUCED_BLOCKS = reduced_grid_size/BLOCK_SIZE;
dim3 dimBlock(BLOCK_SIZE, 1, 1);
dim3 dimReducedGrid(TOTAL_REDUCED_BLOCKS, 1, 1);
hipLaunchKernel(HIP_KERNEL_NAME(benchmark_func< short, BLOCK_SIZE, 0 >), dim3(dimReducedGrid), dim3(dimBlock ), 0, 0, (short)1, (short*)cd);
CUDA_SAFE_CALL( hipGetLastError() );
CUDA_SAFE_CALL( hipDeviceSynchronize() );
}
//---------------------------------------------
//GPU memory operations
//---------------------------------------------
char *D_MALLOC(size_t size)
{
char *buf = NULL;
CUDA_SAFE_CALL(cudaMalloc((void**)&buf, size));
CUDA_SAFE_CALL(cudaMemset(buf, 0, size));
#ifdef __DEBUG__
# ifdef __ALLOC__
BenLog("+d%d bytes\n", size);
# endif //__ALLOC__
d_dmemUsage += size;
#endif
return buf;
}
void ParticleListCPUSorted::copy_from(const ParticleList* list_in)
{
ispecies = list_in -> ispecies;
// Free realkind arrays
if(list_in -> device_type == 0){
for(int i=0;i<ParticleList_nfloats;i++)
{
memcpy(*get_float(i),*(list_in->get_float(i)),nptcls*sizeof(realkind));
}
// Allocate int arrays
for(int i=0;i<ParticleList_nints;i++)
{
memcpy(*get_int(i),*(list_in->get_int(i)),nptcls*sizeof(int));
}
// allocate short ints for cluster id's
memcpy(cluster_id,list_in->cluster_id,nptcls*sizeof(int));
// memcpy(num_subcycles,list_in->num_subcycles,nptcls*sizeof(int));
//
// memcpy(num_piccard,list_in->num_piccard,nptcls*sizeof(double));
// memcpy(num_piccard2,list_in->num_piccard2,nptcls*sizeof(double));
}
else if(list_in->device_type == 1)
{
#ifndef NO_CUDA
enum cudaMemcpyKind kind = cudaMemcpyDeviceToHost;
// Free realkind arrays
for(int i=0;i<ParticleList_nfloats;i++)
{
CUDA_SAFE_CALL(cudaMemcpyAsync(*get_float(i),*(list_in->get_float(i)),nptcls*sizeof(realkind),kind));
}
// Allocate int arrays
for(int i=0;i<ParticleList_nints;i++)
{
CUDA_SAFE_CALL(cudaMemcpyAsync(*get_int(i),*(list_in->get_int(i)),nptcls*sizeof(int),kind));
}
// allocate short ints for cluster id's
CUDA_SAFE_CALL(cudaMemcpyAsync(cluster_id,(list_in->cluster_id),nptcls*sizeof(int),kind));
CUDA_SAFE_CALL(cudaDeviceSynchronize());
#endif
}
}
CudaGridMap::CudaGridMap(const Vec3i &numGridPoints, const Vec3i &numGridPointsPadded, const double *inputEnergies, cudaStream_t stream)
: stream(stream), numGridPoints(numGridPoints), numGridPointsPadded(numGridPointsPadded)
{
// Allocate the padded grid in global memory
CUDA_SAFE_CALL(cudaMalloc((void**)&energiesDevice, sizeof(float) * numGridPointsPadded.Cube()));
// Convert doubles to floats and save them in page-locked memory
int numGridPointsPerMap = numGridPoints.Cube();
CUDA_SAFE_CALL(cudaMallocHost((void**)&energiesHost, sizeof(float) * numGridPointsPerMap));
std::transform(inputEnergies, inputEnergies + numGridPointsPerMap, energiesHost, typecast<float, double>);
// Copy the initial energies from the original grid to the padded one in global memory
// Elements in the area of padding will stay uninitialized
copyGridMapPadded(energiesDevice, numGridPointsPadded, energiesHost, numGridPoints, cudaMemcpyHostToDevice);
}
void CudaUVMSpace::deallocate( void * const arg_alloc_ptr , const size_t /* arg_alloc_size */ ) const
{
try {
Kokkos::Impl::num_uvm_allocations -= 1;
CUDA_SAFE_CALL( cudaFree( arg_alloc_ptr ) );
} catch(...) {}
}
::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 ;
}
int main( int, char ** )
{
do_main();
CUDA_SAFE_CALL( cudaDeviceReset() );
return 0;
}
void CudaGridMap::copyGridMapPadded(float *dst, const Vec3i &numGridPointsDst,
const float *src, const Vec3i &numGridPointsSrc,
cudaMemcpyKind kind)
{
Vec3i numGridPointsMin = Vec3i(Mathi::Min(numGridPointsDst.x, numGridPointsSrc.x),
Mathi::Min(numGridPointsDst.y, numGridPointsSrc.y),
Mathi::Min(numGridPointsDst.z, numGridPointsSrc.z));
int numGridPointsDstXMulY = numGridPointsDst.x * numGridPointsDst.y;
int numGridPointsSrcXMulY = numGridPointsSrc.x * numGridPointsSrc.y;
for (int z = 0; z < numGridPointsMin.z; z++)
{
// Set the base of output indices from z
int outputIndexZBaseDst = z * numGridPointsDstXMulY;
int outputIndexZBaseSrc = z * numGridPointsSrcXMulY;
for (int y = 0; y < numGridPointsMin.y; y++)
{
// Set the base of output indices from (z,y)
int outputIndexZYBaseDst = outputIndexZBaseDst + y * numGridPointsDst.x;
int outputIndexZYBaseSrc = outputIndexZBaseSrc + y * numGridPointsSrc.x;
// Copy one row in axis X
CUDA_SAFE_CALL(cudaMemcpyAsync(dst + outputIndexZYBaseDst, src + outputIndexZYBaseSrc, sizeof(float) * numGridPointsMin.x, kind, stream));
}
}
}
void NodeFieldData::allocate(PlasmaData* _pdata)
{
pdata = _pdata;
nx = pdata->nx;
ny = pdata->ny;
nz = pdata->nz;
cpu_fields = new FieldDataCPU();
cpu_fields -> allocate(pdata);
if(pdata->node_info->nGPU > 0)
{
gpu_fields = (FieldDataGPU*)malloc(pdata->node_info->nGPU * sizeof(FieldDataGPU));
#pragma omp parallel for
for(int i=0;i<pdata->node_info->nGPU;i++)
{
CUDA_SAFE_CALL(cudaSetDevice(pdata->thread_info[pdata->node_info->nspecies+i]->gpu_info->igpu));
gpu_fields[i] = *(new FieldDataGPU());
gpu_fields[i].allocate(pdata);
}
}
if(pdata->node_info->nMIC > 0)
{
mic_fields = new FieldDataMIC();
mic_fields -> allocate(pdata);
}
bcast_timer = new CPUTimer();
}
std::vector<int> host::QueryDevices() {
int device_count = 0;
CUDA_SAFE_CALL(cudaGetDeviceCount(&device_count));
if (device_count < 1) {
fprintf(stderr, "No suitable CUDA devices found!\n");
exit(EXIT_FAILURE);
}
std::vector<int> device_ids;
for (int i = 0; i < device_count; i++) {
cudaDeviceProp device_prop;
CUDA_SAFE_CALL(cudaGetDeviceProperties(&device_prop, i));
int compute_cap_major = device_prop.major;
int compute_cap_minor = device_prop.minor;
int core_count = ConvertSMVer2Cores(compute_cap_major, compute_cap_minor) * device_prop.multiProcessorCount;
float clock_speed = device_prop.clockRate * 1e-6f;
float mem_size = device_prop.totalGlobalMem / 1024.0f / 1024.0f;
if (compute_cap_major >= 2) {
device_ids.push_back(i);
printf("\t[%d] %s (%d.%d, %d cores, %.2f GHz, %.2f MB)\n",
i,
device_prop.name,
compute_cap_major,
compute_cap_minor,
core_count,
clock_speed,
mem_size);
} else {
printf("\t[%d] %s (%d.%d not usable)\n",
i,
device_prop.name,
compute_cap_major,
compute_cap_minor);
}
}
if (device_ids.size() == 0) {
fprintf(stderr, "No suitable CUDA devices found!\n");
exit(EXIT_FAILURE);
}
return device_ids;
}
void Matrix::allocate_rows(int num_rows)
{
deallocate_rows();
set_num_rows(num_rows);
if( num_rows == 0 )
return;
CUDA_SAFE_CALL( cudaMalloc((void**) &m_rows, (num_rows+1)*sizeof(int)) );
}
void Matrix::allocate_vals(int num_vals)
{
deallocate_vals();
set_num_vals(num_vals);
if( num_vals == 0 )
return;
CUDA_SAFE_CALL( cudaMalloc((void**) &m_vals, 16*num_vals*sizeof(double)) );
}
void * CudaUVMSpace::allocate( const size_t arg_alloc_size ) const
{
void * ptr = NULL;
CUDA_SAFE_CALL( cudaMallocManaged( &ptr, arg_alloc_size , cudaMemAttachGlobal ) );
return ptr ;
}
void * CudaHostPinnedSpace::allocate( const size_t arg_alloc_size ) const
{
void * ptr = NULL;
CUDA_SAFE_CALL( cudaHostAlloc( &ptr, arg_alloc_size , cudaHostAllocDefault ) );
return ptr ;
}
void * CudaSpace::allocate( const size_t arg_alloc_size ) const
{
void * ptr = NULL;
CUDA_SAFE_CALL( cudaMalloc( &ptr, arg_alloc_size ) );
return ptr ;
}
void Matrix::allocate_cols(int num_cols)
{
deallocate_cols();
set_num_cols(num_cols);
if( num_cols == 0 )
return;
CUDA_SAFE_CALL( cudaMalloc((void**) &m_cols, num_cols*sizeof(int)) );
}
CudaFloatTexture1D::CudaFloatTexture1D(int width, const double *data, CudaAction action, cudaStream_t stream, CudaInternalAPI *api)
{
channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
// Allocate the texture on the GPU...
CUDA_SAFE_CALL(cudaMallocArray(&deviceArray, &channelDesc, width, 1));
// ... and in page-locked system memory
CUDA_SAFE_CALL(cudaMallocHost((void**)&hostMem, sizeof(float) * width));
// Convert doubles to floats and save them to page-locked system memory
std::transform(data, data + width, hostMem, typecast<float, double>);
// Copy floats from the page-locked memory to the GPU
CUDA_SAFE_CALL(cudaMemcpyToArrayAsync(deviceArray, 0, 0, hostMem, sizeof(float) * width, cudaMemcpyHostToDevice, stream));
if (action == BindToKernel)
api->setDistDepDielTexture(deviceArray, &channelDesc);
}
void InitCUDA(int device)
{
///////////////////////////
// CUDA initialisation
///////////////////////////
int deviceCount;
CUDA_SAFE_CALL(cudaGetDeviceCount(&deviceCount));
if (deviceCount == 0) std::cout << "There is no device supporting CUDA" << std::endl;
CUDA_SAFE_CALL(cudaSetDevice(device));
cudaDeviceProp deviceProp;
CUDA_SAFE_CALL(cudaGetDeviceProperties(&deviceProp, device));
std::cout << "Device " << device << ": " << deviceProp.name << std::endl;
// or
// CUT_DEVICE_INIT(); // with --device=1 (num device chosen)
}
void CompactingHashTable::Release() {
HashTable::Release();
CUDA_SAFE_CALL(cudaFree(d_unique_keys_));
CUDA_SAFE_CALL(cudaFree(d_scratch_cuckoo_keys_));
CUDA_SAFE_CALL(cudaFree(d_scratch_counts_));
CUDA_SAFE_CALL(cudaFree(d_scratch_unique_ids_));
d_unique_keys_ = NULL;
d_scratch_cuckoo_keys_ = NULL;
d_scratch_counts_ = NULL;
d_scratch_unique_ids_ = NULL;
if (scanplan_) {
cudppDestroyPlan(scanplan_);
}
scanplan_ = 0;
unique_keys_size_ = 0;
}
CTfactory( const VolumeGPU<T>& src,
U& texRef,
const cudaTextureFilterMode fm = cudaFilterModePoint,
const cudaTextureAddressMode am = cudaAddressModeClamp,
const int norm = false ) : dca_data(NULL) {
// Check for valid input
if( src.d_data.ptr == NULL ) {
std::cerr << __FUNCTION__
<< ": Source has no data"
<< std::endl;
abort();
}
// Allocate memory
cudaChannelFormatDesc cd = cudaCreateChannelDesc<T>();
cudaExtent tmpExtent = ExtentFromDims( src.dims );
CUDA_SAFE_CALL( cudaMalloc3DArray( &(this->dca_data),
&cd,
tmpExtent ) );
// Do the copy
cudaMemcpy3DParms cp = {0};
cp.srcPtr = src.d_data;
cp.dstArray = this->dca_data;
cp.extent = tmpExtent;
cp.kind = cudaMemcpyDeviceToDevice;
CUDA_SAFE_CALL( cudaMemcpy3D( &cp ) );
// Bind the texture
texRef.normalized = norm;
texRef.addressMode[0] = am;
texRef.addressMode[1] = am;
texRef.addressMode[2] = am;
texRef.filterMode = fm;
CUDA_SAFE_CALL( cudaBindTextureToArray( texRef, this->dca_data ) );
}
//------------------------------------------------
//free memory on device and set the pointer to NULL
//
//param : buf
//------------------------------------------------
void D_FREE(void *buf, size_t size)
{
CUDA_SAFE_CALL(cudaFree(buf));
buf = NULL;
#ifdef __DEBUG__
# ifdef __ALLOC__
BenLog("-d%d bytes\n", size);
# endif //__ALLOC__
d_dmemUsage -= size;
#endif
}
bool CompactingHashTable::Initialize(const unsigned max_table_entries,
const float space_usage,
const unsigned num_functions)
{
bool success = HashTable::Initialize(max_table_entries, space_usage,
num_functions);
unsigned slots_to_allocate = table_size_ + kStashSize;
CUDA_SAFE_CALL(cudaMalloc( (void**)&d_scratch_cuckoo_keys_,
sizeof(unsigned) * slots_to_allocate ));
CUDA_SAFE_CALL(cudaMalloc( (void**)&d_scratch_counts_,
sizeof(unsigned) * slots_to_allocate ));
CUDA_SAFE_CALL(cudaMalloc( (void**)&d_scratch_unique_ids_,
sizeof(unsigned) * slots_to_allocate ));
success &= d_scratch_cuckoo_keys_ != NULL;
success &= d_scratch_counts_ != NULL;
success &= d_scratch_unique_ids_ != NULL;
return success;
}
unsigned int cSystem::getNumGPUs(void)
{
int nGPU;
#ifdef __GEM_USE_CUDA__
CUDA_SAFE_CALL(cudaGetDeviceCount(&nGPU));
#else
nGPU = 0;
#endif
return (unsigned int) nGPU;
}
float finalizeEvents(hipEvent_t start, hipEvent_t stop){
CUDA_SAFE_CALL( hipGetLastError() );
CUDA_SAFE_CALL( hipEventRecord(stop, 0) );
CUDA_SAFE_CALL( hipEventSynchronize(stop) );
float kernel_time;
CUDA_SAFE_CALL( hipEventElapsedTime(&kernel_time, start, stop) );
CUDA_SAFE_CALL( hipEventDestroy(start) );
CUDA_SAFE_CALL( hipEventDestroy(stop) );
return kernel_time;
}
