diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index 3bbcccc1bb..15ca1c9935 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/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/hipamd/cmake/FindHIP.cmake b/hipamd/cmake/FindHIP.cmake index afa92de097..d2377e9adb 100644 --- a/hipamd/cmake/FindHIP.cmake +++ b/hipamd/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/hipamd/docker/dockerfile-build-ubuntu-16.04 b/hipamd/docker/dockerfile-build-ubuntu-16.04 index 8f655f7c78..2b6de904e7 100644 --- a/hipamd/docker/dockerfile-build-ubuntu-16.04 +++ b/hipamd/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/hipamd/tests/src/compiler/hipClassKernel.cpp b/hipamd/tests/src/compiler/hipClassKernel.cpp new file mode 100644 index 0000000000..b3502ca9ae --- /dev/null +++ b/hipamd/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/hipamd/tests/src/compiler/hipClassKernel.h b/hipamd/tests/src/compiler/hipClassKernel.h new file mode 100644 index 0000000000..8538276662 --- /dev/null +++ b/hipamd/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); +};