Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into skudchad-hip
Tento commit je obsažen v:
Saleel Kudchadker
vendorováno
+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 )
|
||||
*/
|
||||
}
|
||||
}*/
|
||||
}
|
||||
|
||||
@@ -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()
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
}
|
||||
@@ -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);
|
||||
};
|
||||
Odkázat v novém úkolu
Zablokovat Uživatele