From a06b21c5008b8c11e8a2527475726369ad8aee06 Mon Sep 17 00:00:00 2001 From: Laurent Morichetti Date: Wed, 21 Mar 2018 11:07:21 -0700 Subject: [PATCH 1/3] Add HIP_KERNEL_NAME/HIP_SYMBOL definitions for HIP-clang, and rename hipLaunch->hipLaunchByPtr. --- include/hip/hcc_detail/hip_runtime.h | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 58f741188d..6da14653ef 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -527,9 +527,8 @@ do {\ #elif defined(__clang__) && defined(__HIP__) -#define hipConfigureCall cudaConfigureCall -#define hipSetupArgument cudaSetupArgument -#define hipLaunch cudaLaunch +#define HIP_KERNEL_NAME(...) __VA_ARGS__ +#define HIP_SYMBOL(X) #X typedef int hipLaunchParm ; @@ -551,7 +550,7 @@ extern "C" { hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); hipError_t hipSetupArgument(const void *arg, size_t size, size_t offset); -hipError_t hipLaunch(const void *func); +hipError_t hipLaunchByPtr(const void *func); #if defined(__cplusplus) } From 21e0adc7005722a4adb077b10b311cb8d84dbabe Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 26 Mar 2018 15:21:45 +0530 Subject: [PATCH 2/3] [ci] Get CI builds to work again (#377) * [ci] Enable tests against HCC from ROCm1.7 as well * [ci] rocm 1.7 requires nodes with dkms label * [ci] Temporarily drop hcc_1_6 and hcc_1_7 builds --- Jenkinsfile | 41 +++++++++++++++++++++++++++++++++++++++-- 1 file changed, 39 insertions(+), 2 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 4ecad7facf..cfbe6cd883 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -359,7 +359,7 @@ if( params.hcc_integration_test ) node('docker && rocm') { - hcc_integration_testing( '--device=/dev/kfd', 'hcc-ctu', build_config ) + hcc_integration_testing( '--device=/dev/kfd --device=/dev/dri --group-add=video', 'hcc-ctu', build_config ) } return @@ -407,9 +407,10 @@ parallel hcc_ctu: docker_clean_images( job_name, hip_image_name ) } }, +/* hcc_1_6: { - node('docker && rocm && !dkms') + node('docker && rocm') { String hcc_ver = 'hcc-1.6' String from_image = 'rocm/dev-ubuntu-16.04:1.6.4' @@ -442,6 +443,42 @@ hcc_1_6: // 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 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 ) + + // 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.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') From 0d8b97274ac79733f4ee9ad1218655875f4122d4 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 26 Mar 2018 18:13:59 -0400 Subject: [PATCH 3/3] Change constant address space to 4 when compiling with a newer hcc. This is due to a recent change to address space mapping in the amdgpu compiler backend. --- include/hip/hcc_detail/texture_functions.h | 65 ++++++++++++---------- 1 file changed, 36 insertions(+), 29 deletions(-) diff --git a/include/hip/hcc_detail/texture_functions.h b/include/hip/hcc_detail/texture_functions.h index 7ab84695f3..e165c3c051 100644 --- a/include/hip/hcc_detail/texture_functions.h +++ b/include/hip/hcc_detail/texture_functions.h @@ -34,14 +34,21 @@ union TData { }; #define __TEXTURE_FUNCTIONS_DECL__ static __inline__ __device__ -#define ADDRESS_SPACE_2 __attribute__((address_space(2))) + + +#if (__hcc_workweek__ >= 18115) +#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4))) +#else +#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(2))) +#endif + #define TEXTURE_PARAMETERS_INIT \ - unsigned int ADDRESS_SPACE_2* i = (unsigned int ADDRESS_SPACE_2*)textureObject; \ - unsigned int ADDRESS_SPACE_2* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ + unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \ + unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ TData texel; -#define TEXTURE_REF_PARAMETERS_INIT \ - unsigned int ADDRESS_SPACE_2* i = (unsigned int ADDRESS_SPACE_2*)texRef.textureObject; \ - unsigned int ADDRESS_SPACE_2* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ +#define TEXTURE_REF_PARAMETERS_INIT \ + unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)texRef.textureObject; \ + unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \ TData texel; #define TEXTURE_SET_FLOAT *retVal = texel.f.x; @@ -146,90 +153,90 @@ union TData { #define TEXTURE_RETURN_FLOAT_XYZW return float4(texel.f.x, texel.f.y, texel.f.z, texel.f.w); extern "C" { -hc::short_vector::float4::vector_value_type __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_2* i, - unsigned int ADDRESS_SPACE_2* s, +hc::short_vector::float4::vector_value_type __ockl_image_sample_1D(unsigned int ADDRESS_SPACE_CONSTANT* i, + unsigned int ADDRESS_SPACE_CONSTANT* s, float c)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_1Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_2D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_2Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c)[[hc]]; -float __ockl_image_sample_2Dad(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c)[[hc]]; -float __ockl_image_sample_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_3D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_1D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, float c, float dx, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, float c, float dx, float dy)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_1Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, float dx, float dy)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_2D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, hc::short_vector::float2::vector_value_type dx, hc::short_vector::float2::vector_value_type dy)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_2Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, hc::short_vector::float2::vector_value_type dx, hc::short_vector::float2::vector_value_type dy)[[hc]]; -float __ockl_image_sample_grad_2Dad(unsigned int ADDRESS_SPACE_2* i, - unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_grad_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i, + unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, hc::short_vector::float2::vector_value_type dx, hc::short_vector::float2::vector_value_type dy)[[hc]]; -float __ockl_image_sample_grad_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_grad_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, hc::short_vector::float2::vector_value_type dx, hc::short_vector::float2::vector_value_type dy)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_grad_3D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, hc::short_vector::float4::vector_value_type dx, hc::short_vector::float4::vector_value_type dy)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_1D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, float c, float l)[[hc]]; + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, float c, float l)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_1Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, float l)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_2D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, float l)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_2Da( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, float l)[[hc]]; -float __ockl_image_sample_lod_2Dad(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_lod_2Dad(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, float l)[[hc]]; -float __ockl_image_sample_lod_2Dd(unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, +float __ockl_image_sample_lod_2Dd(unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float2::vector_value_type c, float l)[[hc]]; hc::short_vector::float4::vector_value_type __ockl_image_sample_lod_3D( - unsigned int ADDRESS_SPACE_2* i, unsigned int ADDRESS_SPACE_2* s, + unsigned int ADDRESS_SPACE_CONSTANT* i, unsigned int ADDRESS_SPACE_CONSTANT* s, hc::short_vector::float4::vector_value_type c, float l)[[hc]]; }