replace "report_fatal_error" in IGC/Compiler by IGC_ASSERT_EXIT

lukaszgotszaldintel · sys_zuul · commit 031c5c6e8683 · 2020-05-18T13:26:07.000-07:00
Change-Id: I4b30478676af1bfbfcefae9b43b35ee7252cef46
diff --git a/IGC/Compiler/Optimizer/OpenCLPasses/DeviceEnqueueFuncs/TransformBlocks.cpp b/IGC/Compiler/Optimizer/OpenCLPasses/DeviceEnqueueFuncs/TransformBlocks.cpp
@@ -718,7 +718,7 @@ namespace //Anonymous
             , _waitEventsList(nullptr)
             , _retEvent(nullptr)
         {
-            if (_calledFunc == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != _calledFunc, "indirect calls are not supported");
         }
 
     public:
@@ -746,7 +746,7 @@ namespace //Anonymous
         virtual llvm::Function* getEnqueuedFunction()
         {
             auto invokeFunc = dyn_cast<llvm::Function>(getInvokeArg()->stripPointerCasts());
-            if (invokeFunc == nullptr) report_fatal_error("Device enqueue invoke param is not a function");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != invokeFunc, "Device enqueue invoke param is not a function");
             return invokeFunc;
         }
 
@@ -823,7 +823,7 @@ namespace //Anonymous
                 }
                 else
                 {
-                    report_fatal_error("enqueue_kernel signature does not match");
+                    IGC_ASSERT_EXIT_MESSAGE(0, "enqueue_kernel signature does not match");
                 }
             }
         }
@@ -884,7 +884,7 @@ namespace //Anonymous
         SPIRVOpEnqueueKernelCallArgs(llvm::CallInst& call, DataContext& dataContext)
             : DeviceExecCallArgs(call, dataContext)
         {
-            if (_call.getNumArgOperands() < 8) report_fatal_error("OpEnqueueKernel signature does not match");
+            IGC_ASSERT_EXIT_MESSAGE(8 <= _call.getNumArgOperands(), "OpEnqueueKernel signature does not match");
 
             _queue = _call.getArgOperand(0);
             _flags = _call.getArgOperand(1);
@@ -908,19 +908,23 @@ namespace //Anonymous
             for (unsigned i = localSizesStartArgNum; i < argsNum; i++)
             {
                 auto arg = _call.getArgOperand(i);
+                IGC_ASSERT(nullptr != arg);
+                IGC_ASSERT(nullptr != arg->getType());
+
                 if (arg->getType()->isPointerTy()) {
                     IRBuilder<> builder(&_call);
                     arg = builder.CreateLoad(arg);
                 }
 
-                if (!arg->getType()->isIntegerTy(64) && !arg->getType()->isIntegerTy(32))
-                    report_fatal_error("OpEnqueueKernel signature does not match");
-
                 if (arg->getType()->isIntegerTy(64))
                 {
                     Type* int32Ty = Type::getInt32Ty(_call.getParent()->getContext());
                     arg = CastInst::CreateTruncOrBitCast(arg, int32Ty, "", &_call);
                 }
+                else
+                {
+                    IGC_ASSERT_EXIT_MESSAGE(arg->getType()->isIntegerTy(32), "OpEnqueueKernel signature does not match");
+                }
                 _local_sizes.push_back(arg);
             }
         }
@@ -964,11 +968,11 @@ namespace //Anonymous
         llvm::Value* AdjustNDRangeType(llvm::Value* nDRange)
         {
             auto expectedNdrangeTy = getNDRangeType();
-
+            IGC_ASSERT(nullptr != nDRange);
+            IGC_ASSERT(nullptr != nDRange->getType());
             IGC_ASSERT(nDRange->getType()->isPointerTy());
             auto actualNDRangeTy = KindQuery::toStructType(nDRange->getType());
-            IGC_ASSERT(actualNDRangeTy);
-            if (actualNDRangeTy == nullptr) report_fatal_error("NDRange type mismatch");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != actualNDRangeTy, "NDRange type mismatch");
 
             if (actualNDRangeTy->isLayoutIdentical(expectedNdrangeTy))
             {
@@ -1016,7 +1020,7 @@ namespace //Anonymous
 
             const auto newName = "__builtin_IB_get_block_simd_size";
             auto calledFunction = _deviceExecCall->getCalledFunction();
-            if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");
             return CreateNewCall(newName, calledFunction->getReturnType(), blockIdValue);
         }
     };
