Merge 'master' into 'amd-master'
Change-Id: I0f076ee4d0d5a725856f689383600820544bde18
This commit is contained in:
Jenkins
Vendored
+39
-2
@@ -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')
|
||||
|
||||
@@ -517,9 +517,8 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
|
||||
|
||||
#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;
|
||||
|
||||
@@ -541,7 +540,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)
|
||||
}
|
||||
|
||||
@@ -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]];
|
||||
}
|
||||
|
||||
|
||||
Reference in New Issue
Block a user