diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 14b8ba650c..f796b0bf95 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/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; } diff --git a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index c817f4e42c..8b237d7d24 100644 --- a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" extern texture 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); diff --git a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 322bd1370e..2485d455ef 100644 --- a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -27,7 +27,7 @@ THE SOFTWARE. #include #include -#define fileName "tex2dKernel.code" +#define fileName "tex2dKernel.code.adipose" texture 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; diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index c471cb8f28..24c6eef3af 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/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;