From 494f7f6a382d1d3061e2b37827ef7f104ebb9d58 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 25 May 2018 12:22:16 -0400 Subject: [PATCH 1/9] Add math function declaration for hip-clang [ROCm/clr commit: f53e86e1b805ae11795a082a50d276da612082e7] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) 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..d682d21dbd 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -514,7 +514,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__ @@ -628,6 +630,8 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, } #pragma push_macro("__DEVICE__") +#include + #endif #endif // HIP_HCC_DETAIL_RUNTIME_H From c7d481a81263231fc73faba545ceb31b0920b72c Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 25 May 2018 16:18:46 -0400 Subject: [PATCH 2/9] Add more function declarations to hip-clang [ROCm/clr commit: 9c43970a27df175efea7247618a3ca3da596ec34] --- projects/clr/hipamd/bin/hipcc | 2 +- .../include/hip/hcc_detail/hip_runtime.h | 30 +++++++++++++++++-- 2 files changed, 29 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index b162de2dce..c1446d0f7d 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") { 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 d682d21dbd..1e469f5d03 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_CLANG__ + +#if defined __HCC__ + template < typename std::common_type::type f> @@ -628,6 +651,9 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, // Ignore all the args for now. __device_trap(); } + +__DEVICE__ void __syncthreads(); + #pragma push_macro("__DEVICE__") #include From b04a576c12a784f00dae4c7c9f5b831404c37b4c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 30 May 2018 11:24:48 +0530 Subject: [PATCH 3/9] Add signature for hipFuncGetAttributes to nvcc_details Change-Id: I06d5b50028fcfe3ea9d81f749b02be6b094c0977 [ROCm/clr commit: a923cac5d5de1d79f8fc01c11428079f1859897a] --- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 5 +++++ 1 file changed, 5 insertions(+) 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)); From cd912ecf7f9b41fb8d4f253197428838157c5f17 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 30 May 2018 11:26:11 +0530 Subject: [PATCH 4/9] [dtests] Fix nvcc path build errors in hipFuncGetAttributes.tst Change-Id: I000bed69162897f7b6edd733ed6e9acc93beb0ed [ROCm/clr commit: e2842f2f2892a6e0a59901b68ccd28437dcff503] --- .../hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 */ From 6126733ebe074f12172c52b636731f2bd5fa6afe Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 30 May 2018 11:42:01 +0530 Subject: [PATCH 5/9] Fixed tex1dFetch test [ROCm/clr commit: 3b0842e70adee91ef3b40002e0d7091d04e0307f] --- .../tests/src/texture/hipBindTexObj1D.cpp | 100 ------------------ .../src/texture/hipBindTexRef1DFetch.cpp | 95 +++++++++++++++++ 2 files changed, 95 insertions(+), 100 deletions(-) delete mode 100644 projects/clr/hipamd/tests/src/texture/hipBindTexObj1D.cpp create mode 100644 projects/clr/hipamd/tests/src/texture/hipBindTexRef1DFetch.cpp 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 Date: Mon, 28 May 2018 10:46:23 -0400 Subject: [PATCH 6/9] Drop --amdgpu-target= options for hip-clang They are replaced by --cuda-gpu-arch= options elsewhere [ROCm/clr commit: 3609e6f3b2768a72cb8eb8951c75550971cd6e7f] --- projects/clr/hipamd/bin/hipcc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index c1446d0f7d..7fafc30053 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -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)) { From 999df03229d7a811dc56332c21ac0b7a20f8b1a6 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 30 May 2018 06:38:41 -0400 Subject: [PATCH 7/9] Add more function declarations for hip-clang [ROCm/clr commit: 27fe2283d600ad6f44f4743b2db3d3f035d79129] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) 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 1e469f5d03..9737ecb19f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -390,7 +390,7 @@ __device__ void __threadfence_system(void); * @} */ -#endif // __HCC_OR_CLANG__ +#endif // __HCC_OR_HIP_CLANG__ #if defined __HCC__ @@ -437,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); @@ -472,7 +474,9 @@ static inline __device__ void printf(const char* format, All... all) {} #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) +#endif //__HCC_OR_HIP_CLANG__ +#ifdef __HCC__ #define HIP_KERNEL_NAME(...) (__VA_ARGS__) #define HIP_SYMBOL(X) #X @@ -652,8 +656,6 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, __device_trap(); } -__DEVICE__ void __syncthreads(); - #pragma push_macro("__DEVICE__") #include From b58a81f80ed10ce0c98426a3e20c1f2a33af9d0e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 30 May 2018 21:54:51 +0530 Subject: [PATCH 8/9] Fixed texture obj 1Dfetch test [ROCm/clr commit: 7cfd0599193923f2699f48bbbe716121aafbeeff] --- .../src/texture/hipTextureObj1DFetch.cpp | 99 ++++++++++++++++ .../tests/src/texture/tex1Dfetch_linear.cpp | 112 ------------------ 2 files changed, 99 insertions(+), 112 deletions(-) create mode 100644 projects/clr/hipamd/tests/src/texture/hipTextureObj1DFetch.cpp delete mode 100644 projects/clr/hipamd/tests/src/texture/tex1Dfetch_linear.cpp diff --git a/projects/clr/hipamd/tests/src/texture/hipTextureObj1DFetch.cpp b/projects/clr/hipamd/tests/src/texture/hipTextureObj1DFetch.cpp new file mode 100644 index 0000000000..faf6541ce1 --- /dev/null +++ b/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(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); -} From bd44b8c22ade3acc97d5eb37d7039642ca4d6b8e Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 30 May 2018 16:33:18 -0400 Subject: [PATCH 9/9] Fix __syncthreads for hip-clang [ROCm/clr commit: 0e0b028846c6eb47c4ee467fdf497253411481b2] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) 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 9737ecb19f..1a6b0f7dda 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -471,12 +471,12 @@ static inline __device__ void printf(const char* format, All... all) {} #endif #endif - - -#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) #endif //__HCC_OR_HIP_CLANG__ #ifdef __HCC__ + +#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) + #define HIP_KERNEL_NAME(...) (__VA_ARGS__) #define HIP_SYMBOL(X) #X @@ -656,6 +656,8 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, __device_trap(); } +extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); + #pragma push_macro("__DEVICE__") #include