diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index fbb135757b..2a87271a60 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1646,13 +1646,13 @@ sub transformKernelLaunch { } sub transformHostFunctions { - my $m = 0; + my $k = 0; foreach $func ( "hipMemcpyToSymbol", "hipMemcpyToSymbolAsync" ) { - $m += s/(?\($2\),/g + $k += s/(?\($2\),/g } foreach $func ( "hipFuncGetAttributes" ) { - $m += s/(?\($3\)$4/g; + $k += s/(?\($3\)$4/g; } - return $m; + return $k; } sub countSupportedDeviceFunctions { - my $m = 0; + my $k = 0; foreach $func ( "__brev", "__brevll", @@ -1935,6 +1935,17 @@ sub countSupportedDeviceFunctions { "atanf", "atanh", "atanhf", + "atomicAdd", + "atomicAnd", + "atomicCAS", + "atomicDec", + "atomicExch", + "atomicInc", + "atomicMax", + "atomicMin", + "atomicOr", + "atomicSub", + "atomicXor", "cbrt", "cbrtf", "ceil", @@ -2134,15 +2145,15 @@ sub countSupportedDeviceFunctions { my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g; my $mt = m/($func)\s*\(\s*.*\s*\)/g; if ($mt && !$mt_namespace) { - $m += $mt; + $k += $mt; } } - return $m; + return $k; } sub warnUnsupportedDeviceFunctions { my $line_num = shift; - my $m = 0; + my $k = 0; foreach $func ( "_Pow_int", "__brkpt", @@ -2250,7 +2261,6 @@ sub warnUnsupportedDeviceFunctions { "__vsubus4", "_fdsign", "_ldsign", - "atomicAdd", "float2int", "float_as_int", "float_as_uint", @@ -2276,11 +2286,11 @@ sub warnUnsupportedDeviceFunctions { my $mt_namespace = m/(\w+)::($func)\s*\(\s*.*\s*\)/g; my $mt = m/($func)\s*\(\s*.*\s*\)/g; if ($mt && !$mt_namespace) { - $m += $mt; + $k += $mt; print STDERR " warning: $fileName:$line_num: unsupported device function \"$func\": $_\n"; } } - return $m; + return $k; } # Count of transforms in all files diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp index b61e4bd73c..bd555156c3 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Device_functions.cpp @@ -599,9 +599,21 @@ const std::map CUDA_DEVICE_FUNC_MAP{ {"h2exp10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"h2cos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, {"h2sin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, - {"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_down_sync",{"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, - {"atomicAdd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"__shfl_down_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + {"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}}, + // atomic functions + {"atomicAdd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicSub", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicExch", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicMin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicMax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicInc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicAnd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicOr", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicXor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, + {"atomicCAS", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}}, }; diff --git a/projects/hip/tests/hipify-clang/unit_tests/device/atomics.cu b/projects/hip/tests/hipify-clang/unit_tests/device/atomics.cu new file mode 100644 index 0000000000..d1c1fad401 --- /dev/null +++ b/projects/hip/tests/hipify-clang/unit_tests/device/atomics.cu @@ -0,0 +1,284 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args "-Xclang" "-fcuda-allow-variadic-functions" + +/* +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. +*/ + +// CHECK: #include "hip/hip_runtime.h" +#include "cuda_runtime.h" + +#include +#include +#include +#include +#include +#include + +#define EXIT_WAIVED 2 + +const char* sampleName = "hipSimpleAtomicsTest"; + +using namespace std; +// Auto-Verification Code +bool testResult = true; + +bool computeGoldBitwise(...) { + return true; +} + +template{}>::type* = nullptr> +bool computeGoldBitwise(T* gpuData, int len) { + T val = 0xff; + for (int i = 0; i < len; ++i) { + // 9th element should be 1 + val &= (2 * i + 7); + } + if (val != gpuData[8]) { + printf("atomicAnd failed\n"); + return false; + } + + val = 0; + for (int i = 0; i < len; ++i) { + // 10th element should be 0xff + val |= (1 << i); + } + if (val != gpuData[9]) { + printf("atomicOr failed\n"); + return false; + } + + val = 0xff; + for (int i = 0; i < len; ++i) { + // 11th element should be 0xff + val ^= i; + } + if (val != gpuData[10]) { + printf("atomicXor failed\n"); + return false; + } + + return true; +} + +template +bool computeGold(T* gpuData, int len) { + T val = 0; + for (int i = 0; i < len; ++i) { + val += 10; + } + if (val != gpuData[0]) { + printf("atomicAdd failed\n"); + return false; + } + + val = 0; + for (int i = 0; i < len; ++i) { + val -= 10; + } + if (val != gpuData[1]) { + printf("atomicSub failed\n"); + return false; + } + + bool found = false; + for (T i = 0; i < len; ++i) { + // third element should be a member of [0, len) + if (i == gpuData[2]) { + found = true; + break; + } + } + if (!found) { + printf("atomicExch failed\n"); + return false; + } + + val = -(1 << 8); + for (T i = 0; i < len; ++i) { + // fourth element should be len-1 + val = max(val, i); + } + if (val != gpuData[3]) { + printf("atomicMax failed\n"); + return false; + } + + val = 1 << 8; + for (T i = 0; i < len; ++i) { + val = min(val, i); + } + if (val != gpuData[4]) { + printf("atomicMin failed\n"); + return false; + } + + int limit = 17; + val = 0; + for (int i = 0; i < len; ++i) { + val = (val >= limit) ? 0 : val + 1; + } + if (val != gpuData[5]) { + printf("atomicInc failed\n"); + return false; + } + + limit = 137; + val = 0; + for (int i = 0; i < len; ++i) { + val = ((val == 0) || (val > limit)) ? limit : val - 1; + } + if (val != gpuData[6]) { + printf("atomicDec failed\n"); + return false; + } + + found = false; + for (T i = 0; i < len; ++i) { + // eighth element should be a member of [0, len) + if (i == gpuData[7]) { + found = true; + break; + } + } + if (!found) { + printf("atomicCAS failed\n"); + return false; + } + + return computeGoldBitwise(gpuData, len); +} + +__device__ +void testKernelExch(...) {} + +template{}>::type* = nullptr> +__device__ +void testKernelExch(T* g_odata) { + // access thread id + const T tid = blockDim.x * blockIdx.x + threadIdx.x; + // Atomic exchange + atomicExch(&g_odata[2], tid); +} + +__device__ +void testKernelSub(...) {} + +template< + typename T, + typename enable_if< + is_same{} || is_same{}>::type* = nullptr> +__device__ +void testKernelSub(T* g_odata) { + // Atomic subtraction (final should be 0) + atomicSub(&g_odata[1], 10); +} + +__device__ +void testKernelIntegral(...) {} + +template{}>::type* = nullptr> +__device__ +void testKernelIntegral(T* g_odata) { + // access thread id + const T tid = blockDim.x * blockIdx.x + threadIdx.x; + // Atomic maximum + atomicMax(&g_odata[3], tid); + // Atomic minimum + atomicMin(&g_odata[4], tid); + // Atomic increment (modulo 17+1) + atomicInc((unsigned int*)&g_odata[5], 17); + // Atomic decrement + atomicDec((unsigned int*)&g_odata[6], 137); + // Atomic compare-and-swap + atomicCAS(&g_odata[7], tid - 1, tid); + // Atomic AND + atomicAnd(&g_odata[8], 2 * tid + 7); + // Atomic OR + atomicOr(&g_odata[9], 1 << tid); + // Atomic XOR + atomicXor(&g_odata[10], tid); + testKernelSub(g_odata); +} + +template +__global__ void testKernel(T* g_odata) { + // Atomic addition + atomicAdd(&g_odata[0], 10); + testKernelIntegral(g_odata); + testKernelExch(g_odata); +} + +template +void runTest() { + // CHECK: hipDeviceProp_t deviceProp; + cudaDeviceProp deviceProp; + deviceProp.major = 0; + deviceProp.minor = 0; + int dev = 0; + // CHECK: hipGetDeviceProperties(&deviceProp, dev); + cudaGetDeviceProperties(&deviceProp, dev); + // Statistics about the GPU device + printf( + "> GPU device has %d Multi-Processors, " + "SM %d.%d compute capabilities\n\n", + deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor); + unsigned int numThreads = 256; + unsigned int numBlocks = 64; + unsigned int numData = 11; + unsigned int memSize = sizeof(T) * numData; + + // Allocate mem for the result on host side + T* hOData = (T*)malloc(memSize); + // Initialize the memory + for (unsigned int i = 0; i < numData; i++) hOData[i] = 0; + // To make the AND and XOR tests generate something other than 0... + hOData[8] = hOData[10] = 0xff; + // Allocate device memory for result + T* dOData; + // CHECK: hipMalloc((void**)&dOData, memSize); + cudaMalloc((void**)&dOData, memSize); + // Copy host memory to device to initialize to zero + // CHECK: hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice); + cudaMemcpy(dOData, hOData, memSize, cudaMemcpyHostToDevice); + // Execute the kernel + // CHECK: hipLaunchKernelGGL((testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData); + testKernel<<>>(dOData); + // Copy result from device to host + // CHECK: hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost); + cudaMemcpy(hOData, dOData, memSize, cudaMemcpyDeviceToHost); + // Compute reference solution + testResult = computeGold(hOData, numThreads * numBlocks); + // Cleanup memory + free(hOData); + // CHECK: hipFree(dOData); + cudaFree(dOData); +} + +int main(int argc, char** argv) { + printf("%s starting...\n", sampleName); + runTest(); + runTest(); + runTest(); + runTest(); + runTest(); + // CHECK: hipDeviceReset(); + cudaDeviceReset(); + printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!"); + exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); +}