Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Problems with making HEMI_DEFINE_CONSTANT static on host #2

Closed
harrism opened this issue Oct 15, 2012 · 1 comment
Closed

Problems with making HEMI_DEFINE_CONSTANT static on host #2

harrism opened this issue Oct 15, 2012 · 1 comment
Assignees
Labels

Comments

@harrism
Copy link
Owner

harrism commented Oct 15, 2012

#define HEMI_DEFINE_CONSTANT(def, value) __constant__ def ## _devconst = value; static def ## _hostconst = value;

With separate compilation, an extern linkage constant variable defined in one translation unit may be referenced by another translation unit. Forcing the host version of the constant to have internal linkage (“static”) won’t work for this scenario. Also, “static” linkage variables cannot be template arguments, which would break the following code:

__constant__ int XXX = 200;

template <int *addr>
__device__  int getValue(void) { return *addr; }

__global__ void kern(int *ptr) {  *ptr = getValue<&XXX>(); }

The snippet above compiles fine, but changing the linkage of “XXX” to static and recompiling fails:

j1.cu(6): error: a template argument may not reference a non-external entity

Perhaps the macro should have an argument to specify the linkage?

@ghost ghost assigned harrism Oct 15, 2012
@harrism
Copy link
Owner Author

harrism commented Oct 18, 2012

Fixed by cd163b8

@harrism harrism closed this as completed Oct 18, 2012
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

No branches or pull requests

1 participant