From fd73443cb0a44a361be3ebc4a86320daee2eae44 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Tue, 27 Sep 2016 15:06:48 -0400
Subject: [PATCH] P4 to Git Change 1319465 by lmoriche@lmoriche_opencl_dev on
 2016/09/27 14:56:10

	SWDEV-103418 - [ROCm CQE][OCLonLC][Fiji] OCLCreateBuffer results in Segmentation fault
	Fix KernelBlitManager::createProgram to report an error if one or more blit kernel could not be created. Re-enable the image blit kernels in the blit program.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/blitcl.cpp#10 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#4 edit
---
 rocclr/runtime/device/blitcl.cpp       |  4 ----
 rocclr/runtime/device/rocm/rocblit.cpp | 29 +++++++++++++-------------
 2 files changed, 14 insertions(+), 19 deletions(-)

diff --git a/rocclr/runtime/device/blitcl.cpp b/rocclr/runtime/device/blitcl.cpp
index 78000884b7..853ebab465 100644
--- a/rocclr/runtime/device/blitcl.cpp
+++ b/rocclr/runtime/device/blitcl.cpp
@@ -80,9 +80,6 @@ __kernel void fillBuffer(
 {
     __amd_fillBuffer(bufUChar, bufUInt, pattern, patternSize, offset, size);
 }
-)
-#if !defined(WITH_LIGHTNING_COMPILER)
-BLIT_KERNELS(
 extern void __amd_copyBufferToImage(
     __global uint*, __write_only image2d_array_t, ulong4,
     int4, int4, uint4, ulong4);
@@ -164,7 +161,6 @@ __kernel void fillImage(
                       origin, size, type);
 }
 )
-#endif // !defined(WITH_LIGHTNING_COMPILER)
 ;
 
 } // namespace device
diff --git a/rocclr/runtime/device/rocm/rocblit.cpp b/rocclr/runtime/device/rocm/rocblit.cpp
index c4824d8bf7..5e43543890 100644
--- a/rocclr/runtime/device/rocm/rocblit.cpp
+++ b/rocclr/runtime/device/rocm/rocblit.cpp
@@ -1509,22 +1509,21 @@ KernelBlitManager::createProgram(Device& device)
     program_ = device.blitProgram()->program_;
     program_->retain();
 
-    bool result = false;
-    do {
-        // Create kernel objects for all blits
-        for (uint i = 0; i < BlitTotal; ++i) {
-            const amd::Symbol* symbol = program_->findSymbol(BlitName[i]);
-            if (symbol == NULL) {
-                break;
-            }
-            kernels_[i] = new amd::Kernel(*program_, *symbol, BlitName[i]);
-            if (kernels_[i] == NULL) {
-                break;
-            }
-        }
+    bool result = true;
 
-        result = true;
-    } while(!result);
+    // Create kernel objects for all blits
+    for (uint i = 0; i < BlitTotal; ++i) {
+        const amd::Symbol* symbol = program_->findSymbol(BlitName[i]);
+        if (symbol == NULL) {
+            result = false;
+            continue;
+        }
+        kernels_[i] = new amd::Kernel(*program_, *symbol, BlitName[i]);
+        if (kernels_[i] == NULL) {
+            result = false;
+            continue;
+        }
+    }
 
     return result;
 }