AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I280092c2c7ab06588eac29b646e9df81c44fa6a5
This commit is contained in:
Jenkins
2019-09-18 05:10:18 -04:00
parent 828e18eff7 313eb28379
commit 89b9ace88a
24 changed files with 986 additions and 311 deletions
Download Patch File Download Diff File Expand all files Collapse all files
bin/hipify-perl
+137 -9
View File
@@ -762,6 +762,7 @@ while (@ARGV) {
$ft{'include'} += s/\bcaffe2\/core\/common_cudnn.h\b/caffe2\/core\/hip\/common_miopen.h/g;
$ft{'include'} += s/\bcaffe2\/operators\/spatial_batch_norm_op.h\b/caffe2\/operators\/hip\/spatial_batch_norm_op_miopen.hip/g;
$ft{'include'} += s/\bchannel_descriptor.h\b/hip\/channel_descriptor.h/g;
$ft{'include'} += s/\bcooperative_groups.h\b/hip\/hip_cooperative_groups.h/g;
$ft{'include'} += s/\bcuda_fp16.h\b/hip\/hip_fp16.h/g;
$ft{'include'} += s/\bcuda_profiler_api.h\b/hip\/hip_profile.h/g;
$ft{'include'} += s/\bcuda_runtime_api.h\b/hip\/hip_runtime_api.h/g;
@@ -1442,6 +1443,7 @@ while (@ARGV) {
$ft{'numeric_literal'} += s/\bcudaComputeModeExclusive\b/hipComputeModeExclusive/g;
$ft{'numeric_literal'} += s/\bcudaComputeModeExclusiveProcess\b/hipComputeModeExclusiveProcess/g;
$ft{'numeric_literal'} += s/\bcudaComputeModeProhibited\b/hipComputeModeProhibited/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrCanMapHostMemory\b/hipDeviceAttributeCanMapHostMemory/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrClockRate\b/hipDeviceAttributeClockRate/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrComputeCapabilityMajor\b/hipDeviceAttributeComputeCapabilityMajor/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrComputeCapabilityMinor\b/hipDeviceAttributeComputeCapabilityMinor/g;
@@ -1449,9 +1451,11 @@ while (@ARGV) {
$ft{'numeric_literal'} += s/\bcudaDevAttrConcurrentKernels\b/hipDeviceAttributeConcurrentKernels/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrCooperativeLaunch\b/hipDeviceAttributeCooperativeLaunch/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrCooperativeMultiDeviceLaunch\b/hipDeviceAttributeCooperativeMultiDeviceLaunch/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrEccEnabled\b/hipDeviceAttributeEccEnabled/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrGlobalMemoryBusWidth\b/hipDeviceAttributeMemoryBusWidth/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrIntegrated\b/hipDeviceAttributeIntegrated/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrIsMultiGpuBoard\b/hipDeviceAttributeIsMultiGpuBoard/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrKernelExecTimeout\b/hipDeviceAttributeKernelExecTimeout/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrL2CacheSize\b/hipDeviceAttributeL2CacheSize/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxBlockDimX\b/hipDeviceAttributeMaxBlockDimX/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxBlockDimY\b/hipDeviceAttributeMaxBlockDimY/g;
@@ -1459,6 +1463,7 @@ while (@ARGV) {
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxGridDimX\b/hipDeviceAttributeMaxGridDimX/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxGridDimY\b/hipDeviceAttributeMaxGridDimY/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxGridDimZ\b/hipDeviceAttributeMaxGridDimZ/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxPitch\b/hipDeviceAttributeMaxPitch/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxRegistersPerBlock\b/hipDeviceAttributeMaxRegistersPerBlock/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxSharedMemoryPerBlock\b/hipDeviceAttributeMaxSharedMemoryPerBlock/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrMaxSharedMemoryPerMultiprocessor\b/hipDeviceAttributeMaxSharedMemoryPerMultiprocessor/g;
@@ -1474,6 +1479,7 @@ while (@ARGV) {
$ft{'numeric_literal'} += s/\bcudaDevAttrMultiProcessorCount\b/hipDeviceAttributeMultiprocessorCount/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrPciBusId\b/hipDeviceAttributePciBusId/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrPciDeviceId\b/hipDeviceAttributePciDeviceId/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrTextureAlignment\b/hipDeviceAttributeTextureAlignment/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrTotalConstantMemory\b/hipDeviceAttributeTotalConstantMemory/g;
$ft{'numeric_literal'} += s/\bcudaDevAttrWarpSize\b/hipDeviceAttributeWarpSize/g;
$ft{'numeric_literal'} += s/\bcudaErrorAssert\b/hipErrorAssert/g;
@@ -1821,20 +1827,142 @@ sub countSupportedDeviceFunctions
sub warnUnsupportedDeviceFunctions
{
my $line_num = shift;
my $line_num = shift;
my $m = 0;
# ToDo: list all of the supported functions
foreach $func (
"__syncthreads_and",
"__syncthreads_or",
"__prof_trigger"
"_Pow_int",
"__brkpt",
"__finite",
"__finitef",
"__finitel",
"__isinf",
"__isinff",
"__isinfl",
"__isnan",
"__isnanf",
"__isnanl",
"__pm0",
"__pm1",
"__pm2",
"__pm3",
"__prof_trigger",
"__shfl_down_sync",
"__shfl_sync",
"__shfl_up_sync",
"__shfl_xor_sync",
"__signbit",
"__signbitf",
"__signbitl",
"__trap",
"__vabs2",
"__vabs4",
"__vabsdiffs2",
"__vabsdiffs4",
"__vabsdiffu2",
"__vabsdiffu4",
"__vabsss2",
"__vabsss4",
"__vadd2",
"__vadd4",
"__vaddss2",
"__vaddss4",
"__vaddus2",
"__vaddus4",
"__vavgs2",
"__vavgs4",
"__vavgu2",
"__vavgu4",
"__vcmpeq2",
"__vcmpeq4",
"__vcmpges2",
"__vcmpges4",
"__vcmpgeu2",
"__vcmpgeu4",
"__vcmpgts2",
"__vcmpgts4",
"__vcmpgtu2",
"__vcmpgtu4",
"__vcmples2",
"__vcmples4",
"__vcmpleu4",
"__vcmplts2",
"__vcmplts4",
"__vcmpltu2",
"__vcmpltu4",
"__vcmpne2",
"__vcmpne4",
"__vhaddu2",
"__vhaddu4",
"__vmaxs2",
"__vmaxs4",
"__vmaxu2",
"__vmaxu4",
"__vmins2",
"__vmins4",
"__vminu2",
"__vminu4",
"__vneg2",
"__vneg4",
"__vnegss2",
"__vnegss4",
"__vsads2",
"__vsads4",
"__vsadu2",
"__vsadu4",
"__vseteq2",
"__vseteq4",
"__vsetges2",
"__vsetges4",
"__vsetgeu2",
"__vsetgeu4",
"__vsetgts2",
"__vsetgts4",
"__vsetgtu4",
"__vsetles2",
"__vsetles4",
"__vsetleu2",
"__vsetleu4",
"__vsetlts2",
"__vsetlts4",
"__vsetltu2",
"__vsetltu4",
"__vsetne2",
"__vsetne4",
"__vsub2",
"__vsub4",
"__vsubss2",
"__vsubss4",
"__vsubus2",
"__vsubus4",
"_fdsign",
"_ldsign",
"float2int",
"float_as_int",
"float_as_uint",
"int2float",
"int_as_float",
"llmax",
"llmin",
"mul24",
"mul64hi",
"mulhi",
"saturate",
"uint2float",
"uint_as_float",
"ullmax",
"ullmin",
"umax",
"umin",
"umul24"
)
{
# match math at the beginning of a word, but not if it already has a namespace qualifier ('::') :
my $mt = m/[:]?[:]?\b($func)\b(\w*\()/g;
if ($mt) {
# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));
# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);
my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g;
my $mt = m/($func)\s*\(\s*.*\s*\)/g;
if ($mt && !$mt_namespace) {
$m += $mt;
print STDERR " warning: $fileName:#$line_num : unsupported device function : $_\n";
print STDERR " warning: $fileName:$line_num: unsupported device function \"$func\": $_\n";
}
}
return $m;
docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md
+5 -5
View File
@@ -78,16 +78,16 @@
| 8 |*`CU_DEVICE_ATTRIBUTE_SHARED_MEMORY_PER_BLOCK`* |*`hipDeviceAttributeMaxSharedMemoryPerBlock`* |
| 9 |*`CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY`* |*`hipDeviceAttributeTotalConstantMemory`* |
| 10 |*`CU_DEVICE_ATTRIBUTE_WARP_SIZE`* |*`hipDeviceAttributeWarpSize`* |
| 11 |*`CU_DEVICE_ATTRIBUTE_MAX_PITCH`* | |
| 11 |*`CU_DEVICE_ATTRIBUTE_MAX_PITCH`* |*`hipDeviceAttributeMaxPitch`* |
| 12 |*`CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK`* |*`hipDeviceAttributeMaxRegistersPerBlock`* |
| 12 |*`CU_DEVICE_ATTRIBUTE_REGISTERS_PER_BLOCK`* |*`hipDeviceAttributeMaxRegistersPerBlock`* |
| 13 |*`CU_DEVICE_ATTRIBUTE_CLOCK_RATE`* |*`hipDeviceAttributeClockRate`* |
| 14 |*`CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT`* | |
| 14 |*`CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT`* |*`hipDeviceAttributeTextureAlignment`* |
| 15 |*`CU_DEVICE_ATTRIBUTE_GPU_OVERLAP`* | |
| 16 |*`CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT`* |*`hipDeviceAttributeMultiprocessorCount`* |
| 17 |*`CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT`* | |
| 17 |*`CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT`* |*`hipDeviceAttributeKernelExecTimeout`* |
| 18 |*`CU_DEVICE_ATTRIBUTE_INTEGRATED`* |*`hipDeviceAttributeIntegrated`* |
| 19 |*`CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY`* | |
| 19 |*`CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY`* |*`hipDeviceAttributeCanMapHostMemory`* |
| 20 |*`CU_DEVICE_ATTRIBUTE_COMPUTE_MODE`* |*`hipDeviceAttributeComputeMode`* |
| 21 |*`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH`* | |
| 22 |*`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH`* | |
@@ -103,7 +103,7 @@
| 29 |*`CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_ARRAY_NUMSLICES`* | |
| 30 |*`CU_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT`* | |
| 31 |*`CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS`* |*`hipDeviceAttributeConcurrentKernels`* |
| 32 |*`CU_DEVICE_ATTRIBUTE_ECC_ENABLED`* | |
| 32 |*`CU_DEVICE_ATTRIBUTE_ECC_ENABLED`* |*`hipDeviceAttributeEccEnabled`* |
| 33 |*`CU_DEVICE_ATTRIBUTE_PCI_BUS_ID`* |*`hipDeviceAttributePciBusId`* |
| 34 |*`CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID`* |*`hipDeviceAttributePciDeviceId`* |
| 35 |*`CU_DEVICE_ATTRIBUTE_TCC_DRIVER`* | |
docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+5 -5
View File
@@ -510,15 +510,15 @@
| 8 |*`cudaDevAttrMaxSharedMemoryPerBlock`* | |*`hipDeviceAttributeMaxSharedMemoryPerBlock`* |
| 9 |*`cudaDevAttrTotalConstantMemory`* | |*`hipDeviceAttributeTotalConstantMemory`* |
| 10 |*`cudaDevAttrWarpSize`* | |*`hipDeviceAttributeWarpSize`* |
| 11 |*`cudaDevAttrMaxPitch`* | | |
| 11 |*`cudaDevAttrMaxPitch`* | |*`hipDeviceAttributeMaxPitch`* |
| 12 |*`cudaDevAttrMaxRegistersPerBlock`* | |*`hipDeviceAttributeMaxRegistersPerBlock`* |
| 13 |*`cudaDevAttrClockRate`* | |*`hipDeviceAttributeClockRate`* |
| 14 |*`cudaDevAttrTextureAlignment`* | | |
| 14 |*`cudaDevAttrTextureAlignment`* | |*`hipDeviceAttributeTextureAlignment`* |
| 15 |*`cudaDevAttrGpuOverlap`* | | |
| 16 |*`cudaDevAttrMultiProcessorCount`* | |*`hipDeviceAttributeMultiprocessorCount`* |
| 17 |*`cudaDevAttrKernelExecTimeout`* | | |
| 17 |*`cudaDevAttrKernelExecTimeout`* | |*`hipDeviceAttributeKernelExecTimeout`* |
| 18 |*`cudaDevAttrIntegrated`* | |*`hipDeviceAttributeIntegrated`* |
| 19 |*`cudaDevAttrCanMapHostMemory`* | | |
| 19 |*`cudaDevAttrCanMapHostMemory`* | |*`hipDeviceAttributeCanMapHostMemory`* |
| 20 |*`cudaDevAttrComputeMode`* | |*`hipDeviceAttributeComputeMode`* |
| 21 |*`cudaDevAttrMaxTexture1DWidth`* | |*`hipDeviceAttributeMaxTexture1DWidth`* |
| 22 |*`cudaDevAttrMaxTexture2DWidth`* | |*`hipDeviceAttributeMaxTexture2DWidth`* |
@@ -531,7 +531,7 @@
| 29 |*`cudaDevAttrMaxTexture2DLayeredLayers`* | | |
| 30 |*`cudaDevAttrSurfaceAlignment`* | | |
| 31 |*`cudaDevAttrConcurrentKernels`* | |*`hipDeviceAttributeConcurrentKernels`* |
| 32 |*`cudaDevAttrEccEnabled`* | | |
| 32 |*`cudaDevAttrEccEnabled`* | |*`hipDeviceAttributeEccEnabled`* |
| 33 |*`cudaDevAttrPciBusId`* | |*`hipDeviceAttributePciBusId`* |
| 34 |*`cudaDevAttrPciDeviceId`* | |*`hipDeviceAttributePciDeviceId`* |
| 35 |*`cudaDevAttrTccDriver`* | | |
hipify-clang/README.md
+65 -62
View File
@@ -110,6 +110,7 @@ To run it:
cmake \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_TARGETS_TO_BUILD="NVPTX" \
-DCMAKE_BUILD_TYPE=Release \
../llvm
make -j install
@@ -121,6 +122,7 @@ To run it:
-G "Visual Studio 15 2017 Win64" \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_TARGETS_TO_BUILD="NVPTX" \
-DCMAKE_BUILD_TYPE=Release \
-Thost=x64 \
../llvm
@@ -252,67 +254,68 @@ Linux 5.2.0 - Platform OS
64 - hipify-clang binary bitness
64 - python 2.7.12 binary bitness
========================================
-- Testing: 58 tests, 12 threads --
PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 58)
PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 58)
PASS: hipify :: unit_tests/headers/headers_test_10.cu (3 of 58)
PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 58)
PASS: hipify :: unit_tests/headers/headers_test_01.cu (5 of 58)
PASS: hipify :: unit_tests/headers/headers_test_11.cu (6 of 58)
PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 58)
PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 58)
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 58)
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (12 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (14 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (15 of 58)
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (16 of 58)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (17 of 58)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (18 of 58)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (19 of 58)
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (20 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (21 of 58)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (22 of 58)
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (23 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (24 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (25 of 58)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (29 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (30 of 58)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (31 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (32 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (33 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (34 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (35 of 58)
PASS: hipify :: unit_tests/headers/headers_test_09.cu (36 of 58)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (38 of 58)
PASS: hipify :: unit_tests/samples/MallocManaged.cpp (39 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (40 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (41 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (42 of 58)
PASS: hipify :: unit_tests/samples/allocators.cu (43 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (44 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (45 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (46 of 58)
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (47 of 58)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (48 of 58)
PASS: hipify :: unit_tests/samples/coalescing.cu (49 of 58)
PASS: hipify :: unit_tests/samples/square.cu (50 of 58)
PASS: hipify :: unit_tests/samples/vec_add.cu (51 of 58)
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (52 of 58)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (53 of 58)
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (54 of 58)
PASS: hipify :: unit_tests/samples/intro.cu (55 of 58)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (56 of 58)
PASS: hipify :: unit_tests/samples/axpy.cu (57 of 58)
PASS: hipify :: unit_tests/samples/cudaRegister.cu (58 of 58)
Testing Time: 2.81s
Expected Passes : 58
-- Testing: 59 tests, 12 threads --
PASS: hipify :: unit_tests/headers/headers_test_01.cu (1 of 59)
PASS: hipify :: unit_tests/headers/headers_test_05.cu (2 of 59)
PASS: hipify :: unit_tests/headers/headers_test_03.cu (3 of 59)
PASS: hipify :: unit_tests/headers/headers_test_10.cu (4 of 59)
PASS: hipify :: unit_tests/headers/headers_test_11.cu (5 of 59)
PASS: hipify :: unit_tests/headers/headers_test_02.cu (6 of 59)
PASS: hipify :: unit_tests/headers/headers_test_06.cu (7 of 59)
PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 59)
PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 59)
PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (11 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (12 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (14 of 59)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 59)
PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (16 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (17 of 59)
PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (18 of 59)
PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (19 of 59)
PASS: hipify :: unit_tests/device/math_functions.cu (20 of 59)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (21 of 59)
PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (22 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (23 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (24 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (25 of 59)
PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 59)
PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (29 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (30 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (31 of 59)
PASS: hipify :: unit_tests/headers/headers_test_09.cu (32 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (33 of 59)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (34 of 59)
PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (35 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (36 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (37 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (38 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (39 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (40 of 59)
PASS: hipify :: unit_tests/samples/MallocManaged.cpp (41 of 59)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (42 of 59)
PASS: hipify :: unit_tests/samples/allocators.cu (43 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (44 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (45 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (46 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (47 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (48 of 59)
PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (49 of 59)
PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 59)
PASS: hipify :: unit_tests/samples/intro.cu (51 of 59)
PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (52 of 59)
PASS: hipify :: unit_tests/samples/axpy.cu (53 of 59)
PASS: hipify :: unit_tests/samples/static_shared_memory.cu (54 of 59)
PASS: hipify :: unit_tests/samples/vec_add.cu (55 of 59)
PASS: hipify :: unit_tests/samples/square.cu (56 of 59)
PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (57 of 59)
PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (58 of 59)
PASS: hipify :: unit_tests/samples/cudaRegister.cu (59 of 59)
Testing Time: 3.23s
Expected Passes : 59
[100%] Built target test-hipify
```
### <a name="windows"></a >Windows
@@ -327,7 +330,7 @@ LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.3.
Build system for the above configurations:
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.3).
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.5).
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
hipify-clang/src/CUDA2HIP_Driver_API_types.cpp
+5 -5
View File
@@ -305,7 +305,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaDevAttrWarpSize
{"CU_DEVICE_ATTRIBUTE_WARP_SIZE", {"hipDeviceAttributeWarpSize", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 10
// cudaDevAttrMaxPitch
{"CU_DEVICE_ATTRIBUTE_MAX_PITCH", {"hipDeviceAttributeMaxPitch", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 11
{"CU_DEVICE_ATTRIBUTE_MAX_PITCH", {"hipDeviceAttributeMaxPitch", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 11
// cudaDevAttrMaxRegistersPerBlock
{"CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK", {"hipDeviceAttributeMaxRegistersPerBlock", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 12
// no analogue
@@ -313,18 +313,18 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaDevAttrClockRate
{"CU_DEVICE_ATTRIBUTE_CLOCK_RATE", {"hipDeviceAttributeClockRate", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 13
// cudaDevAttrTextureAlignment
{"CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT", {"hipDeviceAttributeTextureAlignment", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 14
{"CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT", {"hipDeviceAttributeTextureAlignment", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 14
// cudaDevAttrGpuOverlap
// NOTE: Deprecated, use instead CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT
{"CU_DEVICE_ATTRIBUTE_GPU_OVERLAP", {"hipDeviceAttributeAsyncEngineCount", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 15
// cudaDevAttrMultiProcessorCount
{"CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT", {"hipDeviceAttributeMultiprocessorCount", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 16
// cudaDevAttrKernelExecTimeout
{"CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT", {"hipDeviceAttributeKernelExecTimeout", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 17
{"CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT", {"hipDeviceAttributeKernelExecTimeout", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 17
// cudaDevAttrIntegrated
{"CU_DEVICE_ATTRIBUTE_INTEGRATED", {"hipDeviceAttributeIntegrated", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 18
// cudaDevAttrCanMapHostMemory
{"CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY", {"hipDeviceAttributeCanMapHostMemory", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 19
{"CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY", {"hipDeviceAttributeCanMapHostMemory", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 19
// cudaDevAttrComputeMode
{"CU_DEVICE_ATTRIBUTE_COMPUTE_MODE", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 20
// cudaDevAttrMaxTexture1DWidth
@@ -359,7 +359,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DRIVER_TYPE_NAME_MAP{
// cudaDevAttrConcurrentKernels
{"CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS", {"hipDeviceAttributeConcurrentKernels", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 31
// cudaDevAttrEccEnabled
{"CU_DEVICE_ATTRIBUTE_ECC_ENABLED", {"hipDeviceAttributeEccEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 32
{"CU_DEVICE_ATTRIBUTE_ECC_ENABLED", {"hipDeviceAttributeEccEnabled", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 32
// cudaDevAttrPciBusId
{"CU_DEVICE_ATTRIBUTE_PCI_BUS_ID", {"hipDeviceAttributePciBusId", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 33
// cudaDevAttrPciDeviceId
hipify-clang/src/CUDA2HIP_Perl.cpp
+141
View File
@@ -0,0 +1,141 @@
/*
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.
*/
#include <sstream>
#include <regex>
#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/Support/Path.h"
#include "CUDA2HIP.h"
#include "CUDA2HIP_Scripting.h"
#include "ArgParse.h"
#include "StringUtils.h"
#include "LLVMCompat.h"
#include "Statistics.h"
using namespace llvm;
namespace perl {
const std::string space = " ";
const std::string double_space = space + space;
const std::string triple_space = double_space + space;
void generateUnsupportedDeviceFunctions(std::unique_ptr<std::ostream>& perlStreamPtr) {
unsigned int num = 0;
std::stringstream sUnsupported;
for (auto& ma : CUDA_DEVICE_FUNC_MAP) {
if (Statistics::isUnsupported(ma.second)) {
sUnsupported << (num ? ",\n" : "") << double_space << "\"" << ma.first.str() << "\"";
num++;
}
}
if (num) {
*perlStreamPtr.get() << "\nsub warnUnsupportedDeviceFunctions\n" << "{\n" << space << "my $line_num = shift;\n" << space << "my $m = 0;\n" << space << "foreach $func (\n";
*perlStreamPtr.get() << sUnsupported.str() << "\n" << space << ")\n";
*perlStreamPtr.get() << space << "{\n";
*perlStreamPtr.get() << double_space << "# match device function from the list of unsupported, except those, which have a namespace prefix (aka somenamespace::umin(...));\n";
*perlStreamPtr.get() << double_space << "# function with only global namespace qualifier '::' (aka ::umin(...)) should be treated as a device function (and warned as well as without such qualifier);\n";
*perlStreamPtr.get() << double_space << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
*perlStreamPtr.get() << double_space << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;\n";
*perlStreamPtr.get() << double_space << "if ($mt && !$mt_namespace) {\n";
*perlStreamPtr.get() << triple_space << "$m += $mt;\n";
*perlStreamPtr.get() << triple_space << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";\n";
*perlStreamPtr.get() << double_space << "}\n";
*perlStreamPtr.get() << space << "}\n";
*perlStreamPtr.get() << space << "return $m;\n";
*perlStreamPtr.get() << "}\n";
}
}
bool generate(bool Generate) {
if (!Generate) {
return true;
}
std::string dstPerlMap = OutputPerlMapFilename, dstPerlMapDir = OutputPerlMapDir;
if (dstPerlMap.empty()) {
dstPerlMap = "hipify-perl-map";
}
std::error_code EC;
if (!dstPerlMapDir.empty()) {
std::string sOutputPerlMapDirAbsPath = getAbsoluteDirectoryPath(OutputPerlMapDir, EC, "output hipify-perl map");
if (EC) {
return false;
}
dstPerlMap = sOutputPerlMapDirAbsPath + "/" + dstPerlMap;
}
SmallString<128> tmpFile;
StringRef ext = "hipify-tmp";
EC = sys::fs::createTemporaryFile(dstPerlMap, ext, tmpFile);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
return false;
}
std::unique_ptr<std::ostream> perlStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
std::string sConv = "my $conversions = ";
*perlStreamPtr.get() << "@statNames = (";
for (int i = 0; i < NUM_CONV_TYPES - 1; i++) {
*perlStreamPtr.get() << "\"" << counterNames[i] << "\", ";
sConv += "$ft{'" + std::string(counterNames[i]) + "'} + ";
}
*perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n";
*perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n";
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
if (Statistics::isUnsupported(ma.second)) {
continue;
}
if (i == ma.second.type) {
std::string sCUDA = ma.first.str();
std::string sHIP = ma.second.hipName.str();
sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/");
sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/");
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n";
}
}
}
else {
for (auto& ma : CUDA_RENAMES_MAP()) {
if (Statistics::isUnsupported(ma.second)) {
continue;
}
if (i == ma.second.type) {
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;\n";
}
}
}
}
generateUnsupportedDeviceFunctions(perlStreamPtr);
perlStreamPtr.get()->flush();
bool ret = true;
EC = sys::fs::copy_file(tmpFile, dstPerlMap);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPerlMap << "\n";
ret = false;
}
if (!SaveTemps) {
sys::fs::remove(tmpFile);
}
return ret;
}
}
hipify-clang/src/CUDA2HIP_Python.cpp
+103
View File
@@ -0,0 +1,103 @@
/*
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.
*/
#include <sstream>
#include "llvm/ADT/StringRef.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/Support/Path.h"
#include "CUDA2HIP.h"
#include "CUDA2HIP_Scripting.h"
#include "ArgParse.h"
#include "StringUtils.h"
#include "LLVMCompat.h"
#include "Statistics.h"
using namespace llvm;
namespace python {
bool generate(bool Generate) {
if (!Generate) {
return true;
}
std::string dstPythonMap = "cuda_to_hip_mappings.py", dstPythonMapDir = OutputPythonMapDir;
std::error_code EC;
if (!dstPythonMapDir.empty()) {
std::string sOutputPythonMapDirAbsPath = getAbsoluteDirectoryPath(OutputPythonMapDir, EC, "output hipify-python map");
if (EC) {
return false;
}
dstPythonMap = sOutputPythonMapDirAbsPath + "/" + dstPythonMap;
}
SmallString<128> tmpFile;
StringRef ext = "hipify-tmp";
EC = sys::fs::createTemporaryFile(dstPythonMap, ext, tmpFile);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
return false;
}
std::unique_ptr<std::ostream> pythonStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
*pythonStreamPtr.get() << "import collections\n\n";
*pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n";
*pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n";
const std::string sHIP_UNS = ", HIP_UNSUPPORTED";
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
sUnsupported = sHIP_UNS;
}
StringRef repName = Statistics::isToRoc(ma.second) ? ma.second.rocName : ma.second.hipName;
*pythonStreamPtr.get() << " (\"" << ma.first.str() << "\", (\"" << repName.str() << "\", " << counterTypes[i] << ", " << apiTypes[ma.second.apiType] << sUnsupported << ")),\n";
}
}
}
else {
for (auto& ma : CUDA_RENAMES_MAP()) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
sUnsupported = sHIP_UNS;
}
StringRef repName = Statistics::isToRoc(ma.second) ? ma.second.rocName : ma.second.hipName;
*pythonStreamPtr.get() << " (\"" << ma.first.str() << "\", (\"" << repName.str() << "\", " << counterTypes[i] << ", " << apiTypes[ma.second.apiType] << sUnsupported << ")),\n";
}
}
}
}
*pythonStreamPtr.get() << "])\n\n";
*pythonStreamPtr.get() << "CUDA_TO_HIP_MAPPINGS = [CUDA_RENAMES_MAP, C10_MAPPINGS, PYTORCH_SPECIFIC_MAPPINGS]\n";
pythonStreamPtr.get()->flush();
bool ret = true;
EC = sys::fs::copy_file(tmpFile, dstPythonMap);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPythonMap << "\n";
ret = false;
}
if (!SaveTemps) {
sys::fs::remove(tmpFile);
}
return true;
}
}
hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp
+5 -5
View File
@@ -234,24 +234,24 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// CU_DEVICE_ATTRIBUTE_WARP_SIZE
{"cudaDevAttrWarpSize", {"hipDeviceAttributeWarpSize", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 10
// CU_DEVICE_ATTRIBUTE_MAX_PITCH
{"cudaDevAttrMaxPitch", {"hipDeviceAttributeMaxPitch", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 11
{"cudaDevAttrMaxPitch", {"hipDeviceAttributeMaxPitch", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 11
// CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK
{"cudaDevAttrMaxRegistersPerBlock", {"hipDeviceAttributeMaxRegistersPerBlock", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 12
// CU_DEVICE_ATTRIBUTE_CLOCK_RATE
{"cudaDevAttrClockRate", {"hipDeviceAttributeClockRate", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 13
// CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT
{"cudaDevAttrTextureAlignment", {"hipDeviceAttributeTextureAlignment", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 14
{"cudaDevAttrTextureAlignment", {"hipDeviceAttributeTextureAlignment", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 14
// CU_DEVICE_ATTRIBUTE_GPU_OVERLAP
// NOTE: Is not deprecated as CUDA Driver's API analogue CU_DEVICE_ATTRIBUTE_GPU_OVERLAP
{"cudaDevAttrGpuOverlap", {"hipDeviceAttributeGpuOverlap", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 15
// CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT
{"cudaDevAttrMultiProcessorCount", {"hipDeviceAttributeMultiprocessorCount", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 16
// CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT
{"cudaDevAttrKernelExecTimeout", {"hipDeviceAttributeKernelExecTimeout", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 17
{"cudaDevAttrKernelExecTimeout", {"hipDeviceAttributeKernelExecTimeout", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 17
// CU_DEVICE_ATTRIBUTE_INTEGRATED
{"cudaDevAttrIntegrated", {"hipDeviceAttributeIntegrated", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 18
// CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY
{"cudaDevAttrCanMapHostMemory", {"hipDeviceAttributeCanMapHostMemory", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 19
{"cudaDevAttrCanMapHostMemory", {"hipDeviceAttributeCanMapHostMemory", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 19
// CU_DEVICE_ATTRIBUTE_COMPUTE_MODE
{"cudaDevAttrComputeMode", {"hipDeviceAttributeComputeMode", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 20
// CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH
@@ -277,7 +277,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
// CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS
{"cudaDevAttrConcurrentKernels", {"hipDeviceAttributeConcurrentKernels", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 31
// CU_DEVICE_ATTRIBUTE_ECC_ENABLED
{"cudaDevAttrEccEnabled", {"hipDeviceAttributeEccEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 32
{"cudaDevAttrEccEnabled", {"hipDeviceAttributeEccEnabled", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 32
// CU_DEVICE_ATTRIBUTE_PCI_BUS_ID
{"cudaDevAttrPciBusId", {"hipDeviceAttributePciBusId", "", CONV_NUMERIC_LITERAL, API_RUNTIME}}, // 33
// CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID
hipify-clang/src/CUDA2HIP_Scripting.h
+33
View File
@@ -0,0 +1,33 @@
/*
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.
*/
#pragma once
namespace perl {
bool generate(bool Generate = true);
}
namespace python {
bool generate(bool Generate = true);
}
hipify-clang/src/StringUtils.cpp
+60
View File
@@ -21,6 +21,10 @@ THE SOFTWARE.
*/
#include "StringUtils.h"
#include "LLVMCompat.h"
#include "llvm/ADT/SmallString.h"
using namespace llvm;
llvm::StringRef unquoteStr(llvm::StringRef s) {
if (s.size() > 1 && s.front() == '"' && s.back() == '"') {
@@ -35,3 +39,59 @@ void removePrefixIfPresent(std::string &s, const std::string& prefix) {
}
s.erase(0, prefix.size());
}
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
if (sFile.empty()) {
return sFile;
}
if (!sys::fs::exists(sFile)) {
llvm::errs() << "\n" << sHipify << sError << "source file: " << sFile << " doesn't exist\n";
EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
return "";
}
SmallString<256> fileAbsPath;
EC = llcompat::real_path(sFile, fileAbsPath, true);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": source file: " << sFile << "\n";
return "";
}
EC = std::error_code();
return fileAbsPath.c_str();
}
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
const std::string& sDirType, bool bCreateDir) {
if (sDir.empty()) {
return sDir;
}
EC = std::error_code();
SmallString<256> dirAbsPath;
if (sys::fs::exists(sDir)) {
if (sys::fs::is_regular_file(sDir)) {
llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n";
EC = std::error_code(static_cast<int>(std::errc::not_a_directory), std::generic_category());
return "";
}
}
else {
if (bCreateDir) {
EC = sys::fs::create_directory(sDir);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
return "";
}
}
else {
llvm::errs() << "\n" << sHipify << sError << sDirType << " directory: " << sDir << " doesn't exist\n";
EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
return "";
}
}
EC = llcompat::real_path(sDir, dirAbsPath, true);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
return "";
}
return dirAbsPath.c_str();
}
hipify-clang/src/StringUtils.h
+12
View File
@@ -34,3 +34,15 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
* If `s` starts with `prefix`, remove it. Otherwise, does nothing.
*/
void removePrefixIfPresent(std::string &s, const std::string& prefix);
/**
* Returns Absolute File Path based on filename, otherwise - error.
*/
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC);
/**
* Returns Absolute Directory Path based on directory name, otherwise - error;
* by default the directory is temporary and created.
*/
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
const std::string& sDirType = "temporary", bool bCreateDir = true);
hipify-clang/src/main.cpp
+5 -195
View File
@@ -27,208 +27,18 @@ THE SOFTWARE.
#include <chrono>
#include <iomanip>
#include <sstream>
#include <regex>
#include "CUDA2HIP.h"
#include "CUDA2HIP_Scripting.h"
#include "LLVMCompat.h"
#include "HipifyAction.h"
#include "ArgParse.h"
#include "StringUtils.h"
#include "llvm/Support/Debug.h"
#include "llvm/Support/Path.h"
#define DEBUG_TYPE "cuda2hip"
constexpr auto DEBUG_TYPE = "cuda2hip";
namespace ct = clang::tooling;
std::string getAbsoluteFilePath(const std::string& sFile, std::error_code& EC) {
if (sFile.empty()) {
return sFile;
}
if (!sys::fs::exists(sFile)) {
llvm::errs() << "\n" << sHipify << sError << "source file: " << sFile << " doesn't exist\n";
EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
return "";
}
SmallString<256> fileAbsPath;
EC = llcompat::real_path(sFile, fileAbsPath, true);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": source file: " << sFile << "\n";
return "";
}
EC = std::error_code();
return fileAbsPath.c_str();
}
std::string getAbsoluteDirectoryPath(const std::string& sDir, std::error_code& EC,
const std::string& sDirType = "temporary",
bool bCreateDir = true) {
if (sDir.empty()) {
return sDir;
}
EC = std::error_code();
SmallString<256> dirAbsPath;
if (sys::fs::exists(sDir)) {
if (sys::fs::is_regular_file(sDir)) {
llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n";
EC = std::error_code(static_cast<int>(std::errc::not_a_directory), std::generic_category());
return "";
}
} else {
if (bCreateDir) {
EC = sys::fs::create_directory(sDir);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
return "";
}
} else {
llvm::errs() << "\n" << sHipify << sError << sDirType << " directory: " << sDir << " doesn't exist\n";
EC = std::error_code(static_cast<int>(std::errc::no_such_file_or_directory), std::generic_category());
return "";
}
}
EC = llcompat::real_path(sDir, dirAbsPath, true);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n";
return "";
}
return dirAbsPath.c_str();
}
bool generatePerl(bool Generate = true) {
if (!Generate) {
return true;
}
std::string dstPerlMap = OutputPerlMapFilename, dstPerlMapDir = OutputPerlMapDir;
if (dstPerlMap.empty()) {
dstPerlMap = "hipify-perl-map";
}
std::error_code EC;
if (!dstPerlMapDir.empty()) {
std::string sOutputPerlMapDirAbsPath = getAbsoluteDirectoryPath(OutputPerlMapDir, EC, "output hipify-perl map");
if (EC) {
return false;
}
dstPerlMap = sOutputPerlMapDirAbsPath + "/" + dstPerlMap;
}
SmallString<128> tmpFile;
StringRef ext = "hipify-tmp";
EC = sys::fs::createTemporaryFile(dstPerlMap, ext, tmpFile);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
return false;
}
std::unique_ptr<std::ostream> perlStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
std::string sConv = "my $conversions = ";
*perlStreamPtr.get() << "@statNames = (";
for (int i = 0; i < NUM_CONV_TYPES - 1; i++) {
*perlStreamPtr.get() << "\"" << counterNames[i] << "\", ";
sConv += "$ft{'" + std::string(counterNames[i]) + "'} + ";
}
*perlStreamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");\n\n";
*perlStreamPtr.get() << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};\n\n";
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
if (Statistics::isUnsupported(ma.second)) {
continue;
}
if (i == ma.second.type) {
std::string sCUDA = ma.first.str();
std::string sHIP = ma.second.hipName.str();
sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/");
sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/");
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;\n";
}
}
} else {
for (auto& ma : CUDA_RENAMES_MAP()) {
if (Statistics::isUnsupported(ma.second)) {
continue;
}
if (i == ma.second.type) {
*perlStreamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;\n";
}
}
}
}
perlStreamPtr.get()->flush();
bool ret = true;
EC = sys::fs::copy_file(tmpFile, dstPerlMap);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPerlMap << "\n";
ret = false;
}
if (!SaveTemps) {
sys::fs::remove(tmpFile);
}
return ret;
}
bool generatePython(bool Generate = true) {
if (!Generate) {
return true;
}
std::string dstPythonMap = "cuda_to_hip_mappings.py", dstPythonMapDir = OutputPythonMapDir;
std::error_code EC;
if (!dstPythonMapDir.empty()) {
std::string sOutputPythonMapDirAbsPath = getAbsoluteDirectoryPath(OutputPythonMapDir, EC, "output hipify-python map");
if (EC) {
return false;
}
dstPythonMap = sOutputPythonMapDirAbsPath + "/" + dstPythonMap;
}
SmallString<128> tmpFile;
StringRef ext = "hipify-tmp";
EC = sys::fs::createTemporaryFile(dstPythonMap, ext, tmpFile);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << tmpFile << "\n";
return false;
}
std::unique_ptr<std::ostream> pythonStreamPtr = std::unique_ptr<std::ostream>(new std::ofstream(tmpFile.c_str(), std::ios_base::trunc));
*pythonStreamPtr.get() << "import collections\n\n";
*pythonStreamPtr.get() << "from pyHIPIFY.constants import *\n\n";
*pythonStreamPtr.get() << "CUDA_RENAMES_MAP = collections.OrderedDict([\n";
const std::string sHIP_UNS = ", HIP_UNSUPPORTED";
for (int i = 0; i < NUM_CONV_TYPES; i++) {
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
for (auto& ma : CUDA_INCLUDE_MAP) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
sUnsupported = sHIP_UNS;
}
StringRef repName = Statistics::isToRoc(ma.second) ? ma.second.rocName : ma.second.hipName;
*pythonStreamPtr.get() << " (\"" << ma.first.str() << "\", (\"" << repName.str() << "\", " << counterTypes[i] << ", " << apiTypes[ma.second.apiType] << sUnsupported << ")),\n";
}
}
}
else {
for (auto& ma : CUDA_RENAMES_MAP()) {
if (i == ma.second.type) {
std::string sUnsupported;
if (Statistics::isUnsupported(ma.second)) {
sUnsupported = sHIP_UNS;
}
StringRef repName = Statistics::isToRoc(ma.second) ? ma.second.rocName : ma.second.hipName;
*pythonStreamPtr.get() << " (\"" << ma.first.str() << "\", (\"" << repName.str() << "\", " << counterTypes[i] << ", " << apiTypes[ma.second.apiType] << sUnsupported << ")),\n";
}
}
}
}
*pythonStreamPtr.get() << "])\n\n";
*pythonStreamPtr.get() << "CUDA_TO_HIP_MAPPINGS = [CUDA_RENAMES_MAP, C10_MAPPINGS, PYTORCH_SPECIFIC_MAPPINGS]\n";
pythonStreamPtr.get()->flush();
bool ret = true;
EC = sys::fs::copy_file(tmpFile, dstPythonMap);
if (EC) {
llvm::errs() << "\n" << sHipify << sError << EC.message() << ": while copying " << tmpFile << " to " << dstPythonMap << "\n";
ret = false;
}
if (!SaveTemps) {
sys::fs::remove(tmpFile);
}
return true;
}
int main(int argc, const char **argv) {
std::vector<const char*> new_argv(argv, argv + argc);
if (std::find(new_argv.begin(), new_argv.end(), std::string("--")) == new_argv.end()) {
@@ -247,13 +57,13 @@ int main(int argc, const char **argv) {
llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file" << "\n";
return 1;
}
if (!generatePerl(GeneratePerl)) {
if (!perl::generate(GeneratePerl)) {
llvm::errs() << "\n" << sHipify << sError << "hipify-perl generating failed" << "\n";
return 1;
}
bool bToRoc = TranslateToRoc;
TranslateToRoc = true;
bool bToPython = generatePython(GeneratePython);
bool bToPython = python::generate(GeneratePython);
TranslateToRoc = bToRoc;
if (!bToPython) {
llvm::errs() << "\n" << sHipify << sError << "hipify-python generating failed" << "\n";
include/hip/hcc_detail/hip_runtime_api.h
-1
View File
@@ -40,7 +40,6 @@ THE SOFTWARE.
#endif
#include <hip/hcc_detail/host_defines.h>
#include <hip/hip_runtime_api.h>
#include <hip/hcc_detail/driver_types.h>
#include <hip/hcc_detail/hip_texture_types.h>
#include <hip/hcc_detail/hip_surface_types.h>
include/hip/hip_runtime_api.h
+12 -1
View File
@@ -122,6 +122,10 @@ typedef struct hipDeviceProp_t {
int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
size_t memPitch; ///<Maximum pitch in bytes allowed by memory copies
size_t textureAlignment; ///<Alignment requirement for textures
int kernelExecTimeoutEnabled; ///<Run time limit for kernels executed on the device
int ECCEnabled; ///<Device has ECC support enabled
} hipDeviceProp_t;
@@ -311,7 +315,14 @@ typedef enum hipDeviceAttribute_t {
hipDeviceAttributeMaxTexture3DDepth, ///< Maximum dimensions depth of 3D images in image elements
hipDeviceAttributeHdpMemFlushCntl, ///< Address of the HDP_MEM_COHERENCY_FLUSH_CNTL register
hipDeviceAttributeHdpRegFlushCntl ///< Address of the HDP_REG_COHERENCY_FLUSH_CNTL register
hipDeviceAttributeHdpRegFlushCntl, ///< Address of the HDP_REG_COHERENCY_FLUSH_CNTL register
hipDeviceAttributeMaxPitch, ///< Maximum pitch in bytes allowed by memory copies
hipDeviceAttributeTextureAlignment, ///<Alignment requirement for textures
hipDeviceAttributeKernelExecTimeout, ///<Run time limit for kernels executed on the device
hipDeviceAttributeCanMapHostMemory, ///<Device can map host memory into device address space
hipDeviceAttributeEccEnabled ///<Device has ECC support enabled
} hipDeviceAttribute_t;
enum hipComputeMode {
include/hip/nvcc_detail/hip_runtime_api.h
+20
View File
@@ -842,6 +842,11 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
p_prop->maxTexture3D[1] = cdprop.maxTexture3D[1];
p_prop->maxTexture3D[2] = cdprop.maxTexture3D[2];
p_prop->memPitch = cdprop.memPitch;
p_prop->textureAlignment = cdprop.textureAlignment;
p_prop->kernelExecTimeoutEnabled = cdprop.kernelExecTimeoutEnabled;
p_prop->ECCEnabled = cdprop.ECCEnabled;
return hipCUDAErrorTohipError(cerror);
}
@@ -946,6 +951,21 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att
case hipDeviceAttributeMaxTexture3DDepth:
cdattr = cudaDevAttrMaxTexture3DDepth;
break;
case hipDeviceAttributeMaxPitch:
cdattr = cudaDevAttrMaxPitch;
break;
case hipDeviceAttributeTextureAlignment:
cdattr = cudaDevAttrTextureAlignment;
break;
case hipDeviceAttributeKernelExecTimeout:
cdattr = cudaDevAttrKernelExecTimeout;
break;
case hipDeviceAttributeCanMapHostMemory:
cdattr = cudaDevAttrCanMapHostMemory;
break;
case hipDeviceAttributeEccEnabled:
cdattr = cudaDevAttrEccEnabled;
break;
default:
cerror = cudaErrorInvalidValue;
break;
src/hip_device.cpp
+15
View File
@@ -311,6 +311,21 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
case hipDeviceAttributeCooperativeMultiDeviceLaunch:
*pi = prop->cooperativeMultiDeviceLaunch;
break;
case hipDeviceAttributeMaxPitch:
*pi = prop->memPitch;
break;
case hipDeviceAttributeTextureAlignment:
*pi = prop->textureAlignment;
break;
case hipDeviceAttributeKernelExecTimeout:
*pi = prop->kernelExecTimeoutEnabled;
break;
case hipDeviceAttributeCanMapHostMemory:
*pi = prop->canMapHostMemory;
break;
case hipDeviceAttributeEccEnabled:
*pi = prop->ECCEnabled;
break;
default:
e = hipErrorInvalidValue;
break;
src/hip_hcc.cpp
+16 -1
View File
@@ -734,7 +734,6 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) {
err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_NAME, &archName);
prop->gcnArch = atoi(archName + 3);
DeviceErrorCheck(err);
// Get agent node
@@ -912,6 +911,22 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) {
prop->hdpMemFlushCntl = hdpinfo.HDP_MEM_FLUSH_CNTL;
prop->hdpRegFlushCntl = hdpinfo.HDP_REG_FLUSH_CNTL;
prop->memPitch = INT_MAX; //Maximum pitch in bytes allowed by memory copies (hardcoded 128 bytes in hipMallocPitch)
prop->textureAlignment = 0; //Alignment requirement for textures
prop->kernelExecTimeoutEnabled = 0; //no run time limit for running kernels on device
hsa_isa_t isa;
err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AGENT_INFO_ISA, &isa);
DeviceErrorCheck(err);
std::size_t isa_sz = 0u;
hsa_isa_get_info_alt(isa, HSA_ISA_INFO_NAME_LENGTH, &isa_sz);
std::string isa_name(isa_sz, '\0');
hsa_isa_get_info_alt(isa, HSA_ISA_INFO_NAME, &isa_name.front());
if (isa_name.find("sram-ecc") != std::string::npos)
prop->ECCEnabled = 1; //Device has ECC support Enabled
else
prop->ECCEnabled = 0; //Device has ECC support disabled
return e;
}
src/hip_memory.cpp
+30 -9
View File
@@ -1160,12 +1160,17 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToDevice, false);
} catch (ihipException& ex) {
@@ -1179,12 +1184,17 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, void* src, size_t sizeBytes) {
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToHost, false);
} catch (ihipException& ex) {
@@ -1198,12 +1208,17 @@ hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes) {
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyDtoD, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyDeviceToDevice, false);
} catch (ihipException& ex) {
@@ -1216,12 +1231,17 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte
hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) {
HIP_INIT_SPECIAL_API(hipMemcpyHtoH, (TRACE_MCMD), dst, src, sizeBytes);
hipError_t e = hipSuccess;
if (sizeBytes == 0) return ihipLogStatus(e);
if(dst==NULL || src==NULL){
return ihipLogStatus(hipErrorInvalidValue);
}
hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
hc::completion_future marker;
hipError_t e = hipSuccess;
try {
stream->locked_copySync((void*)dst, (void*)src, sizeBytes, hipMemcpyHostToHost, false);
} catch (ihipException& ex) {
@@ -1989,7 +2009,8 @@ hipError_t hipFree(void* ptr) {
#endif
am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, ptr);
if (status == AM_SUCCESS) {
if (amPointerInfo._hostPointer == NULL) {
/*if (amPointerInfo._hostPointer == NULL) */ //TODO: Fix it when there is proper managed memory support
{
if (HIP_SYNC_FREE) {
// Synchronize all devices, all streams
// to ensure all work has finished on all devices.
src/hip_module.cpp
+110
View File
@@ -480,6 +480,116 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim,
return ihipLogStatus(result);
}
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags);
hipError_t result;
if (numDevices > g_deviceCnt || launchParamsList == nullptr) {
return ihipLogStatus(hipErrorInvalidValue);
}
for (int i = 0; i < numDevices; ++i) {
if (!launchParamsList[i].stream->getDevice()->_props.cooperativeMultiDeviceLaunch) {
return ihipLogStatus(hipErrorInvalidConfiguration);
}
}
hipFunction_t* gwsKds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
hipFunction_t* kds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
if (kds == nullptr || gwsKds == nullptr) {
return ihipLogStatus(hipErrorNotInitialized);
}
// prepare all kernel descriptors for initializing the GWS and the main kernels per device
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
if (lp.stream == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorNotInitialized);
}
gwsKds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(&init_gws),
hip_impl::target_agent(lp.stream));
if (gwsKds[i] == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorInvalidValue);
}
hip_impl::kernargs_size_align gwsKargs = hip_impl::get_program_state().get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(&init_gws));
gwsKds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
gwsKargs.getHandle());
kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(lp.func),
hip_impl::target_agent(lp.stream));
if (kds[i] == nullptr) {
free(gwsKds);
free(kds);
return ihipLogStatus(hipErrorInvalidValue);
}
hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align(
reinterpret_cast<std::uintptr_t>(lp.func));
kds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
kargs.getHandle());
}
// lock all streams before launching the blit kernels for initializing the GWS and main kernels to each device
for (int i = 0; i < numDevices; ++i) {
LockedAccessor_StreamCrit_t streamCrit(launchParamsList[i].stream->criticalData(), false);
#if (__hcc_workweek__ >= 19213)
streamCrit->_av.acquire_locked_hsa_queue();
#endif
}
// launch the init_gws kernel to initialize the GWS followed by launching the main kernels for each device
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& lp = launchParamsList[i];
void *gwsKernelParam[1];
uint nwm1 = (lp.gridDim.x * lp.gridDim.y * lp.gridDim.z) - 1;
gwsKernelParam[0] = &nwm1;
result = ihipModuleLaunchKernel(tls, gwsKds[i], 1, 1, 1, 1, 1, 1,
0, lp.stream, gwsKernelParam, nullptr, nullptr, nullptr, 0, true);
if (result != hipSuccess) {
for (int j = 0; j < numDevices; ++j) {
launchParamsList[j].stream->criticalData().unlock();
#if (__hcc_workweek__ >= 19213)
launchParamsList[j].stream->criticalData()._av.release_locked_hsa_queue();
#endif
}
return ihipLogStatus(hipErrorLaunchFailure);
}
result = ihipModuleLaunchKernel(tls, kds[i],
lp.gridDim.x * lp.blockDim.x,
lp.gridDim.y * lp.blockDim.y,
lp.gridDim.z * lp.blockDim.z,
lp.blockDim.x, lp.blockDim.y,
lp.blockDim.z, lp.sharedMem,
lp.stream, lp.args, nullptr, nullptr, nullptr, 0,
true);
}
// unlock all streams
for (int i = 0; i < numDevices; ++i) {
launchParamsList[i].stream->criticalData().unlock();
#if (__hcc_workweek__ >= 19213)
launchParamsList[i].stream->criticalData()._av.release_locked_hsa_queue();
#endif
}
free(gwsKds);
free(kds);
return ihipLogStatus(result);
}
namespace hip_impl {
hsa_executable_t executable_for(hipModule_t hmod) {
return hmod->executable;
tests/hipify-clang/unit_tests/device/math_functions.cu
+15 -3
View File
@@ -1,5 +1,5 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// Test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
// Synthetic test to warn only on device functions umin and umax as unsupported, but not on user defined ones.
// ToDo: change lit testing in order to parse the output.
#define LEN 1024
@@ -8,9 +8,11 @@
#include <algorithm>
namespace my {
// user defined function
unsigned int umin(unsigned int arg1, unsigned int arg2) {
return (arg1 < arg2) ? arg1 : arg2;
}
// user defined function
unsigned int umax(unsigned int arg1, unsigned int arg2) {
return (arg1 > arg2) ? arg1 : arg2;
}
@@ -18,8 +20,16 @@ namespace my {
__global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsigned int u2)
{
unsigned int _umin = umin(u1, u2);
unsigned int _umax = umax(u1, u2);
// device function call (warn if unsupported)
unsigned int _umin = umin ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umax = umax ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umin_global = ::umin ( u1, u2 );
// device function call (warn if unsupported)
unsigned int _umax_global = ::umax(u1, u2);
if (_umin != _umin_global) return;
if (_umax != _umax_global) return;
int i = threadIdx.x;
A[i] = i + _umin;
B[i] = i + _umax;
@@ -29,7 +39,9 @@ __global__ void uint_arithm(float* A, float* B, float* C, unsigned int u1, unsig
int main() {
unsigned int u1 = 33;
unsigned int u2 = 34;
// user defined function call
unsigned int _min = my::umin(u1, u2);
// user defined function call
unsigned int _max = my::umax(u1, u2);
float *A, *B, *C;
// CHECK: hipMalloc((void**)&A, SIZE);
tests/src/deviceLib/hip_threadfence_system.cpp
+11 -2
View File
@@ -71,12 +71,21 @@ int main() {
}
volatile int* data;
HIP_ASSERT(hipHostMalloc(&data, sizeof(int), hipHostMallocCoherent));
if (hipHostMalloc(&data, sizeof(int), hipHostMallocCoherent) != hipSuccess) {
warn("Memory allocation failed. Skip test. Is SVM atomic supported?")
passed();
return 0;
}
constexpr int init_data = 1000;
*data = init_data;
volatile int* flag;
HIP_ASSERT(hipHostMalloc(&flag, sizeof(int), hipHostMallocCoherent));
if (hipHostMalloc(&flag, sizeof(int), hipHostMallocCoherent) != hipSuccess) {
warn("Memory allocation failed. Skip test. Is SVM atomic supported?")
passed();
return 0;
}
*flag = 0;
// number of rounds per device
tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+5
View File
@@ -140,5 +140,10 @@ int main(int argc, char* argv[]) {
CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpMemFlushCntl, props.hdpMemFlushCntl));
CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpRegFlushCntl, props.hdpRegFlushCntl));
#endif
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxPitch, props.memPitch));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeTextureAlignment, props.textureAlignment));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeKernelExecTimeout, props.kernelExecTimeoutEnabled));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCanMapHostMemory, props.canMapHostMemory));
CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeEccEnabled, props.ECCEnabled));
passed();
};
tests/src/runtimeApi/memory/hipHostMalloc.cpp
+12 -8
View File
@@ -176,15 +176,19 @@ int main() {
if (1) {
int* A = nullptr;
HIPCHECK(hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent));
const char* ptrType = "coherent";
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_EVENT, ptrType);
if (hipHostMalloc((void**)&A, sizeBytes, hipHostMallocCoherent) == hipSuccess) {
const char* ptrType = "coherent";
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToDevice, SYNC_EVENT, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_EVENT, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_DEVICE, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_STREAM, ptrType);
CheckHostPointer(numElements, A, hipEventReleaseToSystem, SYNC_EVENT, ptrType);
}
else {
warn("Coherence memory allocation failed. Is SVM atomic supported?")
}
}
tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+164
View File
@@ -0,0 +1,164 @@
/*
Copyright (c) 2019 - 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.
*/
// Simple test for hipLaunchCooperativeKernel API.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include "hip/hcc_detail/device_library_decls.h"
#include "hip/hcc_detail/hip_cooperative_groups.h"
#include <iostream>
#include <chrono>
#include "test_common.h"
using namespace std::chrono;
namespace cg = cooperative_groups;
const static uint BufferSizeInDwords = 448 * 1024 * 1024;
__global__ void test_gws(uint* buf, uint bufSize, long* tmpBuf, long* result)
{
extern __shared__ long tmp[];
uint offset = blockIdx.x * blockDim.x + threadIdx.x;
uint stride = gridDim.x * blockDim.x;
cg::grid_group gg = cg::this_grid();
long sum = 0;
for (uint i = offset; i < bufSize; i += stride) {
sum += buf[i];
}
tmp[threadIdx.x] = sum;
__syncthreads();
if (threadIdx.x == 0) {
sum = 0;
for (uint i = 0; i < blockDim.x; i++) {
sum += tmp[i];
}
tmpBuf[blockIdx.x] = sum;
}
gg.sync();
if (offset == 0) {
for (uint i = 1; i < gridDim.x; ++i) {
sum += tmpBuf[i];
}
*result = sum;
}
}
int main() {
float *A, *B, *Ad, *Bd;
uint* dA;
long* dB;
long* dC;
uint32_t* init = new uint32_t[BufferSizeInDwords];
for (uint32_t i = 0; i < BufferSizeInDwords; ++i) {
init[i] = i;
}
hipDeviceProp_t deviceProp;
hipGetDeviceProperties(&deviceProp, 0);
if (!deviceProp.cooperativeLaunch) {
std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n";
passed();
return 0;
}
std::cout << "info: running on bus 0x" << deviceProp.pciBusID << " " << deviceProp.name << "\n";
size_t SIZE = BufferSizeInDwords * sizeof(uint);
HIPCHECK(hipMalloc((void**)&dA, SIZE));
HIPCHECK(hipHostMalloc((void**)&dC, sizeof(long)));
HIPCHECK(hipMemcpy(dA, init, SIZE, hipMemcpyHostToDevice));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
dim3 dimBlock = dim3(1);
dim3 dimGrid = dim3(1);
int numBlocks = 0;
uint workgroups[2] = {32, 64};
system_clock::time_point start = system_clock::now();
for (uint i = 0; i < 2; ++i) {
dimBlock.x = workgroups[i];
// Calculate the device occupancy to know how many blocks can be run concurrently
hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));
dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
HIPCHECK(hipMalloc((void**)&dB, dimGrid.x * sizeof(long)));
void *params[4];
params[0] = (void*)&dA;
params[1] = (void*)&BufferSizeInDwords;
params[2] = (void*)&dB;
params[3] = (void*)&dC;
std::cout << "Testing with grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n";
HIPCHECK(hipLaunchCooperativeKernel(test_gws, dimGrid, dimBlock, params, dimBlock.x * sizeof(long), stream));
HIPCHECK(hipMemcpy(init, dC, sizeof(long), hipMemcpyDeviceToHost));
if (*dC != (((long)(BufferSizeInDwords) * (BufferSizeInDwords - 1)) / 2)) {
std::cout << "Data validation failed for grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n";
HIPCHECK(hipStreamDestroy(stream));
hipFree(dC);
hipFree(dB);
hipFree(dA);
delete [] init;
std::cout << "Test failed! \n";
return 0;
} else {
std::cout << "info: data validated!\n";
}
}
system_clock::time_point end = system_clock::now();
std::chrono::duration<double> elapsed_seconds = end - start;
std::time_t end_time = std::chrono::system_clock::to_time_t(end);
std::cout << "finished computation at " << std::ctime(&end_time) <<
"elapsed time: " << elapsed_seconds.count() << "s\n";
HIPCHECK(hipStreamDestroy(stream));
hipFree(dC);
hipFree(dB);
hipFree(dA);
delete [] init;
passed();
return 0;
}