From 83fac9281586226bdf97c978fd5e84085c7d8a51 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sun, 11 Mar 2018 00:12:22 +0530 Subject: [PATCH 1/5] Set offset value to 0 [ROCm/clr commit: c3d786d8101d5a2e802dc4761e74957f118a6c3a] --- projects/clr/hipamd/src/hip_texture.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index a69c91df8d..17012ac07c 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -391,6 +391,7 @@ hipError_t ihipBindTextureImpl(int dim, 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; @@ -467,6 +468,7 @@ hipError_t ihipBindTexture2DImpl(int dim, 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; From 765f3913ceb1fecda0e7de028405a49864214452 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 15 Mar 2018 16:46:01 -0400 Subject: [PATCH 2/5] On RHEL and CentOS, don't use libc++ since we are using the devtoolset-7 which has a newer g++ [ROCm/clr commit: 0501f6c58275b27c8df2cb001a511496fa970e20] --- projects/clr/hipamd/bin/hipcc | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) 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; } From c2f4352aade9424cbf7b0585d73fbc0287e666eb Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Mar 2018 12:35:25 +0530 Subject: [PATCH 3/5] Fixed function not found issue [ROCm/clr commit: 2ff9e5d46dcf23a010b8f6c219964ea726d1a145] --- .../hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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..ed223efedf 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(hipLaunchParm lp, 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); From 0a8ae7a07dc6bc005230ace8e562df153baecc19 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Mar 2018 12:54:44 +0530 Subject: [PATCH 4/5] Change co file name [ROCm/clr commit: da66236083cafe5203ad02f04aad9574d8400092] --- .../samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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..9842c17bc6 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; From ec51e3c539221d79f33ea7aab11f2e987dc47128 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 16 Mar 2018 22:50:25 +0530 Subject: [PATCH 5/5] Removed hidden args and hipLaunchParm from HIP/HCC path [ROCm/clr commit: 2184e7f6da90613eb2ba8441ea519bd38391f9a7] --- .../hipamd/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp | 2 +- .../samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp | 1 - 2 files changed, 1 insertion(+), 2 deletions(-) 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 ed223efedf..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; -extern "C" __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 9842c17bc6..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 @@ -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;