Merge 'master' into 'amd-master'

Change-Id: I42f0b83d11945fc03fb75f4a553c37f9f42f28a4


[ROCm/clr commit: 8af72e8ce7]

projects/clr/hipamd/bin/hipcc

@@ -125,9 +125,7 @@ if ($HIP_PLATFORM eq "hcc") {
 
     # Force -stdlib=libc++ on UB14.04
     $HOST_OSVER= `cat /etc/os-release | grep "^VERSION_ID\=" | cut -d= -f2 | tr -d '\n'`;
-    if (($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"14.04\"")
-         or ($HOST_OSNAME eq "\"centos\"" and $HOST_OSVER eq "\"7\"")
-         or ($HOST_OSNAME eq "\"rhel\"" and $HOST_OSVER eq "\"7.4\"")) {
+    if ($HOST_OSNAME eq "ubuntu" and $HOST_OSVER eq "\"14.04\"") {
         $HIPCXXFLAGS .= " -stdlib=libc++";
         $setStdLib = 1;
     }

projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 extern texture<float, 2, hipReadModeElementType> tex;
 
-__global__ void tex2dKernel(hipLaunchParm lp, float* outputData, int width, int height) {
+extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
     int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
     int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
     outputData[y * width + x] = tex2D(tex, x, y);

projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp

@@ -27,7 +27,7 @@ THE SOFTWARE.
 #include <vector>
 #include <hip/hip_hcc.h>
 
-#define fileName "tex2dKernel.code"
+#define fileName "tex2dKernel.code.adipose"
 
 texture<float, 2, hipReadModeElementType> tex;
 bool testResult = false;
@@ -90,7 +90,6 @@ bool runTest(int argc, char** argv) {
 #ifdef __HIP_PLATFORM_HCC__
 
     struct {
-        uint32_t _hidden[6];  // genco path + wrapper-gen pass used hidden arguments.
         void* _Ad;
         unsigned int _Bd;
         unsigned int _Cd;

projects/clr/hipamd/src/hip_texture.cpp

@@ -389,6 +389,7 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t
     enum hipTextureFilterMode filterMode = tex->filterMode;
     int normalizedCoords = tex->normalized;
     hipTextureObject_t& textureObject = tex->textureObject;
+    *offset = 0;
     auto ctx = ihipGetTlsDefaultCtx();
     if (ctx) {
         hc::accelerator acc = ctx->getDevice()->_acc;
@@ -458,6 +459,7 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size
     enum hipTextureFilterMode filterMode = tex->filterMode;
     int normalizedCoords = tex->normalized;
     hipTextureObject_t& textureObject = tex->textureObject;
+    *offset = 0;
     auto ctx = ihipGetTlsDefaultCtx();
     if (ctx) {
         hc::accelerator acc = ctx->getDevice()->_acc;