AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge 'master' into 'amd-master'

فهرست منبع 
Change-Id: I08b1105ecfa1b64f53cb080fb16193c9d5820a66
This commit is contained in:
Jenkins
2018-09-04 04:09:31 -05:00
والد 265aef0fb4 52b8879f12
کامیت cc17091ca3
14فایلهای تغییر یافته به همراه1150 افزوده شده و 227 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
Jenkinsfile
فروخته شده
+39 -165
مشاهده پرونده
@@ -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 )
*/
}
}*/
}
cmake/FindHIP.cmake
+1 -1
مشاهده پرونده
@@ -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()
docker/dockerfile-build-ubuntu-16.04
+2 -17
مشاهده پرونده
@@ -2,28 +2,13 @@
ARG base_image
FROM ${base_image}
MAINTAINER Kent Knox <kent.knox@amd>
MAINTAINER Maneesh Gupta <maneesh.gupta@amd>
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
docs/markdown/hip_porting_guide.md
+26 -18
مشاهده پرونده
@@ -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
include/hip/hcc_detail/device_functions.h
+69 -21
مشاهده پرونده
@@ -32,9 +32,6 @@ THE SOFTWARE.
#include <hip/hcc_detail/llvm_intrinsics.h>
#include <stddef.h>
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__
include/hip/hcc_detail/device_library_decls.h
+10
مشاهده پرونده
@@ -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);
include/hip/hcc_detail/hip_complex.h
+3 -3
مشاهده پرونده
@@ -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)
include/hip/hcc_detail/host_defines.h
+2 -2
مشاهده پرونده
@@ -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
tests/src/compiler/hipClassKernel.cpp
+270
مشاهده پرونده
@@ -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
}
tests/src/compiler/hipClassKernel.h
+185
مشاهده پرونده
@@ -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 <iostream>
#include <string>
#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);
};
tests/src/deviceLib/hipTestClock.cpp
+69
مشاهده پرونده
@@ -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 <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#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();
}
tests/src/deviceLib/hip_bitextract.cpp
+220
مشاهده پرونده
@@ -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 <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <random>
#include "hip/hip_runtime.h"
#include <hip/device_functions.h>
#define HIP_ASSERT(x) (assert((x) == hipSuccess))
#define TEST_DEBUG (0)
// CPU implementation of bitextract
template <typename T>
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_t> uint32_src0_dist;
std::uniform_int_distribution<uint32_t> uint32_src12_dist(0,31);
std::uniform_int_distribution<uint64_t> uint64_src0_dist;
std::uniform_int_distribution<uint32_t> 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<uint32_t>(hostSrc032[i], hostSrc132[i], hostSrc232[i])) {
errors++;
#if TEST_DEBUG
cout << "device: " << hostOut32[i] << " host: "
<< bit_extract<uint32_t>(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<uint64_t>(hostSrc064[i], hostSrc164[i], hostSrc264[i])) {
errors++;
#if TEST_DEBUG
cout << "device: " << hostOut64[i] << " host: "
<< bit_extract<uint64_t>(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;
}
tests/src/deviceLib/hip_bitinsert.cpp
+239
مشاهده پرونده
@@ -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 <assert.h>
#include <stdio.h>
#include <algorithm>
#include <stdlib.h>
#include <iostream>
#include <random>
#include "hip/hip_runtime.h"
#include <hip/device_functions.h>
#define HIP_ASSERT(x) (assert((x) == hipSuccess))
#define TEST_DEBUG (0)
// CPU implementation of bitinsert
template <typename T>
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_t> uint32_src01_dist;
std::uniform_int_distribution<uint32_t> uint32_src23_dist(0,31);
std::uniform_int_distribution<uint64_t> uint64_src01_dist;
std::uniform_int_distribution<uint32_t> 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<uint32_t>(hostSrc032[i], hostSrc132[i],
hostSrc232[i], hostSrc332[i])) {
errors++;
#if TEST_DEBUG
cout << "device: " << hostOut32[i] << " host: "
<< bit_insert<uint32_t>(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<uint64_t>(hostSrc064[i], hostSrc164[i],
hostSrc264[i], hostSrc364[i])) {
errors++;
#if TEST_DEBUG
cout << "device: " << hostOut64[i] << " host: "
<< bit_insert<uint64_t>(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;
}
tests/src/deviceLib/hip_clz.cpp
+15
مشاهده پرونده
@@ -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;