diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 50808563ff..b5fc365955 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -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; diff --git a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 67bf19b07e..cbd4ccb818 100644 --- a/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -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`* | | diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 2b5ac181cc..dfb80a3537 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -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`* | | | diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 5a1178ec01..fdb9e080f7 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -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 ``` ### 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`: diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 7cfe2adc52..f88086f2cb 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -305,7 +305,7 @@ const std::map 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 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 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 diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp new file mode 100644 index 0000000000..75afc0a534 --- /dev/null +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -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 +#include +#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& 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 perlStreamPtr = std::unique_ptr(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; + } +} diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp new file mode 100644 index 0000000000..5ffd84e0b0 --- /dev/null +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Python.cpp @@ -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 +#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 pythonStreamPtr = std::unique_ptr(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; + } +} diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp index 8ebefb38cf..fda9b34b2c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp @@ -234,24 +234,24 @@ const std::map 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 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 diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h b/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h new file mode 100644 index 0000000000..c585a47d47 --- /dev/null +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h @@ -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); +} diff --git a/projects/hip/hipify-clang/src/StringUtils.cpp b/projects/hip/hipify-clang/src/StringUtils.cpp index a8edf3661f..78d0f8113e 100644 --- a/projects/hip/hipify-clang/src/StringUtils.cpp +++ b/projects/hip/hipify-clang/src/StringUtils.cpp @@ -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(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(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(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(); +} + diff --git a/projects/hip/hipify-clang/src/StringUtils.h b/projects/hip/hipify-clang/src/StringUtils.h index b18c864244..39ce726d1b 100644 --- a/projects/hip/hipify-clang/src/StringUtils.h +++ b/projects/hip/hipify-clang/src/StringUtils.h @@ -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); diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index d6a53b170d..4b426b41f7 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -27,208 +27,18 @@ THE SOFTWARE. #include #include #include -#include #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(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(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(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 perlStreamPtr = std::unique_ptr(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 pythonStreamPtr = std::unique_ptr(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 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"; diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index c93aa67994..35b08be2ff 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -40,7 +40,6 @@ THE SOFTWARE. #endif #include -#include #include #include #include diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index 1f7183acf0..e3e5a14aad 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -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; ///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; diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index 2d74985b64..6406b48c2d 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -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; diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 899959a7ee..2a8610db5e 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -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; } diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 2deedbc635..10eb317f26 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -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. diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 2548660a72..4c07b9777b 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -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(malloc(sizeof(hipFunction_t) * numDevices)); + hipFunction_t* kds = reinterpret_cast(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(&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(&init_gws)); + gwsKds[i]->_kernarg_layout = *reinterpret_cast>*>( + gwsKargs.getHandle()); + + + kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast(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(lp.func)); + kds[i]->_kernarg_layout = *reinterpret_cast>*>( + 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; diff --git a/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu b/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu index 3bc1c1e51d..c833ca0182 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu +++ b/projects/hip/tests/hipify-clang/unit_tests/device/math_functions.cu @@ -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 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); diff --git a/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp b/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp index 24732573e6..39b246b7bc 100644 --- a/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp +++ b/projects/hip/tests/src/deviceLib/hip_threadfence_system.cpp @@ -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 diff --git a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index c69eb93a10..3c2bf567d0 100644 --- a/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -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(); }; diff --git a/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp b/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp index 3c902b3784..04dd74b545 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp +++ b/projects/hip/tests/src/runtimeApi/memory/hipHostMalloc.cpp @@ -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?") + } } diff --git a/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp new file mode 100644 index 0000000000..7680229855 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp @@ -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 +#include +#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(&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 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; +}