From 34fc049ad5b0a4b2fca243cb1626e770b6f7cbc3 Mon Sep 17 00:00:00 2001 From: foreman Date: Tue, 14 Mar 2017 16:47:51 -0400 Subject: [PATCH] P4 to Git Change 1385583 by gandryey@gera-w8 on 2017/03/14 16:43:09 SWDEV-79445 - OCL generic changes and code clean-up - replace NULL with nullptr in the ROCm backend Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocappprofile.cpp#5 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#16 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.hpp#7 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompiler.cpp#28 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/roccompilerlib.cpp#5 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#44 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#18 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.cpp#20 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rockernel.hpp#14 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#14 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.hpp#7 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprintf.cpp#6 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprintf.hpp#4 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#59 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#20 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocsettings.cpp#15 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#33 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#9 edit --- rocclr/runtime/device/rocm/rocappprofile.cpp | 4 +- rocclr/runtime/device/rocm/rocblit.cpp | 118 ++++++++-------- rocclr/runtime/device/rocm/rocblit.hpp | 2 +- rocclr/runtime/device/rocm/roccompiler.cpp | 12 +- rocclr/runtime/device/rocm/roccompilerlib.cpp | 6 +- rocclr/runtime/device/rocm/rocdevice.cpp | 126 +++++++++--------- rocclr/runtime/device/rocm/rocdevice.hpp | 34 ++--- rocclr/runtime/device/rocm/rockernel.cpp | 6 +- rocclr/runtime/device/rocm/rockernel.hpp | 2 +- rocclr/runtime/device/rocm/rocmemory.cpp | 76 +++++------ rocclr/runtime/device/rocm/rocmemory.hpp | 6 +- rocclr/runtime/device/rocm/rocprintf.cpp | 14 +- rocclr/runtime/device/rocm/rocprintf.hpp | 2 +- rocclr/runtime/device/rocm/rocprogram.cpp | 68 +++++----- rocclr/runtime/device/rocm/rocprogram.hpp | 2 +- rocclr/runtime/device/rocm/rocsettings.cpp | 4 +- rocclr/runtime/device/rocm/rocvirtual.cpp | 82 ++++++------ rocclr/runtime/device/rocm/rocvirtual.hpp | 4 +- 18 files changed, 284 insertions(+), 284 deletions(-) diff --git a/rocclr/runtime/device/rocm/rocappprofile.cpp b/rocclr/runtime/device/rocm/rocappprofile.cpp index 61418e545b..be864b932c 100644 --- a/rocclr/runtime/device/rocm/rocappprofile.cpp +++ b/rocclr/runtime/device/rocm/rocappprofile.cpp @@ -16,8 +16,8 @@ amd::AppProfile* rocCreateAppProfile() { amd::AppProfile* appProfile = new roc::AppProfile; - if ((appProfile == NULL) || !appProfile->init()) { - return NULL; + if ((appProfile == nullptr) || !appProfile->init()) { + return nullptr; } return appProfile; diff --git a/rocclr/runtime/device/rocm/rocblit.cpp b/rocclr/runtime/device/rocm/rocblit.cpp index 2179ed658a..7cd64d1c6e 100644 --- a/rocclr/runtime/device/rocm/rocblit.cpp +++ b/rocclr/runtime/device/rocm/rocblit.cpp @@ -15,7 +15,7 @@ DmaBlitManager::DmaBlitManager(VirtualGPU& gpu, Setup setup) : HostBlitManager(gpu, setup) , MinSizeForPinnedTransfer(dev().settings().pinnedMinXferSize_) , completeOperation_(false) - , context_(NULL) + , context_(nullptr) { } @@ -87,7 +87,7 @@ DmaBlitManager::readBuffer( // Find the partial size for unaligned copy size_t partial = reinterpret_cast(dstHost) - tmpHost; - amd::Memory* pinned = NULL; + amd::Memory* pinned = nullptr; bool first = true; size_t tmpSize; size_t pinAllocSize; @@ -114,7 +114,7 @@ DmaBlitManager::readBuffer( // Allocate a GPU resource for pinning pinned = pinHostMemory(tmpHost, pinAllocSize, partial2); - if (pinned != NULL) { + if (pinned != nullptr) { // Get device memory for this virtual device Memory* dstMemory = dev().getRocMemory(pinned); @@ -270,7 +270,7 @@ DmaBlitManager::writeBuffer( // Find the partial size for unaligned copy size_t partial = reinterpret_cast(srcHost) - tmpHost; - amd::Memory* pinned = NULL; + amd::Memory* pinned = nullptr; bool first = true; size_t tmpSize; size_t pinAllocSize; @@ -298,7 +298,7 @@ DmaBlitManager::writeBuffer( // Allocate a GPU resource for pinning pinned = pinHostMemory(tmpHost, pinAllocSize, partial2); - if (pinned != NULL) { + if (pinned != nullptr) { // Get device memory for this virtual device Memory* srcMemory = dev().getRocMemory(pinned); @@ -464,7 +464,7 @@ DmaBlitManager::copyBufferRect( hsa_status_t status = hsa_amd_memory_async_copy( (reinterpret_cast
(dst) + dstOffset), dstAgent, (reinterpret_cast(src) + srcOffset), - srcAgent, size[0], 0, NULL, completion_signal_); + srcAgent, size[0], 0, nullptr, completion_signal_); if (status != HSA_STATUS_SUCCESS) { LogPrintfError("DMA buffer failed with code %d", status); return false; @@ -691,7 +691,7 @@ bool DmaBlitManager::hsaCopyStaged( memcpy(hsaBuffer, hostSrc + offset, size); status = hsa_amd_memory_async_copy( hostDst + offset, dev().getBackendDevice(), hsaBuffer, - dev().getCpuAgent(), size, 0, NULL, completion_signal_); + dev().getCpuAgent(), size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { hsa_signal_value_t val = hsa_signal_wait_acquire(completion_signal_, @@ -715,7 +715,7 @@ bool DmaBlitManager::hsaCopyStaged( // Copy data from Device to Host status = hsa_amd_memory_async_copy(hsaBuffer, dev().getCpuAgent(), hostSrc + offset, dev().getBackendDevice(), - size, 0, NULL, completion_signal_); + size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { hsa_signal_value_t val = hsa_signal_wait_acquire( completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), @@ -741,17 +741,17 @@ bool DmaBlitManager::hsaCopyStaged( KernelBlitManager::KernelBlitManager( VirtualGPU& gpu, Setup setup) : DmaBlitManager(gpu, setup) - , program_(NULL) - , constantBuffer_(NULL) + , program_(nullptr) + , constantBuffer_(nullptr) , xferBufferSize_(0) - , lockXferOps_(NULL) + , lockXferOps_(nullptr) { for (uint i = 0; i < BlitTotal; ++i) { - kernels_[i] = NULL; + kernels_[i] = nullptr; } for (uint i = 0; i < MaxXferBuffers; ++i) { - xferBuffers_[i] = NULL; + xferBuffers_[i] = nullptr; } completeOperation_ = false; @@ -760,25 +760,25 @@ KernelBlitManager::KernelBlitManager( KernelBlitManager::~KernelBlitManager() { for (uint i = 0; i < BlitTotal; ++i) { - if (NULL != kernels_[i]) { + if (nullptr != kernels_[i]) { kernels_[i]->release(); } } - if (NULL != program_) { + if (nullptr != program_) { program_->release(); } - if (NULL != context_) { + if (nullptr != context_) { // Release a dummy context context_->release(); } - if (NULL != constantBuffer_) { + if (nullptr != constantBuffer_) { constantBuffer_->release(); } for (uint i = 0; i < MaxXferBuffers; ++i) { - if (NULL != xferBuffers_[i]) { + if (nullptr != xferBuffers_[i]) { xferBuffers_[i]->release(); } } @@ -820,11 +820,11 @@ KernelBlitManager::createProgram(Device& device) // Create kernel objects for all blits for (uint i = 0; i < BlitTotal; ++i) { const amd::Symbol* symbol = program_->findSymbol(BlitName[i]); - if (symbol == NULL) { + if (symbol == nullptr) { break; } kernels_[i] = new amd::Kernel(*program_, *symbol, BlitName[i]); - if (kernels_[i] == NULL) { + if (kernels_[i] == nullptr) { break; } // Validate blit kernels for the scratch memory usage (pre SI) @@ -840,12 +840,12 @@ KernelBlitManager::createProgram(Device& device) constantBuffer_ = new (*context_) amd::Buffer(*context_, CL_MEM_ALLOC_HOST_PTR, 4 * Ki); - if ((constantBuffer_ != NULL) && !constantBuffer_->create(NULL)) { + if ((constantBuffer_ != nullptr) && !constantBuffer_->create(nullptr)) { constantBuffer_->release(); - constantBuffer_ = NULL; + constantBuffer_ = nullptr; return false; } - else if (constantBuffer_ == NULL) { + else if (constantBuffer_ == nullptr) { return false; } @@ -859,12 +859,12 @@ KernelBlitManager::createProgram(Device& device) xferBuffers_[i] = new (*context_) amd::Buffer(*context_, 0, xferBufferSize_); - if ((xferBuffers_[i] != NULL) && !xferBuffers_[i]->create(NULL)) { + if ((xferBuffers_[i] != nullptr) && !xferBuffers_[i]->create(nullptr)) { xferBuffers_[i]->release(); - xferBuffers_[i] = NULL; + xferBuffers_[i] = nullptr; return false; } - else if (xferBuffers_[i] == NULL) { + else if (xferBuffers_[i] == nullptr) { return false; } @@ -882,7 +882,7 @@ KernelBlitManager::createProgram(Device& device) } lockXferOps_ = new amd::Monitor("Transfer Ops Lock", true); - if (NULL == lockXferOps_) { + if (nullptr == lockXferOps_) { return false; } @@ -1065,7 +1065,7 @@ KernelBlitManager::copyBufferToImageKernel( // todo ROC runtime has a problem with a view for this format (dstImage->getImageFormat().image_channel_data_type != CL_UNORM_INT_101010)) { dstView = createView(gpuMem(dstMemory), newFormat, CL_MEM_WRITE_ONLY); - if (dstView != NULL) { + if (dstView != nullptr) { rejected = false; releaseView = true; } @@ -1160,7 +1160,7 @@ KernelBlitManager::copyBufferToImageKernel( // Execute the blit address parameters = captureArguments(kernels_[blitType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); if (releaseView) { // todo SRD programming could be changed to avoid a stall @@ -1269,7 +1269,7 @@ KernelBlitManager::copyImageToBufferKernel( // todo ROC runtime has a problem with a view for this format (srcImage->getImageFormat().image_channel_data_type != CL_UNORM_INT_101010)) { srcView = createView(gpuMem(srcMemory), newFormat, CL_MEM_READ_ONLY); - if (srcView != NULL) { + if (srcView != nullptr) { rejected = false; releaseView = true; } @@ -1369,7 +1369,7 @@ KernelBlitManager::copyImageToBufferKernel( // Execute the blit address parameters = captureArguments(kernels_[blitType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); if (releaseView) { // todo SRD programming could be changed to avoid a stall @@ -1423,9 +1423,9 @@ KernelBlitManager::copyImage( // Attempt to create a view if the format was rejected if (rejected) { srcView = createView(gpuMem(srcMemory), newFormat, CL_MEM_READ_ONLY); - if (srcView != NULL) { + if (srcView != nullptr) { dstView = createView(gpuMem(dstMemory), newFormat, CL_MEM_WRITE_ONLY); - if (dstView != NULL) { + if (dstView != nullptr) { rejected = false; releaseView = true; } @@ -1513,7 +1513,7 @@ KernelBlitManager::copyImage( // Execute the blit address parameters = captureArguments(kernels_[blitType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); if (releaseView) { // todo SRD programming could be changed to avoid a stall @@ -1588,7 +1588,7 @@ KernelBlitManager::readImage( size_t partial; amd::Memory* amdMemory = pinHostMemory(dstHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force SW copy result = DmaBlitManager::readImage(srcMemory, dstHost, origin, size, rowPitch, slicePitch, entire); @@ -1642,7 +1642,7 @@ KernelBlitManager::writeImage( size_t partial; amd::Memory* amdMemory = pinHostMemory(srcHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force SW copy result = DmaBlitManager::writeImage( srcHost, dstMemory, origin, size, rowPitch, slicePitch, entire); @@ -1789,7 +1789,7 @@ KernelBlitManager::copyBufferRect( // Execute the blit address parameters = captureArguments(kernels_[blitType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); synchronize(); @@ -1822,7 +1822,7 @@ KernelBlitManager::readBuffer( size_t partial; amd::Memory* amdMemory = pinHostMemory(dstHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force SW copy result = DmaBlitManager::readBuffer( srcMemory, dstHost, origin, size, entire); @@ -1878,7 +1878,7 @@ KernelBlitManager::readBufferRect( size_t partial; amd::Memory* amdMemory = pinHostMemory(dstHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force SW copy result = DmaBlitManager::readBufferRect( srcMemory, dstHost, bufRect, hostRect, size, entire); @@ -1936,7 +1936,7 @@ KernelBlitManager::writeBuffer( size_t partial; amd::Memory* amdMemory = pinHostMemory(srcHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force SW copy result = DmaBlitManager::writeBuffer( srcHost, dstMemory, origin, size, entire); @@ -1993,7 +1993,7 @@ KernelBlitManager::writeBufferRect( size_t partial; amd::Memory* amdMemory = pinHostMemory(srcHost, pinSize, partial); - if (amdMemory == NULL) { + if (amdMemory == nullptr) { // Force DMA copy with staging result = DmaBlitManager::writeBufferRect( srcHost, dstMemory, hostRect, bufRect, size, entire); @@ -2060,15 +2060,15 @@ KernelBlitManager::fillBuffer( // Program kernels arguments for the fill operation cl_mem mem = as_cl(memory.owner()); if (dwordAligned) { - setArgument(kernels_[fillType], 0, sizeof(cl_mem), NULL); + setArgument(kernels_[fillType], 0, sizeof(cl_mem), nullptr); setArgument(kernels_[fillType], 1, sizeof(cl_mem), &mem); } else { setArgument(kernels_[fillType], 0, sizeof(cl_mem), &mem); - setArgument(kernels_[fillType], 1, sizeof(cl_mem), NULL); + setArgument(kernels_[fillType], 1, sizeof(cl_mem), nullptr); } Memory* gpuCB = dev().getRocMemory(constantBuffer_); - if (gpuCB == NULL) { + if (gpuCB == nullptr) { return false; } void* constBuf = constantBuffer_->getHostMem(); @@ -2091,7 +2091,7 @@ KernelBlitManager::fillBuffer( // Execute the blit address parameters = captureArguments(kernels_[fillType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters, nullptr); releaseArguments(parameters); } @@ -2187,7 +2187,7 @@ KernelBlitManager::copyBuffer( // Execute the blit address parameters = captureArguments(kernels_[blitType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[blitType], parameters, nullptr); releaseArguments(parameters); } else { @@ -2271,7 +2271,7 @@ KernelBlitManager::fillImage( // If the image format was rejected, then attempt to create a view if (rejected) { memView = createView(gpuMem(memory), newFormat, CL_MEM_WRITE_ONLY); - if (memView != NULL) { + if (memView != nullptr) { rejected = false; releaseView = true; } @@ -2355,7 +2355,7 @@ KernelBlitManager::fillImage( // Execute the blit address parameters = captureArguments(kernels_[fillType]); - result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters, NULL); + result = gpu().submitKernelInternal(ndrange, *kernels_[fillType], parameters, nullptr); releaseArguments(parameters); if (releaseView) { // todo SRD programming could be changed to avoid a stall @@ -2391,16 +2391,16 @@ DmaBlitManager::pinHostMemory( amdMemory = gpu().findPinnedMem(tmpHost, pinAllocSize); - if (NULL != amdMemory) { + if (nullptr != amdMemory) { return amdMemory; } amdMemory = new(*context_) amd::Buffer(*context_, CL_MEM_USE_HOST_PTR, pinAllocSize); - if ((amdMemory != NULL) && !amdMemory->create(tmpHost, SysMem)) { + if ((amdMemory != nullptr) && !amdMemory->create(tmpHost, SysMem)) { amdMemory->release(); - return NULL; + return nullptr; } // Get device memory for this virtual device @@ -2408,14 +2408,14 @@ DmaBlitManager::pinHostMemory( amdMemory->setVirtualDevice(&gpu()); Memory* srcMemory = dev().getRocMemory(amdMemory); - if (srcMemory == NULL) { + if (srcMemory == nullptr) { // Release all pinned memory and attempt pinning again gpu().releasePinnedMem(); srcMemory = dev().getRocMemory(amdMemory); - if (srcMemory == NULL) { + if (srcMemory == nullptr) { // Release memory amdMemory->release(); - amdMemory = NULL; + amdMemory = nullptr; } } @@ -2433,23 +2433,23 @@ KernelBlitManager::createView( amd::Image* image = parentImage->createView( parent.owner()->getContext(), format, &gpu(), 0, flags); - if (image == NULL) { + if (image == nullptr) { LogError("[OCL] Fail to allocate view of image object"); - return NULL; + return nullptr; } Image* devImage = new roc::Image(dev(), *image); - if (devImage == NULL) { + if (devImage == nullptr) { LogError("[OCL] Fail to allocate device mem object for the view"); image->release(); - return NULL; + return nullptr; } if (!devImage->createView(parent)) { LogError("[OCL] Fail to create device mem object for the view"); delete devImage; image->release(); - return NULL; + return nullptr; } image->replaceDeviceMemory(&dev_, devImage); diff --git a/rocclr/runtime/device/rocm/rocblit.hpp b/rocclr/runtime/device/rocm/rocblit.hpp index e425598893..c2e6314406 100644 --- a/rocclr/runtime/device/rocm/rocblit.hpp +++ b/rocclr/runtime/device/rocm/rocblit.hpp @@ -41,7 +41,7 @@ public: //! Creates DmaBlitManager object virtual bool create(amd::Device& device) { - if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, NULL, &completion_signal_)) { + if (HSA_STATUS_SUCCESS != hsa_signal_create(0, 0, nullptr, &completion_signal_)) { false; } return true; diff --git a/rocclr/runtime/device/rocm/roccompiler.cpp b/rocclr/runtime/device/rocm/roccompiler.cpp index d21d80c4d5..3355657580 100644 --- a/rocclr/runtime/device/rocm/roccompiler.cpp +++ b/rocclr/runtime/device/rocm/roccompiler.cpp @@ -30,7 +30,7 @@ static std::string llvmBin_(amd::Os::getEnvironment("LLVM_BIN")); #endif // defined(WITH_LIGHTNING_COMPILER) //CLC_IN_PROCESS_CHANGE -extern int openclFrontEnd(const char* cmdline, std::string*, std::string* typeInfo = NULL); +extern int openclFrontEnd(const char* cmdline, std::string*, std::string* typeInfo = nullptr); namespace roc { @@ -57,7 +57,7 @@ HSAILProgram::compileImpl_LC( Data* input = C->NewBufferReference(DT_CL, sourceCode.c_str(), sourceCode.length()); - if (input == NULL) { + if (input == nullptr) { buildLog_ += "Error while creating data from source code"; return false; } @@ -65,7 +65,7 @@ HSAILProgram::compileImpl_LC( inputs.push_back(input); Buffer* output = C->NewBuffer(DT_LLVM_BC); - if (output == NULL) { + if (output == nullptr) { buildLog_ += "Error while creating buffer for the LLVM bitcode"; return false; } @@ -79,7 +79,7 @@ HSAILProgram::compileImpl_LC( std::string driverOptions(ostrstr.str()); const char* xLang = options->oVariables->XLang; - if (xLang != NULL && strcmp(xLang, "cl")) { + if (xLang != nullptr && strcmp(xLang, "cl")) { buildLog_ += "Unsupported OpenCL language.\n"; } @@ -146,7 +146,7 @@ HSAILProgram::compileImpl_LC( f.close(); Data* inc = C->NewFileReference(DT_CL_HEADER, headerFileNames[i]); - if (inc == NULL) { + if (inc == nullptr) { buildLog_ += "Error while creating data from headers"; return false; } @@ -190,7 +190,7 @@ HSAILProgram::compileImpl_LC( } File* pch = C->NewTempFile(DT_CL_HEADER); - if (pch == NULL || !pch->WriteData((const char*) hdr.first, hdr.second)) { + if (pch == nullptr || !pch->WriteData((const char*) hdr.first, hdr.second)) { buildLog_ += "Error while opening the opencl-c header "; return false; } diff --git a/rocclr/runtime/device/rocm/roccompilerlib.cpp b/rocclr/runtime/device/rocm/roccompilerlib.cpp index e933d5c393..d5c63a067a 100644 --- a/rocclr/runtime/device/rocm/roccompilerlib.cpp +++ b/rocclr/runtime/device/rocm/roccompilerlib.cpp @@ -5,7 +5,7 @@ namespace roc { -void* g_complibModule = NULL; +void* g_complibModule = nullptr; struct CompLibApi g_complibApi; // @@ -13,7 +13,7 @@ struct CompLibApi g_complibApi; // #define LOADSYMBOL(api) \ g_complibApi._##api = (pfn_##api) amd::Os::getSymbol(g_complibModule, #api); \ - if( g_complibApi._##api == NULL ) { \ + if( g_complibApi._##api == nullptr ) { \ LogError ("amd::Os::getSymbol() for exported func " #api " failed."); \ amd::Os::unloadLibrary(g_complibModule); \ return false; \ @@ -22,7 +22,7 @@ struct CompLibApi g_complibApi; bool LoadCompLib(bool offline) { g_complibModule = amd::Os::loadLibrary("amdhsacl" LP64_SWITCH(LINUX_SWITCH("32",""), "64")); - if( g_complibModule == NULL ) { + if( g_complibModule == nullptr ) { if (!offline) { LogError( "amd::Os::loadLibrary() for loading of amdhsacl.dll failed."); } diff --git a/rocclr/runtime/device/rocm/rocdevice.cpp b/rocclr/runtime/device/rocm/rocdevice.cpp index 47d13b11ea..4fbbfad0c8 100644 --- a/rocclr/runtime/device/rocm/rocdevice.cpp +++ b/rocclr/runtime/device/rocm/rocdevice.cpp @@ -164,8 +164,8 @@ bool NullDevice::create(const AMDDeviceInfo& deviceInfo) { settings_ = new Settings(); roc::Settings* hsaSettings = static_cast(settings_); - if ((hsaSettings == NULL) || !hsaSettings->create(false)) { - LogError("Error creating settings for NULL HSA device"); + if ((hsaSettings == nullptr) || !hsaSettings->create(false)) { + LogError("Error creating settings for nullptr HSA device"); return false; } // Report the device name @@ -200,8 +200,8 @@ Device::Device(hsa_agent_t bkendDevice) Device::~Device() { // Release cached map targets - for (uint i = 0; mapCache_ != NULL && i < mapCache_->size(); ++i) { - if ((*mapCache_)[i] != NULL) { + for (uint i = 0; mapCache_ != nullptr && i < mapCache_->size(); ++i) { + if ((*mapCache_)[i] != nullptr) { (*mapCache_)[i]->release(); } } @@ -215,32 +215,32 @@ Device::~Device() // Destroy transfer queue if (xferQueue_ && xferQueue_->terminate()) { delete xferQueue_; - xferQueue_ = NULL; + xferQueue_ = nullptr; } if (blitProgram_) { delete blitProgram_; - blitProgram_ = NULL; + blitProgram_ = nullptr; } - if (context_ != NULL) { + if (context_ != nullptr) { context_->release(); } if (info_.extensions_) { delete[]info_.extensions_; - info_.extensions_ = NULL; + info_.extensions_ = nullptr; } if (settings_) { delete settings_; - settings_ = NULL; + settings_ = nullptr; } } bool NullDevice::initCompiler(bool isOffline) { #if !defined(WITH_LIGHTNING_COMPILER) // Initializes g_complibModule and g_complibApi if they were not initialized - if( g_complibModule == NULL ){ + if( g_complibModule == nullptr ){ if (!LoadCompLib(isOffline)) { if (!isOffline) { LogError("Error - could not find the compiler library"); @@ -252,7 +252,7 @@ bool NullDevice::initCompiler(bool isOffline) { //This is destroyed in Device::teardown acl_error error; if (!compilerHandle_) { - compilerHandle_ = g_complibApi._aclCompilerInit(NULL, &error); + compilerHandle_ = g_complibApi._aclCompilerInit(nullptr, &error); if (error != ACL_SUCCESS) { LogError("Error initializing the compiler handle"); return false; @@ -265,16 +265,16 @@ bool NullDevice::initCompiler(bool isOffline) { bool NullDevice::destroyCompiler() { #if defined(WITH_LIGHTNING_COMPILER) delete compilerHandle_; - compilerHandle_ = NULL; + compilerHandle_ = nullptr; #else // !defined(WITH_LIGHTNING_COMPILER) - if (compilerHandle_ != NULL) { + if (compilerHandle_ != nullptr) { acl_error error = g_complibApi._aclCompilerFini(compilerHandle_); if (error != ACL_SUCCESS) { LogError("Error closing the compiler"); return false; } } - if( g_complibModule != NULL ){ + if( g_complibModule != nullptr ){ UnloadCompLib(); } #endif // !defined(WITH_LIGHTNING_COMPILER) @@ -326,12 +326,12 @@ bool NullDevice::init() { NullDevice::~NullDevice() { if (info_.extensions_) { delete[]info_.extensions_; - info_.extensions_ = NULL; + info_.extensions_ = nullptr; } if (settings_) { delete settings_; - settings_ = NULL; + settings_ = nullptr; } } @@ -470,7 +470,7 @@ bool Device::init() sizeof(amd_loader_ext_table), &amd_loader_ext_table); if (HSA_STATUS_SUCCESS != - hsa_iterate_agents(iterateAgentCallback, NULL)) { + hsa_iterate_agents(iterateAgentCallback, nullptr)) { return false; } @@ -608,30 +608,30 @@ Device::create() // Create a dummy context context_ = new amd::Context(devices, info); - if (context_ == NULL) { + if (context_ == nullptr) { return false; } blitProgram_ = new BlitProgram(context_); // Create blit programs - if (blitProgram_ == NULL || !blitProgram_->create(this)) { + if (blitProgram_ == nullptr || !blitProgram_->create(this)) { delete blitProgram_; - blitProgram_ = NULL; + blitProgram_ = nullptr; LogError("Couldn't create blit kernels!"); return false; } mapCacheOps_ = new amd::Monitor("Map Cache Lock", true); - if (NULL == mapCacheOps_) { + if (nullptr == mapCacheOps_) { return false; } mapCache_ = new std::vector(); - if (mapCache_ == NULL) { + if (mapCache_ == nullptr) { return false; } // Use just 1 entry by default for the map cache - mapCache_->push_back(NULL); + mapCache_->push_back(nullptr); if (settings().stagedXferSize_ != 0) { // Initialize staged write buffers @@ -682,7 +682,7 @@ Device::mapHSADeviceToOpenCLDevice(hsa_agent_t dev) // Create HSA settings settings_ = new Settings(); roc::Settings* hsaSettings = static_cast(settings_); - if ((hsaSettings == NULL) || + if ((hsaSettings == nullptr) || !hsaSettings->create((agent_profile_ == HSA_PROFILE_FULL))) { return false; } @@ -731,7 +731,7 @@ Device::mapHSADeviceToOpenCLDevice(hsa_agent_t dev) hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, void* data) { - if (data == NULL) { + if (data == nullptr) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -764,7 +764,7 @@ hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, hsa_status_t Device::iterateCpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, void* data) { - if (data == NULL) { + if (data == nullptr) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -1173,7 +1173,7 @@ Device::populateOCLDeviceConstants() device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) { - bool profiling = (queue != NULL) && + bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE); // Initialization of heap and other resources occur during the command @@ -1182,7 +1182,7 @@ Device::createVirtualDevice(amd::CommandQueue* queue) if (!virtualDevice->create(profiling)) { delete virtualDevice; - return NULL; + return nullptr; } if (profiling) { @@ -1281,7 +1281,7 @@ Device::findMapTarget(size_t size) const // Must be serialised for access amd::ScopedLock lk(*mapCacheOps_); - amd::Memory* map = NULL; + amd::Memory* map = nullptr; size_t minSize = 0; size_t maxSize = 0; uint mapId = mapCache_->size(); @@ -1289,7 +1289,7 @@ Device::findMapTarget(size_t size) const // Find if the list has a map target of appropriate size for (uint i = 0; i < mapCache_->size(); i++) { - if ((*mapCache_)[i] != NULL) { + if ((*mapCache_)[i] != nullptr) { // Requested size is smaller than the entry size if (size < (*mapCache_)[i]->getSize()) { if ((minSize == 0) || @@ -1316,12 +1316,12 @@ Device::findMapTarget(size_t size) const // Check if we found any map target if (mapId < mapCache_->size()) { map = (*mapCache_)[mapId]; - (*mapCache_)[mapId] = NULL; + (*mapCache_)[mapId] = nullptr; } // If cache is full, then release the biggest map target else if (releaseId < mapCache_->size()) { (*mapCache_)[releaseId]->release(); - (*mapCache_)[releaseId] = NULL; + (*mapCache_)[releaseId] = nullptr; } return map; @@ -1339,7 +1339,7 @@ Device::addMapTarget(amd::Memory* memory) const } // Find if the list has a map target of appropriate size for (uint i = 0; i < mapCache_->size(); ++i) { - if ((*mapCache_)[i] == NULL) { + if ((*mapCache_)[i] == nullptr) { (*mapCache_)[i] = memory; return true; } @@ -1361,7 +1361,7 @@ Device::getRocMemory(amd::Memory* mem) const device::Memory* Device::createMemory(amd::Memory &owner) const { - roc::Memory* memory = NULL; + roc::Memory* memory = nullptr; if (owner.asBuffer()) { memory = new roc::Buffer(*this, owner); } @@ -1372,8 +1372,8 @@ Device::createMemory(amd::Memory &owner) const LogError("Unknown memory type"); } - if (memory == NULL) { - return NULL; + if (memory == nullptr) { + return nullptr; } bool result = memory->create(); @@ -1381,7 +1381,7 @@ Device::createMemory(amd::Memory &owner) const if (!result) { LogError("Failed creating memory"); delete memory; - return NULL; + return nullptr; } // Transfer data only if OCL context has one device. @@ -1395,25 +1395,25 @@ Device::createMemory(amd::Memory &owner) const amd::Image* imageView = owner.asImage()->createView( owner.getContext(), owner.asImage()->getImageFormat(), xferQueue()); - if (imageView == NULL) { + if (imageView == nullptr) { LogError("[OCL] Fail to allocate view of image object"); - return NULL; + return nullptr; } Image* devImageView = new roc::Image(static_cast(*this), *imageView); - if (devImageView == NULL) { + if (devImageView == nullptr) { LogError("[OCL] Fail to allocate device mem object for the view"); imageView->release(); - return NULL; + return nullptr; } - if (devImageView != NULL && + if (devImageView != nullptr && !devImageView->createView(static_cast(*memory))) { LogError("[OCL] Fail to create device mem object for the view"); delete devImageView; imageView->release(); - return NULL; + return nullptr; } imageView->replaceDeviceMemory(this, devImageView); @@ -1436,7 +1436,7 @@ Device::createMemory(amd::Memory &owner) const if (!result) { delete memory; - return NULL; + return nullptr; } return memory; @@ -1444,7 +1444,7 @@ Device::createMemory(amd::Memory &owner) const void* Device::hostAlloc(size_t size, size_t alignment, bool atomics) const { - void* ptr = NULL; + void* ptr = nullptr; const hsa_amd_memory_pool_t segment = (!atomics) ? (system_coarse_segment_.handle != 0) ? system_coarse_segment_ @@ -1454,14 +1454,14 @@ Device::hostAlloc(size_t size, size_t alignment, bool atomics) const { hsa_status_t stat = hsa_amd_memory_pool_allocate(segment, size, 0, &ptr); if (stat != HSA_STATUS_SUCCESS) { LogError("Fail allocation host memory"); - return NULL; + return nullptr; } stat = hsa_amd_agents_allow_access(gpu_agents_.size(), &gpu_agents_[0], - NULL, ptr); + nullptr, ptr); if (stat != HSA_STATUS_SUCCESS) { LogError("Fail hsa_amd_agents_allow_access"); - return NULL; + return nullptr; } return ptr; @@ -1477,22 +1477,22 @@ void * Device::deviceLocalAlloc(size_t size) const { if (gpuvm_segment_.handle == 0 || gpuvm_segment_max_alloc_ == 0) { - return NULL; + return nullptr; } - void *ptr = NULL; + void *ptr = nullptr; hsa_status_t stat = hsa_amd_memory_pool_allocate(gpuvm_segment_, size, 0, &ptr); if (stat != HSA_STATUS_SUCCESS) { LogError("Fail allocation local memory"); - return NULL; + return nullptr; } stat = hsa_memory_assign_agent(ptr, _bkendDevice, HSA_ACCESS_PERMISSION_RW); if (stat != HSA_STATUS_SUCCESS) { LogError("Fail assigning local memory to agent"); memFree(ptr, size); - return NULL; + return nullptr; } return ptr; @@ -1511,25 +1511,25 @@ Device::memFree(void *ptr, size_t size) const void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags, void* svmPtr) const { - amd::Memory* mem = NULL; - if (NULL == svmPtr) { + amd::Memory* mem = nullptr; + if (nullptr == svmPtr) { bool atomics = (flags & CL_MEM_SVM_ATOMICS) != 0; void* ptr = hostAlloc(size, alignment, atomics); - if (ptr != NULL) { + if (ptr != nullptr) { // Copy paste from ORCA code. // create a hidden buffer, which will allocated on the device later mem = new (context) amd::Buffer(context, CL_MEM_USE_HOST_PTR, size, ptr); - if (mem == NULL) { + if (mem == nullptr) { LogError("failed to create a svm mem object!"); - return NULL; + return nullptr; } if (!mem->create(ptr)) { LogError("failed to create a svm hidden buffer!"); mem->release(); - return NULL; + return nullptr; } // add the information to context so that we can use it later. @@ -1538,15 +1538,15 @@ Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_me return ptr; } else { - return NULL; + return nullptr; } } else { // Copy paste from ORCA code. // Find the existing amd::mem object mem = amd::SvmManager::FindSvmBuffer(svmPtr); - if (NULL == mem) { - return NULL; + if (nullptr == mem) { + return nullptr; } return svmPtr; @@ -1556,9 +1556,9 @@ Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_me void Device::svmFree(void* ptr) const { - amd::Memory * svmMem = NULL; + amd::Memory * svmMem = nullptr; svmMem = amd::SvmManager::FindSvmBuffer(ptr); - if (NULL != svmMem) { + if (nullptr != svmMem) { svmMem->release(); amd::SvmManager::RemoveSvmBuffer(ptr); hostFree(ptr); diff --git a/rocclr/runtime/device/rocm/rocdevice.hpp b/rocclr/runtime/device/rocm/rocdevice.hpp index 68c14069a3..5481789772 100644 --- a/rocclr/runtime/device/rocm/rocdevice.hpp +++ b/rocclr/runtime/device/rocm/rocdevice.hpp @@ -81,11 +81,11 @@ public: const Settings &settings() const { return reinterpret_cast(*settings_); } //! Construct an HSAIL program object from the ELF assuming it is valid - virtual device::Program *createProgram(amd::option::Options* options = NULL); + virtual device::Program *createProgram(amd::option::Options* options = nullptr); const AMDDeviceInfo& deviceInfo() const { return deviceInfo_; } - //! Gets the backend device for the NULL device type + //! Gets the backend device for the Null device type virtual hsa_agent_t getBackendDevice() const { ShouldNotReachHere(); const hsa_agent_t kInvalidAgent = { 0 }; @@ -105,9 +105,9 @@ public: //! Create a new virtual device environment. virtual device::VirtualDevice* createVirtualDevice( - amd::CommandQueue* queue = NULL) { + amd::CommandQueue* queue = nullptr) { ShouldNotReachHere(); - return NULL; + return nullptr; } virtual bool registerSvmMemory(void* ptr, size_t size) const { @@ -119,10 +119,10 @@ public: ShouldNotReachHere(); } - //! Just returns NULL for the dummy device + //! Just returns nullptr for the dummy device virtual device::Memory* createMemory(amd::Memory& owner) const { ShouldNotReachHere(); - return NULL; } + return nullptr; } //! Sampler object allocation virtual bool createSampler( @@ -134,16 +134,16 @@ public: return true; } - //! Just returns NULL for the dummy device + //! Just returns nullptr for the dummy device virtual device::Memory* createView( amd::Memory& owner, //!< Owner memory object const device::Memory& parent //!< Parent device memory object for the view ) const { ShouldNotReachHere(); - return NULL; + return nullptr; } - //! Just returns NULL for the dummy device + //! Just returns nullptr for the dummy device virtual void* svmAlloc( amd::Context& context, //!< The context used to create a buffer size_t size, //!< size of svm spaces @@ -152,10 +152,10 @@ public: void* svmPtr //!< existing svm pointer for mGPU case ) const { ShouldNotReachHere(); - return NULL; + return nullptr; } - //! Just returns NULL for the dummy device + //! Just returns nullptr for the dummy device virtual void svmFree( void* ptr //!< svm pointer needed to be freed ) const { @@ -322,10 +322,10 @@ public: //! Instantiate a new virtual device virtual device::VirtualDevice *createVirtualDevice( - amd::CommandQueue* queue = NULL); + amd::CommandQueue* queue = nullptr); //! Construct an HSAIL program object from the ELF assuming it is valid - virtual device::Program *createProgram(amd::option::Options* options = NULL); + virtual device::Program *createProgram(amd::option::Options* options = nullptr); virtual device::Memory *createMemory(amd::Memory &owner) const; @@ -336,16 +336,16 @@ public: ) const { //! \todo HSA team has to implement sampler allocation - *sampler = NULL; + *sampler = nullptr; return true; } - //! Just returns NULL for the dummy device + //! Just returns nullptr for the dummy device virtual device::Memory *createView( amd::Memory &owner, //!< Owner memory object const device::Memory &parent //!< Parent device memory object for the view - ) const { return NULL; } + ) const { return nullptr; } //! Reallocates the provided buffer object virtual bool reallocMemory(amd::Memory &owner) const {return true; } @@ -381,7 +381,7 @@ public: void memFree(void *ptr, size_t size) const; - virtual void* svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags = CL_MEM_READ_WRITE, void* svmPtr = NULL) const; + virtual void* svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags = CL_MEM_READ_WRITE, void* svmPtr = nullptr) const; virtual void svmFree(void* ptr) const; diff --git a/rocclr/runtime/device/rocm/rockernel.cpp b/rocclr/runtime/device/rocm/rockernel.cpp index 5c4ae4f978..094dff74cb 100644 --- a/rocclr/runtime/device/rocm/rockernel.cpp +++ b/rocclr/runtime/device/rocm/rockernel.cpp @@ -811,7 +811,7 @@ bool Kernel::init() program_->binaryElf(), RT_ARGUMENT_ARRAY, openClKernelName.c_str(), - NULL, + nullptr, &sizeOfArgList); if (errorCode != ACL_SUCCESS) { return false; @@ -841,7 +841,7 @@ bool Kernel::init() program_->binaryElf(), RT_WORK_GROUP_SIZE, openClKernelName.c_str(), - NULL, + nullptr, &sizeOfWorkGroupSize); if (errorCode != ACL_SUCCESS) { return false; @@ -903,7 +903,7 @@ bool Kernel::init() // Pull out printf metadata from the ELF size_t sizeOfPrintfList; errorCode = g_complibApi._aclQueryInfo(compileHandle, program_->binaryElf(), RT_GPU_PRINTF_ARRAY, - openClKernelName.c_str(), NULL, &sizeOfPrintfList); + openClKernelName.c_str(), nullptr, &sizeOfPrintfList); if (errorCode != ACL_SUCCESS){ return false; } diff --git a/rocclr/runtime/device/rocm/rockernel.hpp b/rocclr/runtime/device/rocm/rockernel.hpp index 75eda35238..7f0c5f8839 100644 --- a/rocclr/runtime/device/rocm/rockernel.hpp +++ b/rocclr/runtime/device/rocm/rockernel.hpp @@ -141,7 +141,7 @@ public: Argument* hsailArgAt(size_t index) const { for (auto arg : hsailArgList_) if (arg->index_ == index) return arg; assert(!"Should not reach here"); - return NULL; + return nullptr; } //! Return printf info array diff --git a/rocclr/runtime/device/rocm/rocmemory.cpp b/rocclr/runtime/device/rocm/rocmemory.cpp index 8be4d80b8b..55627d4c28 100644 --- a/rocclr/runtime/device/rocm/rocmemory.cpp +++ b/rocclr/runtime/device/rocm/rocmemory.cpp @@ -27,7 +27,7 @@ namespace roc { Memory::Memory(const roc::Device &dev, amd::Memory &owner) : device::Memory(owner) , dev_(dev) - , deviceMemory_(NULL) + , deviceMemory_(nullptr) , kind_(MEMORY_KIND_NORMAL) , pinnedMemory_(nullptr) { @@ -36,7 +36,7 @@ Memory::Memory(const roc::Device &dev, amd::Memory &owner) Memory::Memory(const roc::Device &dev, size_t size) : device::Memory(size) , dev_(dev) - , deviceMemory_(NULL) + , deviceMemory_(nullptr) , kind_(MEMORY_KIND_NORMAL) , pinnedMemory_(nullptr) { @@ -58,9 +58,9 @@ Memory::~Memory() bool Memory::allocateMapMemory(size_t allocationSize) { - assert(mapMemory_ == NULL); + assert(mapMemory_ == nullptr); - void *mapData = NULL; + void *mapData = nullptr; amd::Memory* mapMemory = dev().findMapTarget(owner()->getSize()); if (mapMemory == nullptr) { @@ -68,7 +68,7 @@ Memory::allocateMapMemory(size_t allocationSize) mapMemory = new (dev().context()) amd::Buffer( dev().context(), CL_MEM_ALLOC_HOST_PTR, owner()->getSize()); - if ((mapMemory == NULL) || (!mapMemory->create())) { + if ((mapMemory == nullptr) || (!mapMemory->create())) { LogError("[OCL] Fail to allocate map target object"); if (mapMemory) { mapMemory->release(); @@ -112,7 +112,7 @@ Memory::allocMapTarget( // Otherwise, check for host memory. void *hostMem = owner()->getHostMem(); - if (hostMem != NULL) { + if (hostMem != nullptr) { return (static_cast(hostMem) + origin[0]); } @@ -120,14 +120,14 @@ Memory::allocMapTarget( if (indirectMapCount_ == 1) { if (!allocateMapMemory(owner()->getSize())) { decIndMapCount(); - return NULL; + return nullptr; } } else { // Did the map resource allocation fail? - if (mapMemory_ == NULL) { + if (mapMemory_ == nullptr) { LogError("Could not map target resource"); - return NULL; + return nullptr; } } return reinterpret_cast
(mapMemory_->getHostMem()) + origin[0]; @@ -146,7 +146,7 @@ Memory::decIndMapCount() // Decrement the counter and release indirect map if it's the last op if (--indirectMapCount_ == 0 && - mapMemory_ != NULL) { + mapMemory_ != nullptr) { if (!dev().addMapTarget(mapMemory_)) { // Release the buffer object containing the map data. mapMemory_->release(); @@ -168,13 +168,13 @@ Memory::cpuMap( void * mapTarget = allocMapTarget(amd::Coord3D(0), amd::Coord3D(0), 0, rowPitch, slicePitch); - assert(mapTarget != NULL); + assert(mapTarget != nullptr); if (!isHostMemDirectAccess()) { if (!vDev.blitMgr().readBuffer( *this, mapTarget, amd::Coord3D(0), amd::Coord3D(size()), true)) { decIndMapCount(); - return NULL; + return nullptr; } } @@ -220,7 +220,7 @@ bool Memory::createInteropBuffer(GLenum targetType, int miplevel, size_t* metada in.obj=owner()->getInteropObj()->asGLObject()->getGLName(); in.miplevel=miplevel; in.out_driver_data_size=0; - in.out_driver_data=NULL; + in.out_driver_data=nullptr; if(!dev().mesa().Export(in, out)) return false; @@ -234,7 +234,7 @@ bool Memory::createInteropBuffer(GLenum targetType, int miplevel, size_t* metada return false; kind_=MEMORY_KIND_INTEROP; - assert(deviceMemory_!=NULL && "Interop map failed to produce a pointer!"); + assert(deviceMemory_!=nullptr && "Interop map failed to produce a pointer!"); return true; #endif @@ -244,7 +244,7 @@ void Memory::destroyInteropBuffer() { assert(kind_==MEMORY_KIND_INTEROP && "Memory must be interop type."); hsa_amd_interop_unmap_buffer(deviceMemory_); - deviceMemory_=NULL; + deviceMemory_=nullptr; } bool @@ -608,7 +608,7 @@ Buffer::~Buffer() void Buffer::destroy() { - if (owner()->parent() != NULL) { + if (owner()->parent() != nullptr) { return; } @@ -657,7 +657,7 @@ Buffer::create() //Interop buffer if(owner()->isInterop()) - return createInteropBuffer(GL_ARRAY_BUFFER, 0, NULL, NULL); + return createInteropBuffer(GL_ARRAY_BUFFER, 0, nullptr, nullptr); if (nullptr != owner()->parent()) { amd::Memory& parent = *owner()->parent(); @@ -694,12 +694,12 @@ Buffer::create() if (!(memFlags & (CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR))) { deviceMemory_ = dev().deviceLocalAlloc(size()); - if (deviceMemory_ == NULL) { + if (deviceMemory_ == nullptr) { // TODO: device memory is not enabled yet. // Fallback to system memory if exist. flags_ |= HostMemoryDirectAccess; if (dev().agent_profile() == HSA_PROFILE_FULL && - owner()->getHostMem() != NULL) { + owner()->getHostMem() != nullptr) { deviceMemory_ = owner()->getHostMem(); assert( amd::isMultipleOf( @@ -743,9 +743,9 @@ Buffer::create() return ret; } - return deviceMemory_ != NULL; + return deviceMemory_ != nullptr; } - assert(owner()->getHostMem() != NULL); + assert(owner()->getHostMem() != nullptr); flags_ |= HostMemoryDirectAccess; @@ -756,7 +756,7 @@ Buffer::create() hsa_memory_register(deviceMemory_, size()); } - return deviceMemory_ != NULL; + return deviceMemory_ != nullptr; } if (owner()->getSvmPtr() != owner()->getHostMem()) { @@ -775,7 +775,7 @@ Buffer::create() deviceMemory_ = owner()->getHostMem(); } - return deviceMemory_ != NULL; + return deviceMemory_ != nullptr; } /////////////////////////////////roc::Image////////////////////////////// @@ -848,7 +848,7 @@ Image::Image(const roc::Device& dev, amd::Memory& owner) : flags_ &= (~HostMemoryDirectAccess & ~HostMemoryRegistered); populateImageDescriptor(); hsaImageObject_.handle = 0; - originalDeviceMemory_ = NULL; + originalDeviceMemory_ = nullptr; } void @@ -942,9 +942,9 @@ Image::createInteropImage() MAKE_SCOPE_GUARD(BufferGuard, [&](){ destroyInteropBuffer(); }); amdImageDesc_=(hsa_amd_image_descriptor_t*)malloc(size); - if(amdImageDesc_==NULL) + if(amdImageDesc_==nullptr) return false; - MAKE_SCOPE_GUARD(DescGuard, [&](){ free(amdImageDesc_); amdImageDesc_=NULL; }); + MAKE_SCOPE_GUARD(DescGuard, [&](){ free(amdImageDesc_); amdImageDesc_=nullptr; }); memcpy(amdImageDesc_, meta, size); @@ -977,7 +977,7 @@ Image::create() roc::Memory *parent = static_cast(owner()->parent()->getDeviceMemory(dev_)); - if (parent == NULL) { + if (parent == nullptr) { LogError("[OCL] Fail to allocate parent image"); return false; } @@ -1012,7 +1012,7 @@ Image::create() originalDeviceMemory_ = dev().deviceLocalAlloc(alloc_size); } - if (originalDeviceMemory_ == NULL) { + if (originalDeviceMemory_ == nullptr) { originalDeviceMemory_ = dev().hostAlloc(alloc_size, 1, false); } @@ -1040,7 +1040,7 @@ Image::createView(const Memory &parent) { deviceMemory_ = parent.getDeviceMemory(); - originalDeviceMemory_ = (parent.owner()->asBuffer() != NULL) + originalDeviceMemory_ = (parent.owner()->asBuffer() != nullptr) ? deviceMemory_ : static_cast(parent).originalDeviceMemory_; @@ -1112,18 +1112,18 @@ Image::allocMapTarget( size_t offset = origin[0] * elementSize; - if (pHostMem == NULL) { + if (pHostMem == nullptr) { if (indirectMapCount_ == 1) { if (!allocateMapMemory(owner()->getSize())) { decIndMapCount(); - return NULL; + return nullptr; } } else { // Did the map resource allocation fail? - if (mapMemory_ == NULL) { + if (mapMemory_ == nullptr) { LogError("Could not map target resource"); - return NULL; + return nullptr; } } @@ -1139,7 +1139,7 @@ Image::allocMapTarget( else { slicePitchTmp = *rowPitch * region[1]; } - if (slicePitch != NULL) { + if (slicePitch != nullptr) { *slicePitch = slicePitchTmp; } @@ -1153,7 +1153,7 @@ Image::allocMapTarget( offset += image->getSlicePitch() * origin[2]; *rowPitch = image->getRowPitch(); - if (slicePitch != NULL) { + if (slicePitch != nullptr) { *slicePitch = image->getSlicePitch(); } @@ -1174,19 +1174,19 @@ Image::destroy() assert(status == HSA_STATUS_SUCCESS); } - if (owner()->parent() != NULL) { + if (owner()->parent() != nullptr) { return; } if(kind_==MEMORY_KIND_INTEROP) { free(amdImageDesc_); - amdImageDesc_=NULL; + amdImageDesc_=nullptr; destroyInteropBuffer(); return; } - if (originalDeviceMemory_ != NULL) { + if (originalDeviceMemory_ != nullptr) { dev().memFree(originalDeviceMemory_, deviceImageInfo_.size); } } diff --git a/rocclr/runtime/device/rocm/rocmemory.hpp b/rocclr/runtime/device/rocm/rocmemory.hpp index 9b60aaec56..f0b76479f2 100644 --- a/rocclr/runtime/device/rocm/rocmemory.hpp +++ b/rocclr/runtime/device/rocm/rocmemory.hpp @@ -49,7 +49,7 @@ class Memory : public device::Memory { ); // Immediate blocking write from device cache to owners's backing store. - // Marks owner as "current" by resetting the last writer to NULL. + // Marks owner as "current" by resetting the last writer to nullptr. virtual void syncHostFromCache(SyncFlags syncFlags = SyncFlags()); //! Allocates host memory for synchronization with MGPU context @@ -65,8 +65,8 @@ class Memory : public device::Memory { // Optimization for multilayer map/unmap uint startLayer = 0, //!< Start layer for multilayer map uint numLayers = 0, //!< End layer for multilayer map - size_t* rowPitch = NULL,//!< Row pitch for the device memory - size_t* slicePitch = NULL //!< Slice pitch for the device memory + size_t* rowPitch = nullptr, //!< Row pitch for the device memory + size_t* slicePitch = nullptr //!< Slice pitch for the device memory ); //! Unmap the device memory diff --git a/rocclr/runtime/device/rocm/rocprintf.cpp b/rocclr/runtime/device/rocm/rocprintf.cpp index 141fae7bd2..9dc25c810d 100644 --- a/rocclr/runtime/device/rocm/rocprintf.cpp +++ b/rocclr/runtime/device/rocm/rocprintf.cpp @@ -18,7 +18,7 @@ namespace roc { PrintfDbg::PrintfDbg(Device& device, FILE* file) - : dbgBuffer_(NULL), + : dbgBuffer_(nullptr), dbgBuffer_size_(0), dbgFile_(file), gpuDevice_(device) {} @@ -26,7 +26,7 @@ PrintfDbg::PrintfDbg(Device& device, FILE* file) PrintfDbg::~PrintfDbg() { dev().hostFree(dbgBuffer_, dbgBuffer_size_); } bool PrintfDbg::allocate(bool realloc) { - if (NULL == dbgBuffer_) { + if (nullptr == dbgBuffer_) { dbgBuffer_size_ = dev().info().printfBufferSize_; dbgBuffer_ = reinterpret_cast
( dev().hostAlloc(dbgBuffer_size_, sizeof(void*))); @@ -38,7 +38,7 @@ bool PrintfDbg::allocate(bool realloc) { dbgBuffer_ = reinterpret_cast
(dbgBuffer_size_, sizeof(void*)); } - return (NULL != dbgBuffer_) ? true : false; + return (nullptr != dbgBuffer_) ? true : false; } bool PrintfDbg::checkFloat(const std::string& fmt) const { @@ -142,7 +142,7 @@ size_t PrintfDbg::outputArgument(const std::string& fmt, bool printFloat, // Print the argument(except for string ), using standard PrintfDbg() else { - bool hlModifier = (strstr(fmt.c_str(), "hl") != NULL); + bool hlModifier = (strstr(fmt.c_str(), "hl") != nullptr); std::string hlFmt; if (hlModifier) { hlFmt = fmt; @@ -190,7 +190,7 @@ size_t PrintfDbg::outputArgument(const std::string& fmt, bool printFloat, amd::Os::printf(fmt.data(), fArg); } } else { - bool hhModifier = (strstr(fmt.c_str(), "hh") != NULL); + bool hhModifier = (strstr(fmt.c_str(), "hh") != nullptr); if (hhModifier) { // current implementation of printf in gcc 4.5.2 runtime libraries, // doesn`t recognize "hh" modifier ==> @@ -419,7 +419,7 @@ bool PrintfDbg::output(VirtualGPU& gpu, bool printfEnabled, // Get memory pointer to the staged buffer uint32_t* dbgBufferPtr = reinterpret_cast(dbgBuffer_); - if (NULL == dbgBufferPtr) { + if (nullptr == dbgBufferPtr) { return false; } @@ -432,7 +432,7 @@ bool PrintfDbg::output(VirtualGPU& gpu, bool printfEnabled, // Get a pointer to the buffer data dbgBufferPtr = reinterpret_cast(dbgBuffer_ + 2 * sizeof(uint32_t)); - if (NULL == dbgBufferPtr) { + if (nullptr == dbgBufferPtr) { return false; } diff --git a/rocclr/runtime/device/rocm/rocprintf.hpp b/rocclr/runtime/device/rocm/rocprintf.hpp index d0c8e10ca8..bcc2f252b5 100644 --- a/rocclr/runtime/device/rocm/rocprintf.hpp +++ b/rocclr/runtime/device/rocm/rocprintf.hpp @@ -43,7 +43,7 @@ class PrintfDbg : public amd::HeapObject { static const uint WorkitemDebugSize = 4096; //! constructor - PrintfDbg(Device& device, FILE* file = NULL); + PrintfDbg(Device& device, FILE* file = nullptr); //! Destructor ~PrintfDbg(); diff --git a/rocclr/runtime/device/rocm/rocprogram.cpp b/rocclr/runtime/device/rocm/rocprogram.cpp index c37b24eb5a..0fc549ffda 100644 --- a/rocclr/runtime/device/rocm/rocprogram.cpp +++ b/rocclr/runtime/device/rocm/rocprogram.cpp @@ -71,7 +71,7 @@ HSAILProgram::~HSAILProgram() #if !defined(WITH_LIGHTNING_COMPILER) acl_error error; // Free the elf binary - if (binaryElf_ != NULL) { + if (binaryElf_ != nullptr) { error = g_complibApi._aclBinaryFini(binaryElf_); if (error != ACL_SUCCESS) { LogWarning( "Error while destroying the acl binary \n" ); @@ -95,7 +95,7 @@ HSAILProgram::~HSAILProgram() HSAILProgram::HSAILProgram(roc::NullDevice& device) : Program(device), - binaryElf_(NULL) + binaryElf_(nullptr) { memset(&binOpts_, 0, sizeof(binOpts_)); binOpts_.struct_size = sizeof(binOpts_); @@ -113,7 +113,7 @@ HSAILProgram::HSAILProgram(roc::NullDevice& device) hasGlobalStores_ = false; #if defined(WITH_LIGHTNING_COMPILER) - metadata_ = NULL; + metadata_ = nullptr; #endif // defined(WITH_LIGHTNING_COMPILER) } @@ -134,7 +134,7 @@ HSAILProgram::initClBinary(char *binaryIn, size_t size) &decryptedBin, &decryptedSize, &encryptCode)) { return false; } - if (decryptedBin != NULL) { + if (decryptedBin != nullptr) { // It is decrypted binary. bin = decryptedBin; sz = decryptedSize; @@ -143,7 +143,7 @@ HSAILProgram::initClBinary(char *binaryIn, size_t size) // Both 32-bit and 64-bit are allowed! if (!amd::isElfMagic(bin)) { // Invalid binary. - if (decryptedBin != NULL) { + if (decryptedBin != nullptr) { delete[]decryptedBin; } return false; @@ -151,7 +151,7 @@ HSAILProgram::initClBinary(char *binaryIn, size_t size) clBinary()->setFlags(encryptCode); - return clBinary()->setBinary(bin, sz, (decryptedBin != NULL)); + return clBinary()->setBinary(bin, sz, (decryptedBin != nullptr)); } @@ -185,7 +185,7 @@ HSAILProgram::initBuild(amd::option::Options *options) #endif // !defined(WITH_LIGHTNING_COMPILER) if (!clBinary()->setElfOut(useELF64 ? ELFCLASS64 : ELFCLASS32, (outFileName.size() > - 0) ? outFileName.c_str() : NULL)) { + 0) ? outFileName.c_str() : nullptr)) { LogError("Setup elf out for gpu failed"); return false; } @@ -201,7 +201,7 @@ HSAILProgram::finiBuild(bool isBuildGood) if (!isBuildGood) { // Prevent the encrypted binary form leaking out - clBinary()->setBinary(NULL, 0); + clBinary()->setBinary(nullptr, 0); } @@ -225,24 +225,24 @@ HSAILProgram::getCompilationStagesFromBinary(std::vector& completeStage bool containsShaderIsa = (type() == TYPE_EXECUTABLE); bool containsOpts = !(compileOptions_.empty() && linkOptions_.empty()); #if !defined(WITH_LIGHTNING_COMPILER) // !defined(WITH_LIGHTNING_COMPILER) - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, NULL, &containsLlvmirText, &boolSize); + errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_LLVMIR, nullptr, &containsLlvmirText, &boolSize); if (errorCode != ACL_SUCCESS) { containsLlvmirText = false; } // Checking compile & link options in .comment section - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_OPTIONS, NULL, &containsOpts, &boolSize); + errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_OPTIONS, nullptr, &containsOpts, &boolSize); if (errorCode != ACL_SUCCESS) { containsOpts = false; } // Checking HSAIL in .cg section containsHsailText = true; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_HSAIL, NULL, &containsHsailText, &boolSize); + errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_HSAIL, nullptr, &containsHsailText, &boolSize); if (errorCode != ACL_SUCCESS) { containsHsailText = false; } // Checking BRIG sections containsBrig = true; - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_BRIG, NULL, &containsBrig, &boolSize); + errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_BRIG, nullptr, &containsBrig, &boolSize); if (errorCode != ACL_SUCCESS) { containsBrig = false; } @@ -264,7 +264,7 @@ HSAILProgram::getCompilationStagesFromBinary(std::vector& completeStage completeStages.push_back(from); from = ACL_TYPE_HSAIL_TEXT; } - errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_ISA, NULL, &containsShaderIsa, &boolSize); + errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, RT_CONTAINS_ISA, nullptr, &containsShaderIsa, &boolSize); if (errorCode != ACL_SUCCESS) { containsShaderIsa = false; } @@ -328,7 +328,7 @@ HSAILProgram::getNextCompilationStageFromBinary(amd::option::Options* options) aclType continueCompileFrom = ACL_TYPE_DEFAULT; binary_t binary = this->binary(); // If the binary already exists - if ((binary.first != NULL) && (binary.second > 0)) { + if ((binary.first != nullptr) && (binary.second > 0)) { #if defined(WITH_LIGHTNING_COMPILER) void *mem = (void *) binary.first; #else // !defined(WITH_LIGHTNING_COMPILER) @@ -443,7 +443,7 @@ HSAILProgram::saveBinaryAndSetType(type_t type, void* rawBinary, size_t size) //Write binary to memory #if defined(WITH_LIGHTNING_COMPILER) if (type == TYPE_EXECUTABLE) { // handle code object binary - assert(rawBinary != NULL && size != 0 && "must pass in the binary"); + assert(rawBinary != nullptr && size != 0 && "must pass in the binary"); } else { // handle LLVM binary if (llvmBinary_.empty()) { @@ -484,7 +484,7 @@ HSAILProgram::linkImpl_LC( std::vector inputs; for (auto program : (const std::vector&)inputPrograms) { if (program->llvmBinary_.empty()) { - if (program->clBinary() == NULL) { + if (program->clBinary() == nullptr) { buildLog_ += "Internal error: Input program not compiled!\n"; return false; } @@ -591,7 +591,7 @@ HSAILProgram::linkImpl( HSAILProgram *program = (HSAILProgram *)*it; // Check if the program was created with clCreateProgramWIthBinary binary_t binary = program->binary(); - if ((binary.first != NULL) && (binary.second > 0)) { + if ((binary.first != nullptr) && (binary.second > 0)) { // Binary already exists -- we can also check if there is no // opencl source code // Need to check if LLVMIR exists in the binary @@ -614,7 +614,7 @@ HSAILProgram::linkImpl( size_t boolSize = sizeof(bool); bool containsLLLVMIR = false; errorCode = g_complibApi._aclQueryInfo(device().compiler(), binaryElf_, - RT_CONTAINS_LLVMIR, NULL, &containsLLLVMIR, &boolSize); + RT_CONTAINS_LLVMIR, nullptr, &containsLLLVMIR, &boolSize); if (errorCode != ACL_SUCCESS || !containsLLLVMIR) { buildLog_ +="Error while linking : Invalid binary (Missing LLVMIR section)"; return false; @@ -635,16 +635,16 @@ HSAILProgram::linkImpl( &binaries_to_link[1], ACL_TYPE_LLVMIR_BINARY, "-create-library", - NULL); + nullptr); } else { errorCode = g_complibApi._aclLink(device().compiler(), binaries_to_link[0], 0, - NULL, + nullptr, ACL_TYPE_LLVMIR_BINARY, "-create-library", - NULL); + nullptr); } if (errorCode != ACL_SUCCESS) { buildLog_ += "Failed to link programs"; @@ -669,7 +669,7 @@ HSAILProgram::linkImpl( static inline const char* hsa_strerror(hsa_status_t status) { - const char* str = NULL; + const char* str = nullptr; if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) { return str; } @@ -865,7 +865,7 @@ HSAILProgram::setKernels_LC(amd::option::Options *options, void* binary, size_t status = hsa_executable_create_alt( HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, - NULL, &hsaExecutable_ ); + nullptr, &hsaExecutable_ ); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Executable for AMD HSA Code Object isn't created: "; buildLog_ += hsa_strerror(status); @@ -885,7 +885,7 @@ HSAILProgram::setKernels_LC(amd::option::Options *options, void* binary, size_t } status = hsa_executable_load_agent_code_object( - hsaExecutable_, agent, codeObjectReader, NULL, NULL ); + hsaExecutable_, agent, codeObjectReader, nullptr, nullptr ); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: AMD HSA Code Object loading failed: "; buildLog_ += hsa_strerror(status); @@ -896,7 +896,7 @@ HSAILProgram::setKernels_LC(amd::option::Options *options, void* binary, size_t hsa_code_object_reader_destroy(codeObjectReader); // Freeze the executable. - status = hsa_executable_freeze( hsaExecutable_, NULL ); + status = hsa_executable_freeze( hsaExecutable_, nullptr ); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Freezing the executable failed: "; buildLog_ += hsa_strerror(status); @@ -909,7 +909,7 @@ HSAILProgram::setKernels_LC(amd::option::Options *options, void* binary, size_t size_t progvarsWriteSize = 0; // Begin the Elf image from memory - Elf* e = elf_memory((char*) binary, binSize, NULL); + Elf* e = elf_memory((char*) binary, binSize, nullptr); if (elf_kind(e) != ELF_K_ELF) { buildLog_ += "Error while reading the ELF program binary\n"; return false; @@ -1144,7 +1144,7 @@ HSAILProgram::linkImpl(amd::option::Options *options) case ACL_TYPE_ISA: { #if defined(WITH_LIGHTNING_COMPILER) binary_t isaBinary = binary(); - if ((isaBinary.first != NULL) && (isaBinary.second > 0)) { + if ((isaBinary.first != nullptr) && (isaBinary.second > 0)) { return setKernels_LC(options, (void*) isaBinary.first, isaBinary.second ); } else { @@ -1192,7 +1192,7 @@ HSAILProgram::linkImpl(amd::option::Options *options) // Create an executable. hsa_status_t status = hsa_executable_create_alt( HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, - NULL, &hsaExecutable_ + nullptr, &hsaExecutable_ ); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Failed to create executable: "; @@ -1213,7 +1213,7 @@ HSAILProgram::linkImpl(amd::option::Options *options) } status = hsa_executable_load_agent_code_object( - hsaExecutable_, hsaDevice, codeObjectReader, NULL, NULL ); + hsaExecutable_, hsaDevice, codeObjectReader, nullptr, nullptr ); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: AMD HSA Code Object loading failed: "; buildLog_ += hsa_strerror(status); @@ -1224,7 +1224,7 @@ HSAILProgram::linkImpl(amd::option::Options *options) hsa_code_object_reader_destroy(codeObjectReader); // Freeze the executable. - status = hsa_executable_freeze(hsaExecutable_, NULL); + status = hsa_executable_freeze(hsaExecutable_, nullptr); if (status != HSA_STATUS_SUCCESS) { buildLog_ += "Error: Failed to freeze executable: "; buildLog_ += hsa_strerror(status); @@ -1369,9 +1369,9 @@ HSAILProgram::createBinary(amd::option::Options *options) bool HSAILProgram::initClBinary() { - if (clBinary_ == NULL) { + if (clBinary_ == nullptr) { clBinary_ = new ClBinary(static_cast(device())); - if (clBinary_ == NULL) { + if (clBinary_ == nullptr) { return false; } } @@ -1381,9 +1381,9 @@ HSAILProgram::initClBinary() void HSAILProgram::releaseClBinary() { - if (clBinary_ != NULL) { + if (clBinary_ != nullptr) { delete clBinary_; - clBinary_ = NULL; + clBinary_ = nullptr; } } diff --git a/rocclr/runtime/device/rocm/rocprogram.hpp b/rocclr/runtime/device/rocm/rocprogram.hpp index 74a5e3fa68..137b12917f 100644 --- a/rocclr/runtime/device/rocm/rocprogram.hpp +++ b/rocclr/runtime/device/rocm/rocprogram.hpp @@ -147,7 +147,7 @@ private: /* \brief Returns the next stage to compile from, based on sections and options in binary */ aclType getNextCompilationStageFromBinary(amd::option::Options* options); - bool saveBinaryAndSetType(type_t type, void* binary = NULL, size_t size = 0); + bool saveBinaryAndSetType(type_t type, void* binary = nullptr, size_t size = 0); //! Disable default copy constructor HSAILProgram(const HSAILProgram&) = delete; diff --git a/rocclr/runtime/device/rocm/rocsettings.cpp b/rocclr/runtime/device/rocm/rocsettings.cpp index 09c299645b..98e585ce5e 100644 --- a/rocclr/runtime/device/rocm/rocsettings.cpp +++ b/rocclr/runtime/device/rocm/rocsettings.cpp @@ -38,7 +38,7 @@ Settings::Settings() // operates or is programmed to be in Coherent mode. // Users can turn it off for hardware that does not // support this feature naturally - char *nonCoherentMode = NULL; + char *nonCoherentMode = nullptr; nonCoherentMode = getenv("OPENCL_USE_NC_MEMORY_POLICY"); enableNCMode_ = (nonCoherentMode)? true : false; @@ -48,7 +48,7 @@ Settings::Settings() // devices that do not support this feature. // // @note Update appropriate field of device::Settings - char *partialDispatch = NULL; + char *partialDispatch = nullptr; partialDispatch = getenv("OPENCL_DISABLE_PARTIAL_DISPATCH"); enablePartialDispatch_ = (partialDispatch) ? false : true; partialDispatch_ = (partialDispatch) ? false : true; diff --git a/rocclr/runtime/device/rocm/rocvirtual.cpp b/rocclr/runtime/device/rocm/rocvirtual.cpp index 543bcc7872..490f376e47 100644 --- a/rocclr/runtime/device/rocm/rocvirtual.cpp +++ b/rocclr/runtime/device/rocm/rocvirtual.cpp @@ -105,7 +105,7 @@ VirtualGPU::MemoryDependency::create(size_t numMemObj) if (numMemObj > 0) { // Allocate the array of memory objects for dependency tracking memObjectsInQueue_ = new MemoryState[numMemObj]; - if (NULL == memObjectsInQueue_) { + if (nullptr == memObjectsInQueue_) { return false; } memset(memObjectsInQueue_, 0, sizeof(MemoryState) * numMemObj); @@ -241,13 +241,13 @@ VirtualGPU::processMemObjects( size_t execInfoOffset = kernelParams.getExecInfoOffset(); bool sync = true; - amd::Memory* memory = NULL; + amd::Memory* memory = nullptr; //get svm non arugment information void* const* svmPtrArray = reinterpret_cast(params + execInfoOffset); for (size_t i = 0; i < count; i++) { memory = amd::SvmManager::FindSvmBuffer(svmPtrArray[i]); - if (NULL == memory) { + if (nullptr == memory) { if (!supportFineGrainedSystem) { return false; } @@ -262,7 +262,7 @@ VirtualGPU::processMemObjects( } else { Memory* rocMemory = static_cast(memory->getDeviceMemory(dev())); - if (NULL != rocMemory) { + if (nullptr != rocMemory) { // Synchronize data with other memory instances if necessary rocMemory->syncCacheFromHost(*this); @@ -280,9 +280,9 @@ VirtualGPU::processMemObjects( for (size_t i = 0; i < signature.numParameters(); ++i) { const amd::KernelParameterDescriptor& desc = signature.at(i); const Kernel::Argument* arg = hsaKernel.hsailArgAt(i); - Memory* memory = NULL; + Memory* memory = nullptr; bool readOnly = false; - amd::Memory* svmMem = NULL; + amd::Memory* svmMem = nullptr; // Find if current argument is a buffer if ((desc.type_ == T_POINTER) && (arg->addrQual_ != ROC_ADDRESS_LOCAL)) { @@ -300,8 +300,8 @@ VirtualGPU::processMemObjects( } if (*reinterpret_cast - (params + desc.offset_) != NULL) { - if (NULL == svmMem) { + (params + desc.offset_) != nullptr) { + if (nullptr == svmMem) { memory = static_cast((*reinterpret_cast (params + desc.offset_))->getDeviceMemory(dev())); } @@ -316,7 +316,7 @@ VirtualGPU::processMemObjects( } } - if (memory != NULL) { + if (memory != nullptr) { readOnly |= (arg->access_ == ROC_ACCESS_TYPE_RO); // Validate memory for a dependency in the queue memoryDependency().validate(*this, memory, readOnly); @@ -473,14 +473,14 @@ VirtualGPU::VirtualGPU(Device &device) , index_(device.numOfVgpus_++) // Virtual gpu unique index incrementing { gpu_device_ = device.getBackendDevice(); - printfdbg_ = NULL; + printfdbg_ = nullptr; // Initialize the last signal and dispatch flags - timestamp_ = NULL; + timestamp_ = nullptr; hasPendingDispatch_ = false; - tools_lib_ = NULL; + tools_lib_ = nullptr; - kernarg_pool_base_ = NULL; + kernarg_pool_base_ = nullptr; kernarg_pool_size_ = 0; kernarg_pool_cur_offset_ = 0; aqlHeader_ = kDispatchPacketHeaderNoSync; @@ -491,17 +491,17 @@ VirtualGPU::~VirtualGPU() { releasePinnedMem(); - if (timestamp_ != NULL) { + if (timestamp_ != nullptr) { delete timestamp_; - timestamp_ = NULL; + timestamp_ = nullptr; LogError("There was a timestamp that was not used; deleting."); } - if (printfdbg_ != NULL){ + if (printfdbg_ != nullptr){ delete printfdbg_; - printfdbg_ = NULL; + printfdbg_ = nullptr; } - tools_lib_ = NULL; + tools_lib_ = nullptr; --roc_device_.numOfVgpus_; // Virtual gpu unique index decrementing } @@ -536,7 +536,7 @@ VirtualGPU::create(bool profilingEna) uint32_t queue_size = 1024; queue_size = (queue_max_packets < queue_size) ? queue_max_packets : queue_size; while (hsa_queue_create(gpu_device_, - queue_size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, UINT_MAX, UINT_MAX, + queue_size, HSA_QUEUE_TYPE_MULTI, nullptr, nullptr, UINT_MAX, UINT_MAX, &gpu_queue_) != HSA_STATUS_SUCCESS) { queue_size >>= 1; if (queue_size < 64) { @@ -551,7 +551,7 @@ VirtualGPU::create(bool profilingEna) device::BlitManager::Setup blitSetup; blitMgr_ = new KernelBlitManager(*this, blitSetup); - if ((NULL == blitMgr_) || !blitMgr_->create(roc_device_)) { + if ((nullptr == blitMgr_) || !blitMgr_->create(roc_device_)) { LogError("Could not create BlitManager!"); return false; } @@ -559,7 +559,7 @@ VirtualGPU::create(bool profilingEna) // Create signal for the barrier packet. hsa_signal_t signal = { 0 }; if (HSA_STATUS_SUCCESS != - hsa_signal_create(InitSignalValue, 0, NULL, &signal)) { + hsa_signal_create(InitSignalValue, 0, nullptr, &signal)) { return false; } barrier_signal_ = signal; @@ -571,7 +571,7 @@ VirtualGPU::create(bool profilingEna) // Create a object of PrintfDbg printfdbg_ = new PrintfDbg(roc_device_); - if (NULL == printfdbg_) { + if (nullptr == printfdbg_) { LogError("\nCould not create printfDbg Object!"); return false; } @@ -609,7 +609,7 @@ VirtualGPU::terminate() if (tools_lib_) { amd::Os::unloadLibrary(tools_lib_); - tools_lib_ = NULL; + tools_lib_ = nullptr; } destroyPool(); @@ -696,7 +696,7 @@ VirtualGPU::allocKernArg(size_t size, size_t alignment) void VirtualGPU::profilingBegin(amd::Command &command, bool drmProfiling) { if (command.profilingInfo().enabled_) { - if (timestamp_ != NULL) { + if (timestamp_ != nullptr) { LogWarning("Trying to create a second timestamp in VirtualGPU. \ This could have unintended consequences."); return; @@ -718,7 +718,7 @@ void VirtualGPU::profilingEnd(amd::Command &command) timestamp_->end(); } command.setData(reinterpret_cast(timestamp_)); - timestamp_ = NULL; + timestamp_ = nullptr; } } @@ -733,12 +733,12 @@ struct DestroySampler : public std::binary_functiondata() != NULL) { + while (current != nullptr) { + if (current->data() != nullptr) { ts = reinterpret_cast(current->data()); startTimeStamp = ts->getStart(); endTimeStamp = ts->getStart(); @@ -777,16 +777,16 @@ void VirtualGPU::updateCommandsState(amd::Command *list) // with the COMPLETE (end) timestamp of the previous command, A. This is // also true for any command B, which falls between A and C. current = list; - while (current != NULL) { + while (current != nullptr) { if (current->profilingInfo().enabled_) { - if (current->data() != NULL) { + if (current->data() != nullptr) { // Since this is a valid command to get a timestamp, we use the // timestamp provided by the runtime (saved in the data()) ts = reinterpret_cast(current->data()); startTimeStamp = ts->getStart(); endTimeStamp = ts->getEnd(); delete ts; - current->setData(NULL); + current->setData(nullptr); } else { // If we don't have a command that contains a valid timestamp, @@ -990,7 +990,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) profilingBegin(cmd); const std::vector& svmPointers = cmd.svmPointers(); - if (cmd.pfnFreeFunc() == NULL) { + if (cmd.pfnFreeFunc() == nullptr) { // pointers allocated using clSVMAlloc for (cl_uint i = 0; i < svmPointers.size(); i++) { amd::SvmBuffer::free(cmd.context(), svmPointers[i]); @@ -1179,7 +1179,7 @@ void VirtualGPU::submitMapMemory(amd::MapMemoryCommand &cmd) roc::Memory *hsaMemory = static_cast(devMemory); amd::Memory* mapMemory = hsaMemory->mapMemory(); - void *hostPtr = mapMemory == NULL ? + void *hostPtr = mapMemory == nullptr ? hsaMemory->owner()->getHostMem() : mapMemory->getHostMem(); @@ -1603,14 +1603,14 @@ VirtualGPU::submitKernelInternal( gpuKernel.KernargSegmentByteSize(), gpuKernel.KernargSegmentAlignment()); - if (argBuffer == NULL) { + if (argBuffer == nullptr) { LogError("Out of memory"); return false; } address argPtr = argBuffer; for (auto arg : gpuKernel.hsailArgs()) { - const_address srcArgPtr = NULL; + const_address srcArgPtr = nullptr; if (arg->index_ != uint(-1)) { srcArgPtr = parameters + signature.at(arg->index_).offset_; } @@ -1666,7 +1666,7 @@ VirtualGPU::submitKernelInternal( break; } amd::Memory* mem = *reinterpret_cast(srcArgPtr); - if (mem == NULL) { + if (mem == nullptr) { argPtr = addArg(argPtr, srcArgPtr, arg->size_, arg->alignment_); break; } @@ -1684,7 +1684,7 @@ VirtualGPU::submitKernelInternal( } case ROC_ARGTYPE_REFERENCE: { void *mem = allocKernArg(arg->size_, arg->alignment_); - if (mem == NULL) { + if (mem == nullptr) { LogError("Out of memory"); return false; } @@ -1698,7 +1698,7 @@ VirtualGPU::submitKernelInternal( case ROC_ARGTYPE_IMAGE: { amd::Memory* mem = *reinterpret_cast(srcArgPtr); Image* image = static_cast(mem->getDeviceMemory(dev())); - if (image == NULL) { + if (image == nullptr) { LogError("Kernel image argument is not an image object"); return false; } @@ -1722,7 +1722,7 @@ VirtualGPU::submitKernelInternal( } case ROC_ARGTYPE_SAMPLER: { amd::Sampler* sampler = *reinterpret_cast(srcArgPtr); - if (sampler == NULL) { + if (sampler == nullptr) { LogError("Kernel sampler argument is not an sampler object"); return false; } diff --git a/rocclr/runtime/device/rocm/rocvirtual.hpp b/rocclr/runtime/device/rocm/rocvirtual.hpp index b06c9428cc..64f4ace5e3 100644 --- a/rocclr/runtime/device/rocm/rocvirtual.hpp +++ b/rocclr/runtime/device/rocm/rocvirtual.hpp @@ -115,7 +115,7 @@ public: public: //! Default constructor MemoryDependency() - : memObjectsInQueue_(NULL) + : memObjectsInQueue_(nullptr) , numMemObjectsInQueue_(0) , maxMemObjectsInQueue_(0) {} @@ -177,7 +177,7 @@ public: void submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& cmd); void submitPerfCounter(amd::PerfCounterCommand& cmd){}; - void flush(amd::Command* list = NULL, bool wait = false); + void flush(amd::Command* list = nullptr, bool wait = false); void submitFillMemory(amd::FillMemoryCommand& cmd); void submitMigrateMemObjects(amd::MigrateMemObjectsCommand& cmd);