Merge 'master' into 'amd-master'

Change-Id: I7fc7c2321bc4d89c69271a648f8450deb0c53e02


[ROCm/clr commit: a98791744f]

projects/clr/hipamd/bin/hipcc

@@ -90,7 +90,7 @@ if ($HIP_PLATFORM eq "clang") {
     }
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
     $HIPCC="$HIP_CLANG_PATH/clang++";
-    $HIPCXXFLAGS .= " -I$HIP_PATH/include";
+    $HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include";
 
     $HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
 } elsif ($HIP_PLATFORM eq "hcc") {
@@ -303,6 +303,10 @@ foreach $arg (@ARGV)
         $default_amdgpu_target = 0;
     }
 
+    # hip-clang does not accept --amdgpu-target= options.
+    if (($arg =~ /--amdgpu-target=/) and $HIP_PLATFORM eq 'clang' ) {
+        $swallowArg = 1;
+    }
 
     if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
     {

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h

@@ -29,6 +29,20 @@ THE SOFTWARE.
 #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
 
+#if defined(__HCC__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 1
+#define __HIP_CLANG_ONLY__ 0
+#elif defined(__clang__) && defined(__HIP__)
+#define __HCC_OR_HIP_CLANG__ 1
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 1
+#else
+#define __HCC_OR_HIP_CLANG__ 0
+#define __HCC_ONLY__ 0
+#define __HIP_CLANG_ONLY__ 0
+#endif
+
 //---
 // Top part of file can be compiled with any compiler
 
@@ -41,15 +55,16 @@ THE SOFTWARE.
 #include <stddef.h>
 #endif  //__cplusplus
 
-#if __HCC__
+#if __HCC_OR_HIP_CLANG__
 
 // Define NVCC_COMPAT for CUDA compatibility
 #define NVCC_COMPAT
 #define CUDA_SUCCESS hipSuccess
 
 #include <hip/hip_runtime_api.h>
+#endif  // __HCC_OR_HIP_CLANG__
 
-
+#if __HCC__
 // define HIP_ENABLE_PRINTF to enable printf
 #ifdef HIP_ENABLE_PRINTF
 #define HCC_ENABLE_ACCELERATOR_PRINTF 1
@@ -164,6 +179,10 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif
 
+#endif  // defined __HCC__
+
+#if __HCC_OR_HIP_CLANG__
+
 // TODO - hipify-clang - change to use the function call.
 //#define warpSize hc::__wavesize()
 static constexpr int warpSize = 64;
@@ -371,6 +390,10 @@ __device__ void __threadfence_system(void);
  * @}
  */
 
+#endif  // __HCC_OR_HIP_CLANG__
+
+#if defined __HCC__
+
 template <
     typename std::common_type<decltype(hc_get_group_id), decltype(hc_get_group_size),
                               decltype(hc_get_num_groups), decltype(hc_get_workitem_id)>::type f>
@@ -414,6 +437,8 @@ static constexpr Coordinates<hc_get_workitem_id> threadIdx;
 #define hipGridDim_y (hc_get_num_groups(1))
 #define hipGridDim_z (hc_get_num_groups(2))
 
+#endif // defined __HCC__
+#if __HCC_OR_HIP_CLANG__
 extern "C" __device__ void* __hip_hc_memcpy(void* dst, const void* src, size_t size);
 extern "C" __device__ void* __hip_hc_memset(void* ptr, uint8_t val, size_t size);
 extern "C" __device__ void* __hip_hc_malloc(size_t);
@@ -446,7 +471,9 @@ static inline __device__ void printf(const char* format, All... all) {}
 #endif
 
 #endif
+#endif //__HCC_OR_HIP_CLANG__
 
+#ifdef __HCC__
 
 #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE)
 
@@ -514,7 +541,9 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
  *   @}
  */
 
-
+//
+// hip-clang functions
+//
 #elif defined(__clang__) && defined(__HIP__)
 
 #define HIP_KERNEL_NAME(...) __VA_ARGS__
@@ -626,8 +655,13 @@ __DEVICE__ void inline __assert_fail(const char * __assertion,
     // Ignore all the args for now.
     __device_trap();
 }
+
+extern "C" __device__ __attribute__((noduplicate)) void __syncthreads();
+
 #pragma push_macro("__DEVICE__")
 
+#include <hip/hcc_detail/math_functions.h>
+
 #endif
 
 #endif  // HIP_HCC_DETAIL_RUNTIME_H

projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -150,6 +150,7 @@ typedef CUfunction hipFunction_t;
 typedef CUdeviceptr hipDeviceptr_t;
 typedef struct cudaArray hipArray;
 typedef struct cudaArray* hipArray_const_t;
+typedef cudaFuncAttributes hipFuncAttributes;
 #define hipMemcpy3DParms cudaMemcpy3DParms
 #define hipArrayDefault cudaArrayDefault
 
@@ -1107,6 +1108,10 @@ inline static hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule
     return hipCUResultTohipError(cuModuleGetFunction(function, module, kname));
 }
 
+inline static hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) {
+    return hipCUDAErrorTohipError(cudaFuncGetAttributes(attr, func));
+}
+
 inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
                                             const char* name) {
     return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name));

