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<const char*>(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<const char*>(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<address>(dst) + dstOffset), dstAgent,
                     (reinterpret_cast<const_address>(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<amd::Memory>(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<roc::Settings*>(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<amd::Memory*>();
-    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<roc::Settings*>(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<const Device&>(*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<roc::Image&>(*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 &>(*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<char *>(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<address>(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<roc::Memory *>(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<const Image&>(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<address>(
         dev().hostAlloc(dbgBuffer_size_, sizeof(void*)));
@@ -38,7 +38,7 @@ bool PrintfDbg::allocate(bool realloc) {
     dbgBuffer_ = reinterpret_cast<address>(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<uint32_t*>(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<uint32_t*>(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<aclType>& 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<aclType>& 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<Data*> inputs;
     for (auto program : (const std::vector<HSAILProgram*>&)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<const Device &>(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<void* const*>(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*>(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<amd::Memory* const*>
-                    (params + desc.offset_) != NULL) {
-                if (NULL == svmMem) {
+                    (params + desc.offset_) != nullptr) {
+                if (nullptr == svmMem) {
                     memory = static_cast<Memory*>((*reinterpret_cast<amd::Memory* const*>
                             (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<void*>(timestamp_));
-        timestamp_ = NULL;
+        timestamp_ = nullptr;
     }
 }
 
@@ -733,12 +733,12 @@ struct DestroySampler : public std::binary_function<hsa_ext_sampler_t,
 
 void VirtualGPU::updateCommandsState(amd::Command *list)
 {
-    Timestamp *ts = NULL;
+    Timestamp *ts = nullptr;
 
     amd::Command* current = list;
-    amd::Command* next = NULL;
+    amd::Command* next = nullptr;
 
-    if (current == NULL) {
+    if (current == nullptr) {
         return;
     }
 
@@ -754,8 +754,8 @@ void VirtualGPU::updateCommandsState(amd::Command *list)
         // that has one. This timestamp is used below to mark any command that
         // came before it to start and end with this first valid start time.
         current = list;
-        while (current != NULL) {
-            if (current->data() != NULL) {
+        while (current != nullptr) {
+            if (current->data() != nullptr) {
                 ts = reinterpret_cast<Timestamp*>(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<Timestamp*>(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<void*>& 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<roc::Memory *>(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<amd::Memory* const*>(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<amd::Memory* const*>(srcArgPtr);
                 Image* image = static_cast<Image *>(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<amd::Sampler* const*>(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);