diff --git a/Jenkinsfile b/Jenkinsfile index 3bbcccc1bb..15ca1c9935 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -38,6 +38,13 @@ String get_upstream_build_project( ) return upstream_cause.getUpstreamProject() } +//////////////////////////////////////////////////////////////////////// +// Construct the docker build image name +String docker_build_image_name( ) +{ + return "build-ubuntu-16.04" +} + //////////////////////////////////////////////////////////////////////// // Construct the relative path of the build directory String build_directory_rel( String build_config ) @@ -114,7 +121,7 @@ String checkout_and_version( String platform ) // The docker images contains all dependencies, including OS platform, to build def docker_build_image( String platform, String org, String optional_build_parm, String source_hip_rel, String from_image ) { - String build_image_name = "build-ubuntu-16.04" + String build_image_name = docker_build_image_name( ) String dockerfile_name = "dockerfile-build-ubuntu-16.04" def build_image = null @@ -170,7 +177,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - ctest -E hipPrintfKernel + ctest """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard @@ -179,7 +186,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf } // Only create packages from hcc based builds - if( platform.toLowerCase( ).startsWith( 'hcc-' ) ) + if( platform.toLowerCase( ).startsWith( 'rocm-' ) ) { stage("${platform} packaging") { @@ -191,10 +198,10 @@ def docker_build_inside_image( def build_image, String inside_args, String platf // No matter the base platform, all packages have the same name // Only upload 1 set of packages, so we don't have a race condition uploading packages - if( platform.toLowerCase( ).startsWith( 'hcc-ctu' ) ) + if( platform.toLowerCase( ).startsWith( 'rocm-head' ) ) { archiveArtifacts artifacts: "${build_dir_rel}/*.deb", fingerprint: true - // archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true + archiveArtifacts artifacts: "${build_dir_rel}/*.rpm", fingerprint: true } } } @@ -282,103 +289,26 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ } } -//////////////////////////////////////////////////////////////////////// -// hcc_integration_testing -// This function is sets up compilation and testing of HiP on a compiler downloaded from an upstream build -// Integration testing is centered around docker and constructing clean test environments every time - -// NOTES: I have implemeneted integration testing 3 different ways, and I've come to the conclusion nothing is perfect -// 1. I've tried having HCC push the test compiler to artifactory, and having HiP download the test docker image from artifactory -// a. The act of uploading and downloading images from artifactory takes minutes -// b. There is no good way of deleting images from a repository. You have to use an arcane CURL command and I don't know how -// to keep the password secret. These test integration images are meant to be ephemeral. -// 2. I tried 'docker save' to export a docker image into a tarball, and transfering the image through 'copy artifacts plugin' -// a. The HCC docker image uncompressed is over 1GB -// b. Compressing the docker image takes even longer than uploading the image to artifactory -// 3. Download the HCC .deb and dockerfile through 'copy artifacts plugin'. Create a new HCC image on the fly -// a. There is inefficency in building a new ubuntu image and installing HCC twice (once in HCC build, once here) -// b. This solution doesn't scale when we start testing downstream libraries - -// I've implemented solution #3 above, probably transitioning to #2 down the line (probably without compression) -String hcc_integration_testing( String inside_args, String job, String build_config ) -{ - // Attempt to make unique docker image names for each build, to support concurrent builds - // Mangle docker org name with upstream build info - String testing_org_name = 'hcc-test-' + get_upstream_build_project( ).replaceAll('/','-').toLowerCase( ) + '-' + get_upstream_build_num( ) - // String testing_org_name = 'hcc-test-artifacts-download' - - // Tag image name with this build number - String hcc_test_image_name = "hcc:${env.BUILD_NUMBER}" - - def hip_integration_image = null - - dir( 'integration-testing' ) - { - deleteDir( ) - - // This invokes 'copy artifact plugin' to copy archived files from upstream build - step([$class: 'CopyArtifact', filter: 'build/**/*.deb, docker/dockerfile-hcc-lc-*', - fingerprintArtifacts: true, projectName: get_upstream_build_project( ), flatten: true, - selector: [$class: 'TriggeredBuildSelector', allowUpstreamDependencies: false, fallbackToLastSuccessful: false, upstreamFilterStrategy: 'UseGlobalSetting'], - target: '.' ]) - // step([$class: 'CopyArtifact', filter: 'build/**/*.deb, docker/dockerfile-hcc-lc-*', - // fingerprintArtifacts: true, projectName: 'kknox/hcc/test-artifact-download', flatten: true, - // selector: [$class: 'LastCompletedBuildSelector'], - // target: '.' ]) - - docker.build( "${testing_org_name}/${hcc_test_image_name}", "-f dockerfile-hcc-lc-ubuntu-16.04 ." ) - } - - // Checkout source code, dependencies and version files - String source_hip_rel = checkout_and_version( job ) - - // Conctruct a binary directory path based on build config - String build_hip_rel = build_directory_rel( build_config ); - - // Build hip inside of the build environment - hip_integration_image = docker_build_image( job, testing_org_name, '', source_hip_rel, "${testing_org_name}/${hcc_test_image_name}" ) - - docker_build_inside_image( hip_integration_image, inside_args, job, '', build_config, source_hip_rel, build_hip_rel ) - - docker_clean_images( testing_org_name, '*' ) -} - //////////////////////////////////////////////////////////////////////// // -- MAIN // Following this line is the start of MAIN of this Jenkinsfile String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// Integration testing is a special path which implies testing of an upsteam build of hcc, -// but does not need testing across older builds of hcc or cuda. This is more of a compiler -// hcc unit test -// params.hcc_integration_test is set in HCC build -if( params.hcc_integration_test ) -{ - println "HCC integration testing" - - node('docker && rocm') - { - hcc_integration_testing( '--device=/dev/kfd --device=/dev/dri --group-add=video', 'hcc-ctu', build_config ) - } - - return -} - // The following launches 3 builds in parallel: hcc-ctu, hcc-1.6 and cuda -parallel hcc_ctu: +parallel rocm_1_9: { - node('docker && rocm && dkms') + node('hip-rocm') { - String hcc_ver = 'hcc-ctu' - String from_image = 'compute-artifactory:5001/radeonopencompute/hcc/clang_tot_upgrade/hcc-lc-ubuntu-16.04:latest' + String hcc_ver = 'rocm-1.9.x' + String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest' String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) // Create/reuse a docker image that represents the hip build environment - def hip_build_image = docker_build_image( hcc_ver, 'hip', ' --pull', source_hip_rel, from_image ) + def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) // Print system information for the log hip_build_image.inside( inside_args ) @@ -396,67 +326,34 @@ parallel hcc_ctu: // Build hip inside of the build environment docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) - // After a successful build, upload a docker image of the results - String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) + // Clean docker build image + docker_clean_images( 'hip', docker_build_image_name( ) ) + // After a successful build, upload a docker image of the results + /* + String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) if( params.push_image_to_docker_hub ) { docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) docker_clean_images( 'rocm', hip_image_name ) } docker_clean_images( job_name, hip_image_name ) + */ } }, -/* -hcc_1_6: +rocm_head: { - node('docker && rocm') + node('hip-rocm') { - String hcc_ver = 'hcc-1.6' - String from_image = 'rocm/dev-ubuntu-16.04:1.6.4' - String inside_args = '--device=/dev/kfd --device=/dev/dri' - - // Checkout source code, dependencies and version files - String source_hip_rel = checkout_and_version( hcc_ver ) - - // Create/reuse a docker image that represents the hip build environment - def hip_build_image = docker_build_image( hcc_ver, 'hip', ' --pull', source_hip_rel, from_image ) - - // Print system information for the log - hip_build_image.inside( inside_args ) - { - sh """#!/usr/bin/env bash - set -x - /opt/rocm/bin/rocm_agent_enumerator -t ALL - /opt/rocm/bin/hcc --version - """ - } - - // Conctruct a binary directory path based on build config - String build_hip_rel = build_directory_rel( build_config ); - - // Build hip inside of the build environment - docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) - - // Not pushing hip-hcc-1.6 builds at this time; saves a minute and nobody needs? - // String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) - // docker_clean_images( job_name, hip_image_name ) - } -}, -*/ -hcc_1_7: -{ - node('docker && rocm && dkms') - { - String hcc_ver = 'hcc-1.7' - String from_image = 'rocm/dev-ubuntu-16.04:latest' + String hcc_ver = 'rocm-head' + String from_image = 'ci_test_nodes/rocm-head/ubuntu-16.04:latest' String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' // Checkout source code, dependencies and version files String source_hip_rel = checkout_and_version( hcc_ver ) // Create/reuse a docker image that represents the hip build environment - def hip_build_image = docker_build_image( hcc_ver, 'hip', ' --pull', source_hip_rel, from_image ) + def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) // Print system information for the log hip_build_image.inside( inside_args ) @@ -474,41 +371,18 @@ hcc_1_7: // Build hip inside of the build environment docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) - // Not pushing hip-hcc-1.7 builds at this time; saves a minute and nobody needs? - // String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) - // docker_clean_images( job_name, hip_image_name ) - } -}/*, -nvcc: -{ - node('docker && cuda') - { - //////////////////////////////////////////////////////////////////////// - // Block of string constants customizing behavior for cuda - String nvcc_ver = 'nvcc-9.0' - String from_image = 'nvidia/cuda:9.0-devel' - String inside_args = '--runtime=nvidia'; + // Clean docker image + docker_clean_images( 'hip', docker_build_image_name( ) ) - // Checkout source code, dependencies and version files - String source_hip_rel = checkout_and_version( nvcc_ver ) - - // We pull public nvidia images - def hip_build_image = docker_build_image( nvcc_ver, 'hip', ' --pull', source_hip_rel, from_image ) - - // Print system information for the log - hip_build_image.inside( inside_args ) + // After a successful build, upload a docker image of the results + /* + String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) + if( params.push_image_to_docker_hub ) { - sh """#!/usr/bin/env bash - set -x - nvidia-smi - nvcc --version - """ + docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) + docker_clean_images( 'rocm', hip_image_name ) } - - // Conctruct a binary directory path based on build config - String build_hip_rel = build_directory_rel( build_config ); - - // Build hip inside of the build environment - docker_build_inside_image( hip_build_image, inside_args, nvcc_ver, "-DHIP_NVCC_FLAGS=--Wno-deprecated-gpu-targets", build_config, source_hip_rel, build_hip_rel ) + docker_clean_images( job_name, hip_image_name ) + */ } -}*/ +} diff --git a/cmake/FindHIP.cmake b/cmake/FindHIP.cmake index afa92de097..d2377e9adb 100644 --- a/cmake/FindHIP.cmake +++ b/cmake/FindHIP.cmake @@ -380,7 +380,7 @@ endmacro() ############################################################################### macro(HIP_PREPARE_TARGET_COMMANDS _target _format _generated_files _source_files) set(_hip_flags "") - string(TOUPPER _hip_build_configuration "${CMAKE_BUILD_TYPE}") + string(TOUPPER "${CMAKE_BUILD_TYPE}" _hip_build_configuration) if(HIP_HOST_COMPILATION_CPP) set(HIP_C_OR_CXX CXX) else() diff --git a/docker/dockerfile-build-ubuntu-16.04 b/docker/dockerfile-build-ubuntu-16.04 index 8f655f7c78..2b6de904e7 100644 --- a/docker/dockerfile-build-ubuntu-16.04 +++ b/docker/dockerfile-build-ubuntu-16.04 @@ -2,28 +2,13 @@ ARG base_image FROM ${base_image} -MAINTAINER Kent Knox +MAINTAINER Maneesh Gupta ARG user_uid -# Install Packages -# python and libnuma1 are dependencies of rocm_agent_enumerator -RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \ - sudo \ - build-essential \ - cmake \ - git \ - libelf-dev \ - rpm \ - python \ - libnuma1 \ - && \ - apt-get clean && \ - rm -rf /var/lib/apt/lists/* - # docker pipeline runs containers with particular uid # create a jenkins user with this specific uid so it can use sudo priviledges # Grant any member of sudo group password-less sudo privileges -RUN useradd --create-home -u ${user_uid} -G sudo --shell /bin/bash jenkins && \ +RUN useradd --create-home -u ${user_uid} -G sudo,video --shell /bin/bash jenkins && \ mkdir -p /etc/sudoers.d/ && \ echo '%sudo ALL=(ALL) NOPASSWD:ALL' | tee /etc/sudoers.d/sudo-nopasswd diff --git a/docs/markdown/hip_porting_guide.md b/docs/markdown/hip_porting_guide.md index 9810918174..2487d16e09 100644 --- a/docs/markdown/hip_porting_guide.md +++ b/docs/markdown/hip_porting_guide.md @@ -13,7 +13,7 @@ and provides practical suggestions on how to port CUDA code and work through com * [Converting a project "in-place"](#converting-a-project-in-place) - [Distinguishing Compiler Modes](#distinguishing-compiler-modes) * [Identifying HIP Target Platform](#identifying-hip-target-platform) - * [Identifying the Compiler: hcc or nvcc](#identifying-the-compiler-hcc-or-nvcc) + * [Identifying the Compiler: hcc, hip-clang, or nvcc](#identifying-the-compiler-hcc-hip-clang-or-nvcc) * [Identifying Current Compilation Pass: Host or Device](#identifying-current-compilation-pass-host-or-device) * [Compiler Defines: Summary](#compiler-defines-summary) - [Identifying Architecture Features](#identifying-architecture-features) @@ -137,13 +137,18 @@ All HIP projects target either the hcc or nvcc platform. The platform affects wh Many projects use a mixture of an accelerator compiler (hcc or nvcc) and a standard compiler (e.g., g++). These defines are set for both accelerator and standard compilers and thus are often the best option when writing code that uses conditional compilation. -### Identifying the Compiler: hcc or nvcc -Often, it’s useful to know whether the underlying compiler is hcc or nvcc. This knowledge can guard platform-specific code (features that only work on the nvcc or hcc path but not both) or aid in platform-specific performance tuning. +### Identifying the Compiler: hcc, hip-clang or nvcc +Often, it’s useful to know whether the underlying compiler is hcc, hip-clang or nvcc. This knowledge can guard platform-specific code (features that only work on the nvcc, hip-clang or hcc path but not all) or aid in platform-specific performance tuning. ``` #ifdef __HCC__ // Compiled with hcc +``` +``` +#ifdef __HIP__ +// Compiled with hip-clang + ``` ``` @@ -159,9 +164,9 @@ Often, it’s useful to know whether the underlying compiler is hcc or nvcc. This // Compiled with nvcc (Cuda language extensions enabled) ``` -hcc directly generates the host code (using the Clang x86 target) and passes the code to another host compiler. Thus, it lacks the equivalent of the \__CUDA_ACC define. +hcc and hip-clang directly generates the host code (using the Clang x86 target) and passes the code to another host compiler. Thus, it lacks the equivalent of the \__CUDA_ACC define. -The macro `__HIPCC__` is set if either `__HCC__` or `__CUDACC__` is defined. This configuration is useful in determining when code is being compiled using an accelerator-enabled compiler (hcc or nvcc) as opposed to a standard host compiler (GCC, ICC, Clang, etc.). +The macro `__HIPCC__` is set if either `__HCC__`, `__HIP__` or `__CUDACC__` is defined. This configuration is useful in determining when code is being compiled using an accelerator-enabled compiler (hcc or nvcc) as opposed to a standard host compiler (GCC, ICC, Clang, etc.). ### Identifying Current Compilation Pass: Host or Device @@ -176,22 +181,25 @@ Unlike `__CUDA_ARCH__`, the `__HIP_DEVICE_COMPILE__` value is 1 or undefined, an ### Compiler Defines: Summary -|Define | hcc | nvcc | Other (GCC, ICC, Clang, etc.) -|--- | --- | --- |---| +|Define | hcc | hip-clang | nvcc | Other (GCC, ICC, Clang, etc.) +|--- | --- | --- | --- |---| |HIP-related defines:| -|`__HIP_PLATFORM_HCC___`| Defined | Undefined | Defined if targeting hcc platform; undefined otherwise | -|`__HIP_PLATFORM_NVCC___`| Undefined | Defined | Defined if targeting nvcc platform; undefined otherwise | -|`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined -|`__HIPCC__` | Defined | Defined | Undefined -|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 +|`__HIP_PLATFORM_HCC___`| Defined | Defined | Undefined | Defined if targeting hcc platform; undefined otherwise | +|`__HIP_PLATFORM_NVCC___`| Undefined | Undefined | Defined | Defined if targeting nvcc platform; undefined otherwise | +|`__HIP_DEVICE_COMPILE__` | 1 if compiling for device; undefined if compiling for host | 1 if compiling for device; undefined if compiling for host |1 if compiling for device; undefined if compiling for host | Undefined +|`__HIPCC__` | Defined | Defined | Defined | Undefined +|`__HIP_ARCH_*` | 0 or 1 depending on feature support (see below) |0 or 1 depending on feature support (see below) | 0 or 1 depending on feature support (see below) | 0 |nvcc-related defines:| -|`__CUDACC__` | Undefined | Defined if source code is compiled by nvcc; undefined otherwise | Undefined -|`__NVCC__` | Undefined | Defined | Undefined -|`__CUDA_ARCH__` | Undefined | Unsigned representing compute capability (e.g., "130") if in device code; 0 if in host code | Undefined +|`__CUDACC__` | Undefined | Undefined | Defined if source code is compiled by nvcc; undefined otherwise | Undefined +|`__NVCC__` | Undefined | Undefined | Defined | Undefined +|`__CUDA_ARCH__` | Undefined | Undefined | Unsigned representing compute capability (e.g., "130") if in device code; 0 if in host code | Undefined |hcc-related defines:| -|`__HCC__` | Defined | Undefined | Undefined -|`__HCC_ACCELERATOR__` | Nonzero if in device code; otherwise undefined | Undefined | Undefined -|`__clang__` | Defined | Undefined | Defined if using Clang; otherwise undefined +|`__HCC__` | Defined | Undefined | Undefined | Undefined +|`__HCC_ACCELERATOR__` | Nonzero if in device code; otherwise undefined | Undefined | Undefined | Undefined +|hip-clang-related defines:| +|`__HIP__` | Undefined | Defined | Undefined | Undefined +|hcc/hip-clang common defines:| +|`__clang__` | Defined | Defined | Undefined | Defined if using Clang; otherwise undefined ## Identifying Architecture Features diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 5e1c8a4e6b..adf631d782 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -32,9 +32,6 @@ THE SOFTWARE. #include #include -typedef unsigned long ulong; -typedef unsigned int uint; - /* Integer Intrinsics */ @@ -47,20 +44,12 @@ __device__ static inline unsigned int __popcll(unsigned long long int input) { return __builtin_popcountl(input); } -__device__ static inline unsigned int __clz(unsigned int input) { - return input == 0 ? 32 : __builtin_clz(input); +__device__ static inline int __clz(int input) { + return __ockl_clz_u32((uint)input); } -__device__ static inline unsigned int __clzll(unsigned long long int input) { - return input == 0 ? 64 : ( input == 0 ? -1 : __builtin_clzl(input) ); -} - -__device__ static inline unsigned int __clz(int input) { - return input == 0 ? 32 : ( input > 0 ? __builtin_clz(input) : __builtin_clz(~input) ); -} - -__device__ static inline unsigned int __clzll(long long int input) { - return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); +__device__ static inline int __clzll(long long int input) { + return __ockl_clz_u64((ulong)input); } __device__ static inline unsigned int __ffs(unsigned int input) { @@ -79,12 +68,44 @@ __device__ static inline unsigned int __ffsll(long long int input) { return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; } -__device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); } +__device__ static inline unsigned int __brev(unsigned int input) { + return __llvm_bitrev_b32(input); +} __device__ static inline unsigned long long int __brevll(unsigned long long int input) { return __llvm_bitrev_b64(input); } +__device__ static inline unsigned int __lastbit_u32_u64(uint64_t input) { + return input == 0 ? -1 : __builtin_ctzl(input); +} + +__device__ static inline unsigned int __bitextract_u32(unsigned int src0, unsigned int src1, unsigned int src2) { + uint32_t offset = src1 & 31; + uint32_t width = src2 & 31; + return width == 0 ? 0 : (src0 << (32 - offset - width)) >> (32 - width); +} + +__device__ static inline uint64_t __bitextract_u64(uint64_t src0, unsigned int src1, unsigned int src2) { + uint64_t offset = src1 & 63; + uint64_t width = src2 & 63; + return width == 0 ? 0 : (src0 << (64 - offset - width)) >> (64 - width); +} + +__device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int src2, unsigned int src3) { + uint32_t offset = src2 & 31; + uint32_t width = src3 & 31; + uint32_t mask = (1 << width) - 1; + return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); +} + +__device__ static inline uint64_t __bitinsert_u64(uint64_t src0, uint64_t src1, unsigned int src2, unsigned int src3) { + uint64_t offset = src2 & 63; + uint64_t width = src3 & 63; + uint64_t mask = (1 << width) - 1; + return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); +} + __device__ static unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s); __device__ static unsigned int __hadd(int x, int y); __device__ static int __mul24(int x, int y); @@ -611,16 +632,43 @@ __device__ static inline float __ull2float_rz(unsigned long long int x) { return #ifdef __HCC_OR_HIP_CLANG__ +// Clock functions +__device__ long long int __clock64(); +__device__ long long int __clock(); +__device__ long long int clock64(); +__device__ long long int clock(); +// hip.amdgcn.bc - named sync +__device__ void __named_sync(int a, int b); + #ifdef __HIP_DEVICE_COMPILE__ // Clock functions -__device__ -inline -long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); } +#if __HCC__ +extern "C" uint64_t __clock_u64() __HC__; +#endif __device__ -inline -long long int __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); } +inline __attribute((always_inline)) +long long int __clock64() { +// ToDo: Unify HCC and HIP implementation. +#if __HCC__ + return (long long int) __clock_u64(); +#else + return (long long int) __builtin_amdgcn_s_memrealtime(); +#endif +} + +__device__ +inline __attribute((always_inline)) +long long int __clock() { return __clock64(); } + +__device__ +inline __attribute__((always_inline)) +long long int clock64() { return __clock64(); } + +__device__ +inline __attribute__((always_inline)) +long long int clock() { return __clock(); } // hip.amdgcn.bc - named sync __device__ diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 2a14b0b814..d64877d50a 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -30,6 +30,11 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" +typedef unsigned char uchar; +typedef unsigned short ushort; +typedef unsigned int uint; +typedef unsigned long ulong; + extern "C" __device__ __attribute__((const)) bool __ockl_wfany_i32(int); extern "C" __device__ __attribute__((const)) bool __ockl_wfall_i32(int); extern "C" __device__ uint __ockl_activelane_u32(void); @@ -40,6 +45,11 @@ extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint); extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int); extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint); +extern "C" __device__ __attribute__((const)) uchar __ockl_clz_u8(uchar); +extern "C" __device__ __attribute__((const)) ushort __ockl_clz_u16(ushort); +extern "C" __device__ __attribute__((const)) uint __ockl_clz_u32(uint); +extern "C" __device__ __attribute__((const)) ulong __ockl_clz_u64(ulong); + extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float); extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float); extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float); diff --git a/include/hip/hcc_detail/hip_complex.h b/include/hip/hcc_detail/hip_complex.h index 48a2852a5a..d19abd7247 100644 --- a/include/hip/hcc_detail/hip_complex.h +++ b/include/hip/hcc_detail/hip_complex.h @@ -121,7 +121,7 @@ THE SOFTWARE. return ret; \ } #define MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(ComplexT, T) \ - __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ + explicit __device__ __host__ ComplexT(T val) : x(val), y(val) {} \ __device__ __host__ ComplexT(T val1, T val2) : x(val1), y(val2) {} #endif @@ -131,7 +131,7 @@ struct hipFloatComplex { public: typedef float value_type; __device__ __host__ hipFloatComplex() : x(0.0f), y(0.0f) {} - __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} + explicit __device__ __host__ hipFloatComplex(float x) : x(x), y(0.0f) {} __device__ __host__ hipFloatComplex(float x, float y) : x(x), y(y) {} MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, unsigned short) MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipFloatComplex, signed short) @@ -151,7 +151,7 @@ struct hipDoubleComplex { public: typedef double value_type; __device__ __host__ hipDoubleComplex() : x(0.0f), y(0.0f) {} - __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} + explicit __device__ __host__ hipDoubleComplex(double x) : x(x), y(0.0f) {} __device__ __host__ hipDoubleComplex(double x, double y) : x(x), y(y) {} MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, unsigned short) MAKE_COMPONENT_CONSTRUCTOR_TWO_COMPONENT(hipDoubleComplex, signed short) diff --git a/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index af51e9d3fe..5d1c3d8f68 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/include/hip/hcc_detail/host_defines.h @@ -52,7 +52,7 @@ THE SOFTWARE. #endif // GENERIC_GRID_LAUNCH #define __noinline__ __attribute__((noinline)) -#define __forceinline__ __attribute__((always_inline)) +#define __forceinline__ inline __attribute__((always_inline)) /* @@ -71,7 +71,7 @@ THE SOFTWARE. #define __constant__ __attribute__((constant)) #define __noinline__ __attribute__((noinline)) -#define __forceinline__ __attribute__((always_inline)) +#define __forceinline__ inline __attribute__((always_inline)) #else diff --git a/tests/src/compiler/hipClassKernel.cpp b/tests/src/compiler/hipClassKernel.cpp new file mode 100644 index 0000000000..b3502ca9ae --- /dev/null +++ b/tests/src/compiler/hipClassKernel.cpp @@ -0,0 +1,270 @@ +/* +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 "hipClassKernel.h" + +// check sizeof empty class is 1 +__global__ void +emptyClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testClassEmpty ob1,ob2; + result_ecd[tid] = (sizeof(testClassEmpty) == 1) + && (&ob1 != &ob2); +} + +void HipClassTests::TestForEmptyClass(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(emptyClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +// tests for classes >8 bytes +__global__ void + sizeClassBKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeB) == 12) + && (sizeof(testSizeC) == 16) + && (sizeof(testSizeP1) == 6) + && (sizeof(testSizeP2) == 13) + && (sizeof(testSizeP3) == 8); +} + +void HipClassTests::TestForClassBSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassBKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +__global__ void +sizeClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeA) == 16) + && (sizeof(testSizeDerived) == 24) + && (sizeof(testSizeDerived2) == 20); +} + +void HipClassTests::TestForClassSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + +#ifdef ENABLE_VIRTUAL_TESTS +__global__ void + sizeVirtualClassKernel(bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (sizeof(testSizeDV) == 16) + && (sizeof(testSizeDerivedDV) == 16) + && (sizeof(testSizeVirtDerPack) == 24) + && (sizeof(testSizeVirtDer) == 24) + && (sizeof(testSizeDerMulti) == 48) ; + } + +void HipClassTests::TestForVirtualClassSize(void){ + bool *result_ecd, *result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + hipLaunchKernelGGL(sizeVirtualClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} +#endif + +// check pass by value +__global__ void +passByValueKernel(testPassByValue obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj.exI == 10) + && (obj.exC == 'C'); +} + +void HipClassTests::TestForPassByValue(void){ + bool *result_ecd,*result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + testPassByValue exObj; + exObj.exI = 10; + exObj.exC = 'C'; + hipLaunchKernelGGL(passByValueKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + exObj, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); +} + + // check obj created with hipMalloc +__global__ void +mallocObjKernel(testPassByValue *obj, bool* result_ecd) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + result_ecd[tid] = (obj->exI == 100) + && (obj->exC == 'C'); +} + +void HipClassTests::TestForMallocPassByValue(void){ + bool *result_ecd,*result_ech; + result_ech = HipClassTests::AllocateHostMemory(); + result_ecd = HipClassTests::AllocateDeviceMemory(); + + + testPassByValue *exObjM; + HIPCHECK(hipMalloc(&exObjM, sizeof(testPassByValue))); + exObjM->exI = 100; + exObjM->exC = 'C'; + hipLaunchKernelGGL(mallocObjKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0, + exObjM, + result_ecd); + + HipClassTests::VerifyResult(result_ech,result_ecd); + HipClassTests::FreeMem(result_ech,result_ecd); + +} + +// check if constr and destr are accessible from kernel +#ifdef ENABLE_DESTRUCTOR_TEST +__global__ void +testDeviceClassKernel() { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + testDeviceClass ob1; + testDeviceClass ob2; + ob2.iVar = 10; +} + +void HipClassTests::TestForConsrtDesrt(){ + testDeviceClass tDC; + hipLaunchKernelGGL(testDeviceClassKernel, + dim3(BLOCKS), + dim3(THREADS_PER_BLOCK), + 0, + 0); +} +#endif + +bool* HipClassTests::AllocateHostMemory(void){ + bool *result_ech; + HIPCHECK(hipHostMalloc(&result_ech, + NBOOL, + hipHostMallocDefault)); + return result_ech; +} + +bool* HipClassTests::AllocateDeviceMemory(void){ + bool* result_ecd; + HIPCHECK(hipMalloc(&result_ecd, + NBOOL)); + HIPCHECK(hipMemset(result_ecd, + false, + NBOOL)); + return result_ecd; +} + +void HipClassTests::VerifyResult(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipMemcpy(result_ech, + result_ecd, + BLOCKS*sizeof(bool), + hipMemcpyDeviceToHost)); + // validation on host side + for (int i = 0; i < BLOCKS; i++) { + HIPASSERT(result_ech[i] == true); + } +} + +void HipClassTests::FreeMem(bool* result_ech, bool* result_ecd){ + HIPCHECK(hipHostFree(result_ech)); + HIPCHECK(hipFree(result_ecd)); +} + +int main(){ + HipClassTests classTests; + classTests.TestForEmptyClass(); + test_passed(TestForEmptyClass); + classTests.TestForClassBSize(); + test_passed(TestForClassBSize); + classTests.TestForClassSize(); + test_passed(TestForClassSize); + classTests.TestForPassByValue(); + test_passed(TestForPassByValue); +// classTests.TestForMallocPassByValue(); + // test_passed(TestForMallocPassByValue); #this test is crashing + +#ifdef ENABLE_VIRTUAL_TESTS + classTests.TestForVirtualClassSize(); + test_passed(TestForVirtualClassSize); +#endif + + + +#ifdef ENABLE_DESTRUCTOR_TEST + classTests.TestForConsrtDesrt(); + test_passed(TestForConsrtDesrt); +#endif +} diff --git a/tests/src/compiler/hipClassKernel.h b/tests/src/compiler/hipClassKernel.h new file mode 100644 index 0000000000..8538276662 --- /dev/null +++ b/tests/src/compiler/hipClassKernel.h @@ -0,0 +1,185 @@ +/* +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. +*/ +#ifndef _COMPILER_HIPCLASSKERNEL_H_ +#define _COMPILER_HIPCLASSKERNEL_H_ + +#include +#include +#include "hip/hip_runtime.h" +#include "test_common.h" + +static const int BLOCKS = 512; +static const int THREADS_PER_BLOCK = 1; +static const int ENABLE_DESTRUCTOR_TEST = 0; +static const int ENABLE_VIRTUAL_TESTS = 0; +size_t NBOOL = BLOCKS * sizeof(bool); + +#define test_passed(test_name) printf("%s %s PASSED!%s\n", KGRN, #test_name, KNRM); + + +class testClassEmpty {}; + +class testPassByValue{ + public: + int exI; + char exC; +}; + +class testSizeA { + public: + float xa; + int ia; + double da; + static char ca; +}; + +class testSizeDerived : testSizeA { + public: + float fd; +}; + +#pragma pack(push,4) +class testSizeDerived2 : testSizeA { + public: + float fd; +}; +#pragma pack(pop) + +class testSizeB { + public: + char ab; + int ib; + char cb; +}; + +#ifdef ENBABLE_VIRTUAL_TESTS +class testSizeVirtDer : public virtual testSizeB { + public: + int ivd; +}; + +class testSizeVirtDer1 : public virtual testSizeB { + public: + int ivd1; +}; + +class testSizeDerMulti : public testSizeVirtDer, public testSizeVirtDer1 { + public: + int ivd2; +}; + +#pragma pack(push,4) +class testSizeVirtDerPack : public virtual testSizeB { + public: + int ivd; +}; +#pragma pack(pop) +#endif + +class testSizeC { + public: + char ac; + int ic; + int bc[2]; +}; + +#ifdef ENABLE_VIRTUAL_TESTS +class testSizeDV { + public: + virtual void __host__ __device__ func1(); + private: + int iDV; + +}; + +class testSizeDerivedDV : testSizeDV { + public: + virtual void __host__ __device__ funcD1(); + private: + int iDDV; +}; +#endif + +#pragma pack(push, 1) +class testSizeP1 { + public: + char ap; + int ip; + char cp; +}; +#pragma pack(pop) + +#pragma pack(push, 1) +class testSizeP2 { + public: + char ap1; + int ip1; + int bp1[2]; +}; +#pragma pack(pop) + +#pragma pack(push, 2) +class testSizeP3 { + public: + char ap2; + int ip2; + char cp2; +}; +#pragma pack(pop) + +#ifdef ENABLE_DESTRUCTOR_TEST +class testDeviceClass { + public: + int iVar; + __host__ __device__ testDeviceClass(); + __host__ __device__ testDeviceClass(int a); + __host__ __device__ ~testDeviceClass(); +}; + +__host__ __device__ +testDeviceClass::testDeviceClass() { + iVar = 5; +} + +__host__ __device__ +testDeviceClass::testDeviceClass(int a) { + iVar = a; +} +#endif + +#endif // _HIPCLASSKERNEL_H_ + +class HipClassTests{ + public: + void TestForEmptyClass(void); + void TestForClassBSize(void); + void TestForClassSize(void); + void TestForVirtualClassSize(void); + void TestForPassByValue(void); + void TestForMallocPassByValue(void); + void TestForConsrtDesrt(void); + + bool* AllocateHostMemory(void); + bool* AllocateDeviceMemory(void); + void VerifyResult(bool* result_ech, bool* result_ecd); + void FreeMem(bool* result_ech, bool* result_ecd); +}; diff --git a/tests/src/deviceLib/hipTestClock.cpp b/tests/src/deviceLib/hipTestClock.cpp new file mode 100644 index 0000000000..46f64e35a3 --- /dev/null +++ b/tests/src/deviceLib/hipTestClock.cpp @@ -0,0 +1,69 @@ +/* +Copyright (c) 2015-2016 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 +#include +#include +#include "test_common.h" + +#define HIP_ASSERT(status) assert(status == hipSuccess) + +#define LEN 512 +#define SIZE 2048 + +struct TestClock { + + static __global__ void kernel1(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = clock() + clock64() + __clock() + __clock64(); + } + + static __global__ void kernel2(int* Ad) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + Ad[tid] = clock() + clock64() + __clock() + __clock64() - Ad[tid]; + } + + void run() { + int *A, *Ad; + A = new int[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); + hipLaunchKernelGGL(kernel1, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + hipLaunchKernelGGL(kernel2, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(0 != A[i]); + } + } +}; + +int main() { + TestClock().run(); + passed(); +} diff --git a/tests/src/deviceLib/hip_bitextract.cpp b/tests/src/deviceLib/hip_bitextract.cpp new file mode 100644 index 0000000000..34e0be2b57 --- /dev/null +++ b/tests/src/deviceLib/hip_bitextract.cpp @@ -0,0 +1,220 @@ +/* +Copyright (c) 2015-2016 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 +#include +#include +#include +#include +#include +#include "hip/hip_runtime.h" +#include + +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) + +#define TEST_DEBUG (0) + + +// CPU implementation of bitextract +template +T bit_extract(T src0, unsigned int src1, unsigned int src2) { + unsigned int bits = sizeof(T) * 8; + T offset = src1 & (bits - 1); + T width = src2 & (bits - 1); + if (width == 0) { + return 0; + } else { + return (src0 << (bits - width - offset)) >> (bits - width); + } +} + +__global__ void HIP_kernel(hipLaunchParm lp, + unsigned int* out32, unsigned int* in32_0, + unsigned int* in32_1, unsigned int* in32_2, + unsigned long long int* out64, unsigned long long int* in64_0, + unsigned int* in64_1, unsigned int* in64_2) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + + out32[x] = __bitextract_u32(in32_0[x], in32_1[x], in32_2[x]); + out64[x] = __bitextract_u64(in64_0[x], in64_1[x], in64_2[x]); +} + + +using namespace std; + +int main() { + unsigned int* hostOut32; + unsigned int* hostSrc032; + unsigned int* hostSrc132; + unsigned int* hostSrc232; + unsigned long long int* hostOut64; + unsigned long long int* hostSrc064; + unsigned int* hostSrc164; + unsigned int* hostSrc264; + + unsigned int* deviceOut32; + unsigned int* deviceSrc032; + unsigned int* deviceSrc132; + unsigned int* deviceSrc232; + unsigned long long int* deviceOut64; + unsigned long long int* deviceSrc064; + unsigned int* deviceSrc164; + unsigned int* deviceSrc264; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + cout << " System minor " << devProp.minor << endl; + cout << " System major " << devProp.major << endl; + cout << " agent prop name " << devProp.name << endl; + + cout << "hip Device prop succeeded " << endl; + + unsigned int wave_size = devProp.warpSize; + unsigned int num_waves_per_block = 2; + unsigned int num_threads_per_block = wave_size * num_waves_per_block; + unsigned int num_blocks = 2; + unsigned int NUM = num_threads_per_block * num_blocks; + + int i; + int errors; + + hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + + hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int)); + hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int)); + hostSrc164 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + + // initialize the input data + std::random_device rd; + std::uniform_int_distribution uint32_src0_dist; + std::uniform_int_distribution uint32_src12_dist(0,31); + std::uniform_int_distribution uint64_src0_dist; + std::uniform_int_distribution uint64_src12_dist(0,63); + for (i = 0; i < NUM; i++) { + hostOut32[i] = 0; + hostSrc032[i] = uint32_src0_dist(rd); + hostSrc132[i] = uint32_src12_dist(rd); + hostSrc232[i] = uint32_src12_dist(rd); + hostOut64[i] = 0; + hostSrc064[i] = uint64_src0_dist(rd); + hostSrc164[i] = uint64_src12_dist(rd); + hostSrc264[i] = uint64_src12_dist(rd); + } + + HIP_ASSERT(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int))); + + HIP_ASSERT(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int))); + + HIP_ASSERT(hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + + HIP_ASSERT(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int), + hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + + + hipLaunchKernel(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), + 0, 0, + deviceOut32, deviceSrc032, deviceSrc132, deviceSrc232, + deviceOut64, deviceSrc064, deviceSrc164, deviceSrc264); + + + HIP_ASSERT(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost)); + HIP_ASSERT(hipMemcpy(hostOut64, deviceOut64, + NUM * sizeof(unsigned long long int), hipMemcpyDeviceToHost)); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostOut32[i] != bit_extract(hostSrc032[i], hostSrc132[i], hostSrc232[i])) { + errors++; +#if TEST_DEBUG + cout << "device: " << hostOut32[i] << " host: " + << bit_extract(hostSrc032[i], hostSrc132[i], hostSrc232[i]) + << " " << hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] << "\n"; +#endif + } + } + if (errors != 0) { + cout << "__bitextract_u32() FAILED\n" << endl; + return -1; + } else { + cout << "__bitextract_u32() checked!" << endl; + } + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostOut64[i] != bit_extract(hostSrc064[i], hostSrc164[i], hostSrc264[i])) { + errors++; +#if TEST_DEBUG + cout << "device: " << hostOut64[i] << " host: " + << bit_extract(hostSrc064[i], hostSrc164[i], hostSrc264[i]) + << " " << hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] << "\n"; +#endif + } + } + if (errors != 0) { + cout << "__bitextract_u64() FAILED" << endl; + return -1; + } else { + cout << "__bitextract_u64() checked!" << endl; + } + + cout << "__bitextract_u32() and __bitextract_u64() PASSED!" << endl; + + HIP_ASSERT(hipFree(deviceOut32)); + HIP_ASSERT(hipFree(deviceSrc032)); + HIP_ASSERT(hipFree(deviceSrc132)); + HIP_ASSERT(hipFree(deviceSrc232)); + HIP_ASSERT(hipFree(deviceOut64)); + HIP_ASSERT(hipFree(deviceSrc064)); + HIP_ASSERT(hipFree(deviceSrc164)); + HIP_ASSERT(hipFree(deviceSrc264)); + + free(hostOut32); + free(hostSrc032); + free(hostSrc132); + free(hostSrc232); + free(hostOut64); + free(hostSrc064); + free(hostSrc164); + free(hostSrc264); + + return errors; +} diff --git a/tests/src/deviceLib/hip_bitinsert.cpp b/tests/src/deviceLib/hip_bitinsert.cpp new file mode 100644 index 0000000000..063281768c --- /dev/null +++ b/tests/src/deviceLib/hip_bitinsert.cpp @@ -0,0 +1,239 @@ +/* +Copyright (c) 2015-2016 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 +#include +#include +#include +#include +#include +#include "hip/hip_runtime.h" +#include + +#define HIP_ASSERT(x) (assert((x) == hipSuccess)) + +#define TEST_DEBUG (0) + + +// CPU implementation of bitinsert +template +T bit_insert(T src0, T src1, unsigned int src2, unsigned int src3) { + unsigned int bits = sizeof(T) * 8; + T offset = src2 & (bits - 1); + T width = src3 & (bits - 1); + T mask = (1 << width) - 1; + return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset)); +} + +__global__ void HIP_kernel(hipLaunchParm lp, unsigned int* out32, + unsigned int* in32_0, unsigned int* in32_1, + unsigned int* in32_2, unsigned int* in32_3, + unsigned long long int* out64, unsigned long long int* in64_0, + unsigned long long int* in64_1, unsigned int* in64_2, + unsigned int* in64_3) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + + out32[x] = __bitinsert_u32(in32_0[x], in32_1[x], in32_2[x], in32_3[x]); + out64[x] = __bitinsert_u64(in64_0[x], in64_1[x], in64_2[x], in64_3[x]); +} + + +using namespace std; + +int main() { + unsigned int* hostOut32; + unsigned int* hostSrc032; + unsigned int* hostSrc132; + unsigned int* hostSrc232; + unsigned int* hostSrc332; + unsigned long long int* hostOut64; + unsigned long long int* hostSrc064; + unsigned long long int* hostSrc164; + unsigned int* hostSrc264; + unsigned int* hostSrc364; + + unsigned int* deviceOut32; + unsigned int* deviceSrc032; + unsigned int* deviceSrc132; + unsigned int* deviceSrc232; + unsigned int* deviceSrc332; + unsigned long long int* deviceOut64; + unsigned long long int* deviceSrc064; + unsigned long long int* deviceSrc164; + unsigned int* deviceSrc264; + unsigned int* deviceSrc364; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + cout << " System minor " << devProp.minor << endl; + cout << " System major " << devProp.major << endl; + cout << " agent prop name " << devProp.name << endl; + + cout << "hip Device prop succeeded " << endl; + + unsigned int wave_size = devProp.warpSize; + unsigned int num_waves_per_block = 2; + unsigned int num_threads_per_block = wave_size * num_waves_per_block; + unsigned int num_blocks = 2; + unsigned int NUM = num_threads_per_block * num_blocks; + + int i; + int errors; + + hostOut32 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc032 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc132 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc232 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc332 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + + hostOut64 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int)); + hostSrc064 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int)); + hostSrc164 = (unsigned long long int*)malloc(NUM * sizeof(unsigned long long int)); + hostSrc264 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + hostSrc364 = (unsigned int*)malloc(NUM * sizeof(unsigned int)); + + // initialize the input data + std::random_device rd; + std::uniform_int_distribution uint32_src01_dist; + std::uniform_int_distribution uint32_src23_dist(0,31); + std::uniform_int_distribution uint64_src01_dist; + std::uniform_int_distribution uint64_src23_dist(0,63); + for (i = 0; i < NUM; i++) { + hostOut32[i] = 0; + hostSrc032[i] = uint32_src01_dist(rd); + hostSrc132[i] = uint32_src01_dist(rd); + hostSrc232[i] = uint32_src23_dist(rd); + hostSrc232[i] = uint32_src23_dist(rd); + hostOut64[i] = 0; + hostSrc064[i] = uint64_src01_dist(rd); + hostSrc164[i] = uint64_src01_dist(rd); + hostSrc264[i] = uint64_src23_dist(rd); + hostSrc264[i] = uint64_src23_dist(rd); + } + + HIP_ASSERT(hipMalloc((void**)&deviceOut32, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc032, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc132, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc232, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc332, NUM * sizeof(unsigned int))); + + HIP_ASSERT(hipMalloc((void**)&deviceOut64, NUM * sizeof(unsigned long long int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc064, NUM * sizeof(unsigned long long int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc164, NUM * sizeof(unsigned long long int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc264, NUM * sizeof(unsigned int))); + HIP_ASSERT(hipMalloc((void**)&deviceSrc364, NUM * sizeof(unsigned int))); + + HIP_ASSERT(hipMemcpy(deviceSrc032, hostSrc032, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc132, hostSrc132, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc232, hostSrc232, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc332, hostSrc332, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + + HIP_ASSERT(hipMemcpy(deviceSrc064, hostSrc064, NUM * sizeof(unsigned long long int), + hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc164, hostSrc164, NUM * sizeof(unsigned long long int), + hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc264, hostSrc264, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceSrc364, hostSrc364, NUM * sizeof(unsigned int), hipMemcpyHostToDevice)); + + + hipLaunchKernel(HIP_kernel, dim3(num_blocks), dim3(num_threads_per_block), + 0, 0, + deviceOut32, deviceSrc032, deviceSrc132, deviceSrc232, deviceSrc332, + deviceOut64, deviceSrc064, deviceSrc164, deviceSrc264, deviceSrc364); + + + HIP_ASSERT(hipMemcpy(hostOut32, deviceOut32, NUM * sizeof(unsigned int), hipMemcpyDeviceToHost)); + HIP_ASSERT(hipMemcpy(hostOut64, deviceOut64, + NUM * sizeof(unsigned long long int), hipMemcpyDeviceToHost)); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostOut32[i] != bit_insert(hostSrc032[i], hostSrc132[i], + hostSrc232[i], hostSrc332[i])) { + errors++; +#if TEST_DEBUG + cout << "device: " << hostOut32[i] << " host: " + << bit_insert(hostSrc032[i], hostSrc132[i], hostSrc232[i], hostSrc332[i]) + << " " << hostSrc032[i] << " " << hostSrc132[i] << " " << hostSrc232[i] + << " " << hostSrc332[i] << "\n"; +#endif + } + } + if (errors != 0) { + cout << "__bitinsert_u32() FAILED\n" << endl; + return -1; + } else { + cout << "__bitinsert_u32() checked!" << endl; + } + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostOut64[i] != bit_insert(hostSrc064[i], hostSrc164[i], + hostSrc264[i], hostSrc364[i])) { + errors++; +#if TEST_DEBUG + cout << "device: " << hostOut64[i] << " host: " + << bit_insert(hostSrc064[i], hostSrc164[i], hostSrc264[i], hostSrc364[i]) + << " " << hostSrc064[i] << " " << hostSrc164[i] << " " << hostSrc264[i] + << " " << hostSrc364[i] << "\n"; +#endif + } + } + if (errors != 0) { + cout << "__bitinsert_u64() FAILED" << endl; + return -1; + } else { + cout << "__bitinsert_u64() checked!" << endl; + } + + cout << "__bitinsert_u32() and __bitinsert_u64() PASSED!" << endl; + + HIP_ASSERT(hipFree(deviceOut32)); + HIP_ASSERT(hipFree(deviceSrc032)); + HIP_ASSERT(hipFree(deviceSrc132)); + HIP_ASSERT(hipFree(deviceSrc232)); + HIP_ASSERT(hipFree(deviceSrc332)); + HIP_ASSERT(hipFree(deviceOut64)); + HIP_ASSERT(hipFree(deviceSrc064)); + HIP_ASSERT(hipFree(deviceSrc164)); + HIP_ASSERT(hipFree(deviceSrc264)); + HIP_ASSERT(hipFree(deviceSrc364)); + + free(hostOut32); + free(hostSrc032); + free(hostSrc132); + free(hostSrc232); + free(hostSrc332); + free(hostOut64); + free(hostSrc064); + free(hostSrc164); + free(hostSrc264); + free(hostSrc364); + + return errors; +} diff --git a/tests/src/deviceLib/hip_clz.cpp b/tests/src/deviceLib/hip_clz.cpp index 985da04967..604b62c7e2 100644 --- a/tests/src/deviceLib/hip_clz.cpp +++ b/tests/src/deviceLib/hip_clz.cpp @@ -67,6 +67,21 @@ unsigned int firstbit_u64(unsigned long long int a) { return pos; } +// Check implicit conversion will not cause ambiguity. +__device__ void test_ambiguity() { + short s; + unsigned short us; + float f; + int i; + unsigned int ui; + __clz(f); + __clz(s); + __clz(us); + __clzll(f); + __clzll(i); + __clzll(ui); +} + __global__ void HIP_kernel(hipLaunchParm lp, unsigned int* a, unsigned int* b, unsigned int* c, unsigned long long int* d, int width, int height) { int x = blockDim.x * blockIdx.x + threadIdx.x;