diff --git a/projects/hip/tests/catch/unit/rtc/CMakeLists.txt b/projects/hip/tests/catch/unit/rtc/CMakeLists.txt index d52d3f9288..8e2423bef4 100644 --- a/projects/hip/tests/catch/unit/rtc/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/rtc/CMakeLists.txt @@ -3,6 +3,15 @@ set(TEST_SRC saxpy.cc ) +# AMD only tests +set(AMD_TEST_SRC + customOptions.cc +) + +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) +endif() + # Create shared lib of all tests add_library(RTC SHARED EXCLUDE_FROM_ALL ${TEST_SRC}) diff --git a/projects/hip/tests/catch/unit/rtc/customOptions.cc b/projects/hip/tests/catch/unit/rtc/customOptions.cc new file mode 100644 index 0000000000..040b567100 --- /dev/null +++ b/projects/hip/tests/catch/unit/rtc/customOptions.cc @@ -0,0 +1,148 @@ +#include + +#include +#include + + +#include +#include +#include +#include +#include +#include +#include + +static constexpr auto program{ + R"( +extern "C" +__global__ void kernel(int* a) { + // C++17 feature + if (int j = 10; *a % 2 == 0) + *a = 10 + j; + else + *a = 20 + j; +} +)"}; + +TEST_CASE("Unit_hiprtc_cpp17") { + using namespace std; + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + program, // buffer + "program.cu", // name + 0, nullptr, nullptr); + hipDeviceProp_t props; + int device = 0; + HIP_CHECK(hipGetDeviceProperties(&props, device)); +#ifdef __HIP_PLATFORM_AMD__ + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; +#else + std::string sarg = std::string("--fmad=false"); +#endif + const char* options[] = {sarg.c_str(), "-std=c++17", "-Werror"}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 3, options)}; + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + std::cout << log << '\n'; + } + hiprtcDestroyProgram(&prog); + REQUIRE(compileResult == HIPRTC_SUCCESS); +} + +static constexpr const char template_kernel[]{R"( +template struct complex { + public: + typedef T value_type; + inline __host__ __device__ complex(const T& re, const T& im); + __host__ __device__ inline complex& operator*=(const complex z); + __host__ __device__ inline T real() const volatile { return m_data[0]; } + __host__ __device__ inline T imag() const volatile { return m_data[1]; } + __host__ __device__ inline T real() const { return m_data[0]; } + __host__ __device__ inline T imag() const { return m_data[1]; } + __host__ __device__ inline void real(T re) volatile { m_data[0] = re; } + __host__ __device__ inline void imag(T im) volatile { m_data[1] = im; } + __host__ __device__ inline void real(T re) { m_data[0] = re; } + __host__ __device__ inline void imag(T im) { m_data[1] = im; } + + private: + T m_data[2]; +}; +template inline __host__ __device__ complex::complex(const T& re, const T& im) { + real(re); + imag(im); +} +template +__host__ __device__ inline complex& complex::operator*=(const complex z) { + *this = *this * z; + return *this; +} +template +__host__ __device__ inline complex operator*(const complex& lhs, const complex& rhs) { + return complex(lhs.real() * rhs.real() - lhs.imag() * rhs.imag(), + lhs.real() * rhs.imag() + lhs.imag() * rhs.real()); +} + +template __global__ void my_sqrt(T* input, int N) { + unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; + if (x < N) { + input[x] *= input[x]; + } +} +)"}; + +TEST_CASE("Unit_hiprtc_namehandling") { + using namespace std; + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + template_kernel, // buffer + "template_kernel.cu", // name + 0, nullptr, nullptr); + hipDeviceProp_t props; + int device = 0; + HIP_CHECK(hipGetDeviceProperties(&props, device)); +#ifdef __HIP_PLATFORM_AMD__ + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; +#else + std::string sarg = std::string("--fmad=false"); +#endif + const char* options[] = {sarg.c_str()}; + + std::vector name_expressions; + name_expressions.push_back("my_sqrt"); + name_expressions.push_back("my_sqrt"); + name_expressions.push_back("my_sqrt>"); + name_expressions.push_back("my_sqrt >"); + for (size_t i = 0; i < name_expressions.size(); i++) { + REQUIRE(HIPRTC_SUCCESS == hiprtcAddNameExpression(prog, name_expressions[i].c_str())); + } + + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + std::cout << log << '\n'; + } + + std::map mangled_names; + + for (size_t i = 0; i < name_expressions.size(); i++) { + const char* mangled_instantiation_cstr; + REQUIRE(HIPRTC_SUCCESS == hiprtcGetLoweredName(prog, name_expressions[i].c_str(), &mangled_instantiation_cstr)); + + std::string mangled_name_str = mangled_instantiation_cstr; + mangled_names[name_expressions[i]] = mangled_name_str; + REQUIRE(mangled_name_str.size() > 0); + } + + // Match the last two names + REQUIRE(mangled_names["my_sqrt>"] == mangled_names["my_sqrt >"]); + + hiprtcDestroyProgram(&prog); + REQUIRE(compileResult == HIPRTC_SUCCESS); +}