@@ -1031,7 +1035,7 @@ namespace //Anonymous
         {
             const auto newName = "__builtin_IB_get_max_workgroup_size";
             auto calledFunction = _deviceExecCall->getCalledFunction();
-            if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");
             return CreateNewCall(newName, calledFunction->getReturnType(), {});
         }
     };
@@ -1057,7 +1061,7 @@ namespace //Anonymous
 
             const auto newName = "IGIL_calc_sub_group_count_for_ndrange";
             auto calledFunction = _deviceExecCall->getCalledFunction();
-            if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");
             return CreateNewCall(newName, calledFunction->getReturnType(), args);
         }
     };
@@ -1082,7 +1086,7 @@ namespace //Anonymous
 
             const auto newName = "__intel_calc_kernel_local_size_for_sub_group_count";
             auto calledFunction = _deviceExecCall->getCalledFunction();
-            if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");
             return CreateNewCall(newName, calledFunction->getReturnType(), args);
         }
     };
@@ -1105,7 +1109,7 @@ namespace //Anonymous
 
             const auto newName = "__intel_calc_kernel_max_num_subgroups";
             auto calledFunction = _deviceExecCall->getCalledFunction();
-            if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");
             return CreateNewCall(newName, calledFunction->getReturnType(), args);
         }
     };
@@ -1791,7 +1795,7 @@ namespace //Anonymous
 
         // get list of calls which "enqueue" this dispatcher
         auto parentCalls = _dataContext.getCallHandlersFor(invokeFunc);
-        if (parentCalls.size() == 0) report_fatal_error("parent calls for invoke are not set");
+        IGC_ASSERT_EXIT_MESSAGE(parentCalls.size(), "parent calls for invoke are not set");
 
         // lookup for a call which located in a "__kernel"
         // if __kernel call will not found, use the one which is not self
@@ -1810,7 +1814,7 @@ namespace //Anonymous
             }
         }
 
-        if (parentCallHandler == nullptr) report_fatal_error("Fail parent call lookup: possible closed self-enqueue");
+        IGC_ASSERT_EXIT_MESSAGE(nullptr != parentCallHandler, "Fail parent call lookup: possible closed self-enqueue");
 
         auto capturedValues = parentCallHandler->getArgs()->getParamValue()->getCapturedValues(_dataContext.getDispatcherForInvokeFunc(invokeFunc)->getCaptureIndicies());
         auto& capture = capturedValues[captureNum];
@@ -1925,8 +1929,7 @@ namespace //Anonymous
         if (!_invocations[invokeFunc]._dispatcher)
         {
             const Function* parentKernel = getParentKernel(invokeFunc);
-            if (!parentKernel)
-                report_fatal_error("Fail parent kernel lookup: possible closed self-enqueue");
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != parentKernel, "Fail parent kernel lookup: possible closed self-enqueue");
 
             _invocations[invokeFunc]._dispatcher.reset(new Dispatcher(invokeFunc, parentKernel->getName(), _blocksNum++));
         }