projects/clr/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
  * RUN: %t
  * HIT_END
  */

projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp

@@ -1,100 +0,0 @@
-/*
-Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.
-*/
-
-
-/* HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
- * HIT_END
- */
-
-#include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
-#include "test_common.h"
-#include <iostream>
-#include <string.h>
-
-#define N 512
-using namespace std;
-
-texture<int, hipTextureType1D, hipReadModeElementType> tex;
-
-bool testResult = true;
-
-__global__ void kernel(int *out) {
-  int x = blockIdx.x * blockDim.x + threadIdx.x;
-  out[x] = tex1Dfetch(tex, x);
-}
-
-void runTest(void);
-
-int main(int argc, char **argv) {
-  runTest();
-
-  if (testResult) {
-    passed();
-  } else {
-    exit(EXIT_FAILURE);
-  }
-}
-
-void runTest() {
-  string out;
-  int *tex_buf;
-  int val[N], i, output[N];
-  size_t size = 0;
-
-  for (i = 0; i < N; i++) {
-    val[i] = i;
-    output[i] = 0;
-  }
-  hipChannelFormatDesc chan_desc =
-      hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindUnsigned);
-
-  hipMalloc(&tex_buf, N * sizeof(int));
-
-  hipMemcpy(tex_buf, val, N * sizeof(int), hipMemcpyHostToDevice);
-
-  tex.addressMode[0] = hipAddressModeWrap;
-  tex.filterMode = hipFilterModeLinear;
-  tex.normalized = true;
-
-  hipBindTexture(&size, &tex, (void *)tex_buf, &chan_desc, N * sizeof(int));
-
-  dim3 dimBlock(64, 1, 1);
-  dim3 dimGrid(N / dimBlock.x, 1, 1);
-
-  hipLaunchKernelGGL(kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, output);
-
-  hipDeviceSynchronize();
-
-  hipMemcpy(output, tex_buf, N * sizeof(int), hipMemcpyDeviceToHost);
-
-  for (i = 0; i < N; i++) {
-    if (output[i] != val[i]) {
-      testResult = false;
-      return;
-    }
-  }
-  hipUnbindTexture(&tex);
-  hipFree(tex_buf);
-}

projects/clr/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp

