added architecture specific macros

1. Added __HIP_ARCH_GFXNUM__
2. Usage, -D__HIP_ARCH_GFX803__=1

Change-Id: I68b3a85d62cfab3a45d2b7a70cb3518ab2565236


[ROCm/hip commit: 9a1989193a]
This commit is contained in:
Aditya Atluri
2017-02-08 19:45:32 -06:00
vanhempi 036a95a4db
commit d48f7e93cc
2 muutettua tiedostoa jossa 10 lisäystä ja 2 poistoa
+6 -2
Näytä tiedosto
@@ -266,7 +266,7 @@ foreach $arg (@ARGV)
$target_gfx802 = 1;
}
if($arg eq '--amdgpu-target=gfx803')
{
{
$target_gfx803 = 1;
}
if(($trimarg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
@@ -337,18 +337,22 @@ if($HIP_PLATFORM eq "hcc"){
# Handle ROCm target platform
if ($target_gfx701 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx701";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX701__=1 ";
}
if ($target_gfx801 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx801";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX801__=1 ";
}
if ($target_gfx802 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx802";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX802__=1 ";
}
if ($target_gfx803 eq 1) {
$HIPLDFLAGS .= " --amdgpu-target=gfx803";
$HIPCXXFLAGS .= " -D__HIP_ARCH_GFX803__=1 ";
$ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
}
if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0)
if ($target_gfx701 eq 0 and $target_gfx801 eq 0 and $target_gfx802 eq 0 and $target_gfx803 eq 0)
{
$HIPLDFLAGS .= " --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803";
$ENV{HIP_HC_IR_GFX803}="$HIP_PATH/lib/hip_hc_gfx803.ll\n";
@@ -25,6 +25,8 @@ THE SOFTWARE.
#define HALF_SIZE 64*sizeof(__half)
#define HALF2_SIZE 64*sizeof(__half2)
#if __HIP_ARCH_GFX803__ > 0
__global__ void __halfMath(hipLaunchParm lp, __half *A, __half *B, __half *C) {
int tx = hipThreadIdx_x;
__half a = A[tx];
@@ -56,6 +58,8 @@ __global__ void __half2Math(hipLaunchParm lp, __half2 *A, __half2 *B, __half2 *C
c = __hmul2_sat(b, c);
}
#endif
int main(){
__half *A, *B, *C;
hipMalloc(&A, HALF_SIZE);