@@ -1960,7 +1963,7 @@ namespace //Anonymous
 
     const llvm::Function* DataContext::getParentKernelImpl(const llvm::Function* invokeFunc, std::set<const llvm::Function*>& processedFuncs)
     {
-        if (invokeFunc == nullptr) report_fatal_error("invoke function should not be null");
+        IGC_ASSERT_EXIT_MESSAGE(nullptr != invokeFunc, "invoke function should not be null");
 
         processedFuncs.insert(invokeFunc);
 
@@ -2383,7 +2386,7 @@ namespace //Anonymous
 
     Capture::Capture(llvm::Value* val, DataContext& context)
     {
-        if (val == nullptr) report_fatal_error("captured value is null");
+        IGC_ASSERT_EXIT_MESSAGE(nullptr != val, "captured value is null");
         value = val;
         if (auto arg = context.getArgForValue(value))
         {
@@ -2436,7 +2439,7 @@ namespace //Anonymous
         , _block_invokeFunction(nullptr)
         , _dataContext(dataContext)
     {
-        if (!_paramStruct) report_fatal_error("Enqueue param is not a struct");
+        IGC_ASSERT_EXIT_MESSAGE(nullptr != _paramStruct, "Enqueue param is not a struct");
 
         if (_dataContext.getKindQuery().isBlockStructType(_paramStruct->getType())) {
             /// ObjectiveC block value passed to OCL "device execution" call.
@@ -2455,12 +2458,9 @@ namespace //Anonymous
     {
         if (_paramValue == nullptr)
         {
-            if (auto param = getParamArg()) {
-                _paramValue = _dataContext.getDeviceEnqueueParamValue(param->stripPointerCasts());
-            }
-            else {
-                report_fatal_error("Enqueue param is not set");
-            }
+            auto param = getParamArg();
+            IGC_ASSERT_EXIT_MESSAGE(nullptr != param, "Enqueue param is not set");
+            _paramValue = _dataContext.getDeviceEnqueueParamValue(param->stripPointerCasts());
         }
         return _paramValue;
     }
@@ -2610,7 +2610,7 @@ namespace //Anonymous
             case CaptureKind::IMAGE:
             case CaptureKind::SAMPLER:
             {
-                if (Capture::ARG_NUM_NONE == capturedValue.argNum) report_fatal_error("unknown argument number for an object");
+                IGC_ASSERT_EXIT_MESSAGE(Capture::ARG_NUM_NONE != capturedValue.argNum, "unknown argument number for an object");
                 auto objArgNumValue = llvm::ConstantInt::get(int32ty, capturedValue.argNum);
                 auto getObjIDFunc = getOrCreateFunc("__builtin_IB_get_object_id", int32ty, int32ty);
                 auto objIdValue = builder.CreateCall(getObjIDFunc, objArgNumValue, "obj_id");
diff --git a/IGC/Compiler/Optimizer/OpenCLPasses/GenericAddressResolution/GenericAddressDynamicResolution.cpp b/IGC/Compiler/Optimizer/OpenCLPasses/GenericAddressResolution/GenericAddressDynamicResolution.cpp
@@ -220,7 +220,7 @@ bool GenericAddressDynamicResolution::visitLoadStoreInst(Instruction& I)
         pointerAddressSpace = store->getPointerAddressSpace();
     }
     else {
-        report_fatal_error("Unable to resolve generic address space pointer");
+        IGC_ASSERT_EXIT_MESSAGE(0, "Unable to resolve generic address space pointer");
     }
 
     if (pointerAddressSpace == ADDRESS_SPACE_GENERIC) {

Original file line number	Diff line number	Diff line change
`@@ -718,7 +718,7 @@ namespace //Anonymous`
`718`	`718`	`, _waitEventsList(nullptr)`
`719`	`719`	`, _retEvent(nullptr)`
`720`	`720`	`{`
`721`		`- if (_calledFunc == nullptr) report_fatal_error("indirect calls are not supported");`
	`721`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != _calledFunc, "indirect calls are not supported");`
`722`	`722`	`}`
`723`	`723`
`724`	`724`	`public:`
`@@ -746,7 +746,7 @@ namespace //Anonymous`
`746`	`746`	`virtual llvm::Function* getEnqueuedFunction()`
`747`	`747`	`{`
`748`	`748`	`auto invokeFunc = dyn_cast<llvm::Function>(getInvokeArg()->stripPointerCasts());`
`749`		`- if (invokeFunc == nullptr) report_fatal_error("Device enqueue invoke param is not a function");`
	`749`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != invokeFunc, "Device enqueue invoke param is not a function");`
`750`	`750`	`return invokeFunc;`
`751`	`751`	`}`
`752`	`752`
`@@ -823,7 +823,7 @@ namespace //Anonymous`
`823`	`823`	`}`
`824`	`824`	`else`
`825`	`825`	`{`
`826`		`- report_fatal_error("enqueue_kernel signature does not match");`
	`826`	`+ IGC_ASSERT_EXIT_MESSAGE(0, "enqueue_kernel signature does not match");`
`827`	`827`	`}`
`828`	`828`	`}`
`829`	`829`	`}`
`@@ -884,7 +884,7 @@ namespace //Anonymous`
`884`	`884`	`SPIRVOpEnqueueKernelCallArgs(llvm::CallInst& call, DataContext& dataContext)`
`885`	`885`	`: DeviceExecCallArgs(call, dataContext)`
`886`	`886`	`{`
`887`		`- if (_call.getNumArgOperands() < 8) report_fatal_error("OpEnqueueKernel signature does not match");`
	`887`	`+ IGC_ASSERT_EXIT_MESSAGE(8 <= _call.getNumArgOperands(), "OpEnqueueKernel signature does not match");`
`888`	`888`
`889`	`889`	`_queue = _call.getArgOperand(0);`
`890`	`890`	`_flags = _call.getArgOperand(1);`
`@@ -908,19 +908,23 @@ namespace //Anonymous`
`908`	`908`	`for (unsigned i = localSizesStartArgNum; i < argsNum; i++)`
`909`	`909`	`{`
`910`	`910`	`auto arg = _call.getArgOperand(i);`
	`911`	`+ IGC_ASSERT(nullptr != arg);`
	`912`	`+ IGC_ASSERT(nullptr != arg->getType());`
	`913`	`+`
`911`	`914`	`if (arg->getType()->isPointerTy()) {`
`912`	`915`	`IRBuilder<> builder(&_call);`
`913`	`916`	`arg = builder.CreateLoad(arg);`
`914`	`917`	`}`
`915`	`918`
`916`		`- if (!arg->getType()->isIntegerTy(64) && !arg->getType()->isIntegerTy(32))`
`917`		`- report_fatal_error("OpEnqueueKernel signature does not match");`
`918`		`-`
`919`	`919`	`if (arg->getType()->isIntegerTy(64))`
`920`	`920`	`{`
`921`	`921`	`Type* int32Ty = Type::getInt32Ty(_call.getParent()->getContext());`
`922`	`922`	`arg = CastInst::CreateTruncOrBitCast(arg, int32Ty, "", &_call);`
`923`	`923`	`}`
	`924`	`+ else`
	`925`	`+ {`
	`926`	`+ IGC_ASSERT_EXIT_MESSAGE(arg->getType()->isIntegerTy(32), "OpEnqueueKernel signature does not match");`
	`927`	`+ }`
`924`	`928`	`_local_sizes.push_back(arg);`
`925`	`929`	`}`
`926`	`930`	`}`
`@@ -964,11 +968,11 @@ namespace //Anonymous`
`964`	`968`	`llvm::Value* AdjustNDRangeType(llvm::Value* nDRange)`
`965`	`969`	`{`
`966`	`970`	`auto expectedNdrangeTy = getNDRangeType();`
`967`		`-`
	`971`	`+ IGC_ASSERT(nullptr != nDRange);`
	`972`	`+ IGC_ASSERT(nullptr != nDRange->getType());`
`968`	`973`	`IGC_ASSERT(nDRange->getType()->isPointerTy());`
`969`	`974`	`auto actualNDRangeTy = KindQuery::toStructType(nDRange->getType());`
`970`		`- IGC_ASSERT(actualNDRangeTy);`
`971`		`- if (actualNDRangeTy == nullptr) report_fatal_error("NDRange type mismatch");`
	`975`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != actualNDRangeTy, "NDRange type mismatch");`
`972`	`976`
`973`	`977`	`if (actualNDRangeTy->isLayoutIdentical(expectedNdrangeTy))`
`974`	`978`	`{`
`@@ -1016,7 +1020,7 @@ namespace //Anonymous`
`1016`	`1020`
`1017`	`1021`	`const auto newName = "__builtin_IB_get_block_simd_size";`
`1018`	`1022`	`auto calledFunction = _deviceExecCall->getCalledFunction();`
`1019`		`- if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");`
	`1023`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");`
`1020`	`1024`	`return CreateNewCall(newName, calledFunction->getReturnType(), blockIdValue);`
`1021`	`1025`	`}`
`1022`	`1026`	`};`
`@@ -1031,7 +1035,7 @@ namespace //Anonymous`
`1031`	`1035`	`{`
`1032`	`1036`	`const auto newName = "__builtin_IB_get_max_workgroup_size";`
`1033`	`1037`	`auto calledFunction = _deviceExecCall->getCalledFunction();`
`1034`		`- if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");`
	`1038`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");`
`1035`	`1039`	`return CreateNewCall(newName, calledFunction->getReturnType(), {});`
`1036`	`1040`	`}`
`1037`	`1041`	`};`
`@@ -1057,7 +1061,7 @@ namespace //Anonymous`
`1057`	`1061`
`1058`	`1062`	`const auto newName = "IGIL_calc_sub_group_count_for_ndrange";`
`1059`	`1063`	`auto calledFunction = _deviceExecCall->getCalledFunction();`
`1060`		`- if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");`
	`1064`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");`
`1061`	`1065`	`return CreateNewCall(newName, calledFunction->getReturnType(), args);`
`1062`	`1066`	`}`
`1063`	`1067`	`};`
`@@ -1082,7 +1086,7 @@ namespace //Anonymous`
`1082`	`1086`
`1083`	`1087`	`const auto newName = "__intel_calc_kernel_local_size_for_sub_group_count";`
`1084`	`1088`	`auto calledFunction = _deviceExecCall->getCalledFunction();`
`1085`		`- if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");`
	`1089`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");`
`1086`	`1090`	`return CreateNewCall(newName, calledFunction->getReturnType(), args);`
`1087`	`1091`	`}`
`1088`	`1092`	`};`
`@@ -1105,7 +1109,7 @@ namespace //Anonymous`
`1105`	`1109`
`1106`	`1110`	`const auto newName = "__intel_calc_kernel_max_num_subgroups";`
`1107`	`1111`	`auto calledFunction = _deviceExecCall->getCalledFunction();`
`1108`		`- if (calledFunction == nullptr) report_fatal_error("indirect calls are not supported");`
	`1112`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != calledFunction, "indirect calls are not supported");`
`1109`	`1113`	`return CreateNewCall(newName, calledFunction->getReturnType(), args);`
`1110`	`1114`	`}`
`1111`	`1115`	`};`
`@@ -1791,7 +1795,7 @@ namespace //Anonymous`
`1791`	`1795`
`1792`	`1796`	`// get list of calls which "enqueue" this dispatcher`
`1793`	`1797`	`auto parentCalls = _dataContext.getCallHandlersFor(invokeFunc);`
`1794`		`- if (parentCalls.size() == 0) report_fatal_error("parent calls for invoke are not set");`
	`1798`	`+ IGC_ASSERT_EXIT_MESSAGE(parentCalls.size(), "parent calls for invoke are not set");`
