device function pointer as template parameter

up vote
0
down vote

favorite

I have a template struct for some reasons (beyond the scope of this question) :

template<typename T, __device__ retV (*funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

which can be used this way:

__device__ double increment(double x) {

    return x + 1.0;

}



__constant__ func<double, double, &increment> myfunc;



__device__ double apply(double x) 

{

    return myfunc.invoke(x);

}

This works well with nvcc (cuda 10.0), but fails with nvrtc (JIT compilation) with the following error:

error: attributes may not appear here

How should I modify this code to make it work with nvrtc?
Or should I add flags in my command-line?

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

why the downvote please? Questions seems legit
– Regis Portalez
Nov 9 at 15:58

Which CUDA toolkit and host compiler are you using?
– talonmies
Nov 10 at 8:05

Cuda 10.0 on windows. Visual 2017 latest. Same on Ubuntu 18.04. Likely gcc 5.4
– Regis Portalez
Nov 10 at 9:08

I'm really surprised it compiles. Memory space attributes should never be applied to class data. The only time it is valid is for member functions. nvrtc is clang based, as far as I can tell, which might explain the difference, but I'll wager nvrtc is correct in this case
– talonmies
Nov 10 at 9:25

add a comment |

up vote
0
down vote

favorite

I have a template struct for some reasons (beyond the scope of this question) :

template<typename T, __device__ retV (*funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

which can be used this way:

__device__ double increment(double x) {

    return x + 1.0;

}



__constant__ func<double, double, &increment> myfunc;



__device__ double apply(double x) 

{

    return myfunc.invoke(x);

}

This works well with nvcc (cuda 10.0), but fails with nvrtc (JIT compilation) with the following error:

error: attributes may not appear here

How should I modify this code to make it work with nvrtc?
Or should I add flags in my command-line?

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

why the downvote please? Questions seems legit
– Regis Portalez
Nov 9 at 15:58

Which CUDA toolkit and host compiler are you using?
– talonmies
Nov 10 at 8:05

Cuda 10.0 on windows. Visual 2017 latest. Same on Ubuntu 18.04. Likely gcc 5.4
– Regis Portalez
Nov 10 at 9:08

I'm really surprised it compiles. Memory space attributes should never be applied to class data. The only time it is valid is for member functions. nvrtc is clang based, as far as I can tell, which might explain the difference, but I'll wager nvrtc is correct in this case
– talonmies
Nov 10 at 9:25

add a comment |

up vote
0
down vote

favorite

I have a template struct for some reasons (beyond the scope of this question) :

template<typename T, __device__ retV (*funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

which can be used this way:

__device__ double increment(double x) {

    return x + 1.0;

}



__constant__ func<double, double, &increment> myfunc;



__device__ double apply(double x) 

{

    return myfunc.invoke(x);

}

This works well with nvcc (cuda 10.0), but fails with nvrtc (JIT compilation) with the following error:

error: attributes may not appear here

How should I modify this code to make it work with nvrtc?
Or should I add flags in my command-line?

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

I have a template struct for some reasons (beyond the scope of this question) :

template<typename T, __device__ retV (*funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

which can be used this way:

__device__ double increment(double x) {

    return x + 1.0;

}



__constant__ func<double, double, &increment> myfunc;



__device__ double apply(double x) 

{

    return myfunc.invoke(x);

}

This works well with nvcc (cuda 10.0), but fails with nvrtc (JIT compilation) with the following error:

error: attributes may not appear here

How should I modify this code to make it work with nvrtc?
Or should I add flags in my command-line?

c++ templates cuda nvrtc

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

edited Nov 9 at 13:03

asked Nov 9 at 12:33

Regis Portalez

2,80811732

asked Nov 9 at 12:33

Regis Portalez

2,80811732

asked Nov 9 at 12:33

Regis Portalez

2,80811732

why the downvote please? Questions seems legit
– Regis Portalez
Nov 9 at 15:58

Which CUDA toolkit and host compiler are you using?
– talonmies
Nov 10 at 8:05

Cuda 10.0 on windows. Visual 2017 latest. Same on Ubuntu 18.04. Likely gcc 5.4
– Regis Portalez
Nov 10 at 9:08

I'm really surprised it compiles. Memory space attributes should never be applied to class data. The only time it is valid is for member functions. nvrtc is clang based, as far as I can tell, which might explain the difference, but I'll wager nvrtc is correct in this case
– talonmies
Nov 10 at 9:25

add a comment |

why the downvote please? Questions seems legit
– Regis Portalez
Nov 9 at 15:58

Which CUDA toolkit and host compiler are you using?
– talonmies
Nov 10 at 8:05

Cuda 10.0 on windows. Visual 2017 latest. Same on Ubuntu 18.04. Likely gcc 5.4
– Regis Portalez
Nov 10 at 9:08

I'm really surprised it compiles. Memory space attributes should never be applied to class data. The only time it is valid is for member functions. nvrtc is clang based, as far as I can tell, which might explain the difference, but I'll wager nvrtc is correct in this case
– talonmies
Nov 10 at 9:25

why the downvote please? Questions seems legit
– Regis Portalez
Nov 9 at 15:58

Which CUDA toolkit and host compiler are you using?
– talonmies
Nov 10 at 8:05

Cuda 10.0 on windows. Visual 2017 latest. Same on Ubuntu 18.04. Likely gcc 5.4
– Regis Portalez
Nov 10 at 9:08

I'm really surprised it compiles. Memory space attributes should never be applied to class data. The only time it is valid is for member functions. nvrtc is clang based, as far as I can tell, which might explain the difference, but I'll wager nvrtc is correct in this case
– talonmies
Nov 10 at 9:25

add a comment |

1 Answer
1

active

oldest

votes

up vote
1
down vote

Well, answer is quite easy :

__device__ attribute is misplaced (as indicated by the compiler). The funcstruct should look like:

template<typename T, retV (* __device__  funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

But I don't know why nvcc and nvrtc have different expectations on this.

answered Nov 9 at 13:32

Regis Portalez

2,80811732

add a comment |

Your Answer

StackExchange.ifUsing("editor", function () {
StackExchange.using("externalEditor", function () {
StackExchange.using("snippets", function () {
StackExchange.snippets.init();
});
});
}, "code-snippets");

StackExchange.ready(function() {
var channelOptions = {
tags: "".split(" "),
id: "1"
};
initTagRenderer("".split(" "), "".split(" "), channelOptions);

StackExchange.using("externalEditor", function() {
// Have to fire editor after snippets, if snippets enabled
if (StackExchange.settings.snippets.snippetsEnabled) {
StackExchange.using("snippets", function() {
createEditor();
});
}
else {
createEditor();
}
});

function createEditor() {
StackExchange.prepareEditor({
heartbeatType: 'answer',
convertImagesToLinks: true,
noModals: true,
showLowRepImageUploadWarning: true,
reputationToPostImages: 10,
bindNavPrevention: true,
postfix: "",
imageUploader: {
brandingHtml: "Powered by u003ca class="icon-imgur-white" href="https://imgur.com/"u003eu003c/au003e",
contentPolicyHtml: "User contributions licensed under u003ca href="https://creativecommons.org/licenses/by-sa/3.0/"u003ecc by-sa 3.0 with attribution requiredu003c/au003e u003ca href="https://stackoverflow.com/legal/content-policy"u003e(content policy)u003c/au003e",
allowUrls: true
},
onDemand: true,
discardSelector: ".discard-answer"
,immediatelyShowMarkdownHelp:true
});

}
});

draft saved

draft discarded

Sign up or log in

StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});

Post as a guest

Name

Required, but never shown

StackExchange.ready(
function () {
StackExchange.openid.initPostLogin('.new-post-login', 'https%3a%2f%2fstackoverflow.com%2fquestions%2f53225799%2fdevice-function-pointer-as-template-parameter%23new-answer', 'question_page');
}
);

Post as a guest

Name

Required, but never shown

1 Answer
1

active

oldest

votes

1 Answer
1

active

oldest

votes

up vote
1
down vote

Well, answer is quite easy :

__device__ attribute is misplaced (as indicated by the compiler). The funcstruct should look like:

template<typename T, retV (* __device__  funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

But I don't know why nvcc and nvrtc have different expectations on this.

answered Nov 9 at 13:32

Regis Portalez

2,80811732

add a comment |

up vote
1
down vote

Well, answer is quite easy :

__device__ attribute is misplaced (as indicated by the compiler). The funcstruct should look like:

template<typename T, retV (* __device__  funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

But I don't know why nvcc and nvrtc have different expectations on this.

answered Nov 9 at 13:32

Regis Portalez

2,80811732

add a comment |

up vote
1
down vote

Well, answer is quite easy :

__device__ attribute is misplaced (as indicated by the compiler). The funcstruct should look like:

template<typename T, retV (* __device__  funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

But I don't know why nvcc and nvrtc have different expectations on this.

answered Nov 9 at 13:32

Regis Portalez

2,80811732

Well, answer is quite easy :

__device__ attribute is misplaced (as indicated by the compiler). The funcstruct should look like:

template<typename T, retV (* __device__  funcptr)(T)>

struct func 

{

    __device__ inline retV invoke(T i) { funcptr(i); }

};

But I don't know why nvcc and nvrtc have different expectations on this.

answered Nov 9 at 13:32

Regis Portalez

2,80811732

answered Nov 9 at 13:32

Regis Portalez

2,80811732

answered Nov 9 at 13:32

Regis Portalez

2,80811732

answered Nov 9 at 13:32

Regis Portalez

2,80811732

add a comment |

draft saved

draft discarded

draft saved

draft discarded

Sign up or log in

StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});

Post as a guest

Name

Required, but never shown

Post as a guest

Name

Required, but never shown

Sign up or log in

StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});

Post as a guest

Name

Required, but never shown

Sign up or log in

StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});

Post as a guest

Name

Required, but never shown

Sign up or log in

StackExchange.ready(function () {
StackExchange.helpers.onClickDraftSave('#login-link');
});

Post as a guest

Name

Required, but never shown

Name

Required, but never shown

Name

Required, but never shown

This page is only for reference, If you need detailed information, please check here

搜尋此網誌

Agfdhyk