RadeonOpenCompute
HSADevice::CreateKernel(): Unable to create kernel
Sample_Code&logs.zip Below error is seen while executing thrust testing testcases on HIP/ROCm path. HSADevice::CreateKernel(): Unable to create kernel
Issue Description: Unable to create kernel is seen when device kernel is launched from a test framework of thrust.
Non Working case:
While passing thrust::device_vector
Code snippet to reproduce above issue:
#include <unittest/unittest.h> template <class Vector> void test() { typedef typename Vector::value_type T; Vector input(3);
} DECLARE_VECTOR_UNITTEST(test)
Working-case:
Instead of using macro DECLARE_VECTOR_UNITTEST, if we use the below main() definition issue doesn't exists.
Code snippet to reproduce above issue:
#include <unittest/unittest.h> template <class Vector> void test() { typedef typename Vector::value_type T; Vector input(3); }
int main()
{
test<thrust::device_vector
This case works fine for HIP/CUDA using nvcc, but issue is reproduced for HIP/ROCm using hcc
Steps to reproduce working case: $ git clone https://github.com/ROCmSoftwarePlatform/Thrust.git $ cd Thrust $ export HIP_PLATFORM=hcc (For HCC Platform ) $ cd testing $ /opt/rocm/bin/hipcc test_sample1.cpp -I. -I../ -g -o test_sample1.out (sample code test_sample1.cpp is attached) $ ./test_sample.out (Executable to be run on AMD hardware)
Steps to reproduce Non-working case: $ git clone https://github.com/ROCmSoftwarePlatform/Thrust.git $ cd Thrust $ export HIP_PLATFORM=hcc (For HCC Platform ) $ cd testing $ /opt/rocm/bin/hipcc test_sample2.cpp testframework.cpp -I. -I../ -g -o test_sample2.out (sample code test_sample2.cpp is attached) $ ./test_sample.out (Executable to be run on AMD hardware)
Configuration Details Package: rocm-dev Status: install ok installed Priority: optional Section: devel Installed-Size: 13 Maintainer: Advanced Micro Devices Inc. Architecture: amd64 Version: 1.7.60 Depends: hsa-rocr-dev, hsa-ext-rocr-dev, rocm-device-libs, rocm-utils, hcc, hip_base, hip_doc, hip_hcc, hip_samples, rocm-smi, hsakmt-roct, hsakmt-roct-dev, hsa-amd-aqlprofile Description: Radeon Open Compute (ROCm) Runtime software stack Homepage: https://github.com/RadeonOpenCompute/ROCm
HIP version : 1.4.17494 == hipconfig HIP_PATH : /opt/rocm HIP_PLATFORM : hcc CPP_CONFIG : -D__HIP_PLATFORM_HCC__= -I/opt/rocm/include -I/opt/rocm/hcc/include
== hcc HSA_PATH : /opt/rocm/hsa HCC_HOME : /opt/rocm/hcc HCC clang version 6.0.0 (ssh://gerritgit/compute/ec/hcc-tot/clang 42ceed861a212d9bd0aef883ee7981144f3ecc02) (ssh://gerritgit/compute/ec/hcc-tot/llvm 23e086be6f627e6e983c6789d2e77da6bf85ebb6) (based on HCC 1.1.17493-2f85d8a-42ceed8-23e086b ) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /opt/rocm/hcc/bin Can't exec "/opt/rocm/hcc/compiler/bin/llc": No such file or directory at /opt/rocm/bin/hipconfig line 132. HCC-cxxflags : -hc -std=c++amp -I/opt/rocm/hcc-1.0/include -I/opt/rocm/includeHCC-ldflags : -hc -std=c++amp -L/opt/rocm/hcc-1.0/lib -Wl,--rpath=/opt/rocm/hcc-1.0/lib -ldl -lm -lpthread -lunwind -lhc_am -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive
== Linux Kernel Hostname : tcs-amd Linux tcs-amd 4.4.0-104-generic #127-Ubuntu SMP Mon Dec 11 12:16:42 UTC 2017 x86_64 x86_64 x86_64 GNU/Linux No LSB modules are available. Distributor ID: Ubuntu Description: Ubuntu 16.04.3 LTS Release: 16.04 Codename: xenial
@sriharikarnam In the attachment there is no definition of DECLARE_VECTOR_UNITTEST. Please help provide it, or prepare a self-contained example.
@whchung Definition of DECLARE_VECTOR_UNITTEST is available in below file at line # 162 https://github.com/ROCmSoftwarePlatform/Thrust/blob/master/testing/unittest/testframework.h
// Macro to create host and device versions of a
// unit test for a couple data types
#define DECLARE_VECTOR_UNITTEST(VTEST)
void VTEST##Host(void) { VTEST< thrust::host_vector
void VTEST##Device(void) { VTEST< thrust::device_vector
@sriharikarnam Could you help provide a minimal self-contained example? Without any knowledge of your application I fail to see where kernels are produced and launched.
@whchung Thanks for your inputs. We have identified a temporary workaround for the issue. We are investigating further on this issue.
Attached working case log max_element_log.txt