`1795`	`1799`
`1796`	`1800`	`// lookup for a call which located in a "__kernel"`
`1797`	`1801`	`// if __kernel call will not found, use the one which is not self`
`@@ -1810,7 +1814,7 @@ namespace //Anonymous`
`1810`	`1814`	`}`
`1811`	`1815`	`}`
`1812`	`1816`
`1813`		`- if (parentCallHandler == nullptr) report_fatal_error("Fail parent call lookup: possible closed self-enqueue");`
	`1817`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != parentCallHandler, "Fail parent call lookup: possible closed self-enqueue");`
`1814`	`1818`
`1815`	`1819`	`auto capturedValues = parentCallHandler->getArgs()->getParamValue()->getCapturedValues(_dataContext.getDispatcherForInvokeFunc(invokeFunc)->getCaptureIndicies());`
`1816`	`1820`	`auto& capture = capturedValues[captureNum];`
`@@ -1925,8 +1929,7 @@ namespace //Anonymous`
`1925`	`1929`	`if (!_invocations[invokeFunc]._dispatcher)`
`1926`	`1930`	`{`
`1927`	`1931`	`const Function* parentKernel = getParentKernel(invokeFunc);`
`1928`		`- if (!parentKernel)`
`1929`		`- report_fatal_error("Fail parent kernel lookup: possible closed self-enqueue");`
	`1932`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != parentKernel, "Fail parent kernel lookup: possible closed self-enqueue");`
`1930`	`1933`
`1931`	`1934`	`_invocations[invokeFunc]._dispatcher.reset(new Dispatcher(invokeFunc, parentKernel->getName(), _blocksNum++));`
`1932`	`1935`	`}`
`@@ -1960,7 +1963,7 @@ namespace //Anonymous`
`1960`	`1963`
`1961`	`1964`	`const llvm::Function* DataContext::getParentKernelImpl(const llvm::Function* invokeFunc, std::set<const llvm::Function*>& processedFuncs)`
`1962`	`1965`	`{`
`1963`		`- if (invokeFunc == nullptr) report_fatal_error("invoke function should not be null");`
	`1966`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != invokeFunc, "invoke function should not be null");`
`1964`	`1967`
`1965`	`1968`	`processedFuncs.insert(invokeFunc);`
`1966`	`1969`
`@@ -2383,7 +2386,7 @@ namespace //Anonymous`
`2383`	`2386`
`2384`	`2387`	`Capture::Capture(llvm::Value* val, DataContext& context)`
`2385`	`2388`	`{`
`2386`		`- if (val == nullptr) report_fatal_error("captured value is null");`
	`2389`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != val, "captured value is null");`
`2387`	`2390`	`value = val;`
`2388`	`2391`	`if (auto arg = context.getArgForValue(value))`
`2389`	`2392`	`{`
`@@ -2436,7 +2439,7 @@ namespace //Anonymous`
`2436`	`2439`	`, _block_invokeFunction(nullptr)`
`2437`	`2440`	`, _dataContext(dataContext)`
`2438`	`2441`	`{`
`2439`		`- if (!_paramStruct) report_fatal_error("Enqueue param is not a struct");`
	`2442`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != _paramStruct, "Enqueue param is not a struct");`
`2440`	`2443`
`2441`	`2444`	`if (_dataContext.getKindQuery().isBlockStructType(_paramStruct->getType())) {`
`2442`	`2445`	`/// ObjectiveC block value passed to OCL "device execution" call.`
`@@ -2455,12 +2458,9 @@ namespace //Anonymous`
`2455`	`2458`	`{`
`2456`	`2459`	`if (_paramValue == nullptr)`
`2457`	`2460`	`{`
`2458`		`- if (auto param = getParamArg()) {`
`2459`		`- _paramValue = _dataContext.getDeviceEnqueueParamValue(param->stripPointerCasts());`
`2460`		`- }`
`2461`		`- else {`
`2462`		`- report_fatal_error("Enqueue param is not set");`
`2463`		`- }`
	`2461`	`+ auto param = getParamArg();`
	`2462`	`+ IGC_ASSERT_EXIT_MESSAGE(nullptr != param, "Enqueue param is not set");`
	`2463`	`+ _paramValue = _dataContext.getDeviceEnqueueParamValue(param->stripPointerCasts());`
`2464`	`2464`	`}`
`2465`	`2465`	`return _paramValue;`
`2466`	`2466`	`}`
`@@ -2610,7 +2610,7 @@ namespace //Anonymous`
`2610`	`2610`	`case CaptureKind::IMAGE:`
`2611`	`2611`	`case CaptureKind::SAMPLER:`
`2612`	`2612`	`{`
`2613`		`- if (Capture::ARG_NUM_NONE == capturedValue.argNum) report_fatal_error("unknown argument number for an object");`
	`2613`	`+ IGC_ASSERT_EXIT_MESSAGE(Capture::ARG_NUM_NONE != capturedValue.argNum, "unknown argument number for an object");`
`2614`	`2614`	`auto objArgNumValue = llvm::ConstantInt::get(int32ty, capturedValue.argNum);`
`2615`	`2615`	`auto getObjIDFunc = getOrCreateFunc("__builtin_IB_get_object_id", int32ty, int32ty);`
`2616`	`2616`	`auto objIdValue = builder.CreateCall(getObjIDFunc, objArgNumValue, "obj_id");`
Original file line number	Diff line number	Diff line change
`@@ -220,7 +220,7 @@ bool GenericAddressDynamicResolution::visitLoadStoreInst(Instruction& I)`
`220`	`220`	`pointerAddressSpace = store->getPointerAddressSpace();`
`221`	`221`	`}`
`222`	`222`	`else {`
`223`		`- report_fatal_error("Unable to resolve generic address space pointer");`
	`223`	`+ IGC_ASSERT_EXIT_MESSAGE(0, "Unable to resolve generic address space pointer");`
`224`	`224`	`}`
`225`	`225`
`226`	`226`	`if (pointerAddressSpace == ADDRESS_SPACE_GENERIC) {`