Update bugs - Add CreateKernel, new signature for static kerns.

[ROCm/hip commit: 4f363df159]
This commit is contained in:
Ben Sander
2017-04-16 14:22:48 -05:00
parent 2137285ffc
commit dc8636a178
+19
View File
@@ -11,7 +11,13 @@
### Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**`
Some common code practices may lead to hipcc generating a error with the form :
```
undefined reference to `__hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T_
```
Or:
```
error: weak declaration cannot have internal linkage
```
Suggested workarounds:
- Avoid use of static with kernel definition:
@@ -26,6 +32,19 @@ namespace {
}
```
### Can't find kernels inside dynamic linked library
HCC requires use of the "-Bdynamic" flag when creating a dynamic library which contains kernels. The dynamic flag causes the symbols to be created with a signature which allows HCC to discover and load the kernels in the dynamic library. This flag is often not set by default and must be added to the link step of the library. If not done, HCC will be unable to find the kernels defined in the library, and will emit a message such as:
```
HSADevice::CreateKernel(): Unable to create kernel"
```
To correct, add the following flag to hcc or hipcc:
```
$ hipcc -Wl,-Bsymbolic ...
```
### What is the current limitation of HIP Generic Grid Launch method?
1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak)));