@@ -0,0 +1,95 @@
+/*
+Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+
+/* HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+#define N 512
+
+texture<float, 1, hipReadModeElementType> tex;
+
+__global__ void kernel(float *out) {
+  int x = blockIdx.x * blockDim.x + threadIdx.x;
+  if(x<N){
+      out[x] = tex1Dfetch(tex, x);
+  }
+}
+
+int runTest(void);
+
+int main(int argc, char **argv) {
+    int testResult = runTest();
+    if (testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+}
+
+int runTest() {
+    int testResult = 1;
+    float *texBuf;
+    float val[N], output[N];
+    size_t size = 0;
+    float *devBuf;
+    for (int i = 0; i < N; i++) {
+        val[i] = (float)i;
+        output[i] = 0.0;
+    }
+    hipChannelFormatDesc chanDesc =
+        hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
+
+    HIPCHECK(hipMalloc(&texBuf, N * sizeof(float)));
+    HIPCHECK(hipMalloc(&devBuf, N * sizeof(float)));
+    HIPCHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice));
+  
+    tex.addressMode[0] = hipAddressModeClamp;
+    tex.addressMode[1] = hipAddressModeClamp;
+    tex.filterMode = hipFilterModePoint;
+    tex.normalized = 0;
+
+    HIPCHECK(hipBindTexture(&size, tex, (void *)texBuf, chanDesc, N * sizeof(float)));
+
+    dim3 dimBlock(64, 1, 1);
+    dim3 dimGrid(N / dimBlock.x, 1, 1);
+
+    hipLaunchKernelGGL(kernel, dim3(dimGrid), dim3(dimBlock), 0, 0, devBuf);
+    HIPCHECK(hipDeviceSynchronize());
+    HIPCHECK(hipMemcpy(output, devBuf, N * sizeof(float), hipMemcpyDeviceToHost));
+    for (int i = 0; i < N; i++) {
+        if (output[i] != val[i]) {
+            testResult = 0;
+            break;
+        }
+    }
+    HIPCHECK(hipUnbindTexture(&tex));
+    HIPCHECK(hipFree(texBuf));
+    HIPCHECK(hipFree(devBuf));
+    return testResult;
+}

projects/clr/hipamd/tests/src/texture/hipTextureObj1DFetch.cpp

@@ -0,0 +1,99 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+/*HIT_START
+ * BUILD: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+#define N 512
+
+__global__ void tex1dKernel(float *val, hipTextureObject_t obj) {
+    int k = blockIdx.x * blockDim.x + threadIdx.x;
+    if (k < N)
+        val[k] = tex1Dfetch<float>(obj, k);
+}
+
+int runTest(void);
+
+int main(int argc, char **argv) {
+    int testResult = runTest();
+    if(testResult) {
+        passed();
+    } else {
+        exit(EXIT_FAILURE);
+    }
+}
+
+int runTest() {
+    int testResult = 1;
+    // Allocating the required buffer on gpu device
+    float *texBuf, *texBufOut;
+    float val[N], output[N];
+    for (int i = 0; i < N; i++) {
+        val[i] = (i + 1) * (i + 1);
+        output[i] = 0.0;
+    }
+    HIPCHECK(hipMalloc(&texBuf, N * sizeof(float)));
+    HIPCHECK(hipMalloc(&texBufOut, N * sizeof(float)));
+    HIPCHECK(hipMemcpy(texBuf, val, N * sizeof(float), hipMemcpyHostToDevice));
+    HIPCHECK(hipMemset(texBufOut, 0, N * sizeof(float)));
+    hipResourceDesc resDescLinear;
+
+    memset(&resDescLinear, 0, sizeof(resDescLinear));
+    resDescLinear.resType = hipResourceTypeLinear;
+    resDescLinear.res.linear.devPtr = texBuf;
+    resDescLinear.res.linear.desc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
+    resDescLinear.res.linear.sizeInBytes = N * sizeof(float);
+
+    hipTextureDesc texDesc;
+    memset(&texDesc, 0, sizeof(texDesc));
+    texDesc.readMode = hipReadModeElementType;
+
+    // Creating texture object
+    hipTextureObject_t texObj = 0;
+    HIPCHECK(hipCreateTextureObject(&texObj, &resDescLinear, &texDesc, NULL));
+
+    dim3 dimBlock(64, 1, 1);
+    dim3 dimGrid(N / dimBlock.x, 1, 1);
+
+    hipLaunchKernelGGL(tex1dKernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
+                       texBufOut, texObj);
+    HIPCHECK(hipDeviceSynchronize());
+
+    HIPCHECK(hipMemcpy(output, texBufOut, N * sizeof(float), hipMemcpyDeviceToHost));
+
+    for(int i = 0; i < N; i++)
+        if (output[i] != val[i]) {
+            testResult = 0;
+            break;
+        }
+
+    HIPCHECK(hipDestroyTextureObject(texObj));
+    HIPCHECK(hipFree(texBuf));
+    HIPCHECK(hipFree(texBufOut));
+    return testResult;
+}

projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp

@@ -1,112 +0,0 @@
-/*
-Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.
-*/
-
-/*HIT_START
- * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
- * RUN: %t
- * HIT_END
- */
-
-#include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
-#include "test_common.h"
-#include <iostream>
-#include <string.h>
-
-#define N 512
-using namespace std;
-
-bool testResult = true;
-
-__global__ void tex1d_kernel(float *val, hipTextureObject_t obj) {
-  int k = blockIdx.x * blockDim.x + threadIdx.x;
-  val[k] = tex1Dfetch<float>(obj, k);
-}
-
-void runTest(void);
-
-int main(int argc, char **argv) {
-  runTest();
-
-  if (testResult) {
-    passed();
-  } else {
-    exit(EXIT_FAILURE);
-  }
-}
-
-void runTest() {
-
-  // Allocating the required buffer on gpu device
-  float *tex_buf, *tex_buf_check;
-  float val[N], output[N];
-  int i;
-  for (i = 0; i < N; i++)
-    val[i] = (i + 1) * (i + 1);
-  hipMalloc(&tex_buf, N * sizeof(float));
-
-  hipMalloc(&tex_buf_check, N * sizeof(float));
-
-  hipMemcpy(tex_buf, val, N * sizeof(float), hipMemcpyHostToDevice);
-
-  hipMemset(tex_buf_check, 0, N * sizeof(float));
-  hipResourceDesc res_lin;
-
-  memset(&res_lin, 0, sizeof(res_lin));
-
-  res_lin.resType = hipResourceTypeLinear;
-  res_lin.res.linear.devPtr = tex_buf;
-  res_lin.res.linear.desc.f = hipChannelFormatKindFloat;
-  res_lin.res.linear.desc.x = 32;
-  res_lin.res.linear.sizeInBytes = N * sizeof(float);
-
-  hipTextureDesc tex_desc;
-  memset(&tex_desc, 0, sizeof(tex_desc));
-  tex_desc.readMode = hipReadModeElementType;
-
-  // Creating texture object
-
-  hipTextureObject_t tex_obj = 0;
-
-  hipCreateTextureObject(&tex_obj, &res_lin, &tex_desc, NULL);
-
-  dim3 dimBlock(64, 1, 1);
-  dim3 dimGrid(N / dimBlock.x, 1, 1);
-
-  for (i = 0; i < N; i++)
-    output[i] = 0;
-
-  hipLaunchKernelGGL(tex1d_kernel, dim3(dimGrid), dim3(dimBlock), 0, 0,
-                     tex_buf_check, tex_obj);
-  hipDeviceSynchronize();
-
-  hipMemcpy(output, tex_buf_check, N * sizeof(float), hipMemcpyDeviceToHost);
-
-  for (i = 0; i < N; i++)
-    if (output[i] != val[i]) {
-      testResult = false;
-    }
-
-  hipDestroyTextureObject(tex_obj);
-  hipFree(tex_buf);
-  hipFree(tex_buf_check);
-}