diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index b162de2dce..7fafc30053 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/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)) { diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index e1b334aec3..1a6b0f7dda 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/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 #endif //__cplusplus -#if __HCC__ +#if __HCC_OR_HIP_CLANG__ // Define NVCC_COMPAT for CUDA compatibility #define NVCC_COMPAT #define CUDA_SUCCESS hipSuccess #include +#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::type f> @@ -414,6 +437,8 @@ static constexpr Coordinates 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 + #endif #endif // HIP_HCC_DETAIL_RUNTIME_H diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index d54a8de99c..fbff263295 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/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)); diff --git a/projects/clr/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp b/projects/clr/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp index 7d3eff5f73..6af87edb25 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp +++ b/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 */ diff --git a/projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp b/projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp deleted file mode 100644 index 2bcf78b1ca..0000000000 --- a/projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp +++ /dev/null @@ -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 -#include - -#define N 512 -using namespace std; - -texture 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); -} diff --git a/projects/clr/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp b/projects/clr/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp new file mode 100644 index 0000000000..5d9a1b103b --- /dev/null +++ b/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 tex; + +__global__ void kernel(float *out) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + if(x(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; +} diff --git a/projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp b/projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp deleted file mode 100644 index 4f755873f7..0000000000 --- a/projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp +++ /dev/null @@ -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 -#include - -#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(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); -}