Hi,
In pre-Pascal architecture I used the following code to call class member functions via function pointers. It used to work fine and cuda memory checker did not find any problems.
Under Pascal (GTX Titan X or 1080) under CUDA 8.0 (compute_61,sm_61) I get memory errors during debug mode (Windows 10, Visual Studio 2013/15). I get an access violation on assignment of the function pointer (line 19), and another access violation on retrieval of the function pointer (line 32). The actual code runs fine and the function pointer call correctly executes. Is this a bug in memory checker under Pascal or am I doing something illegal? Minimal working example below.
Thanks and Happy New Year!
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
template<typename T> using FunctionHandler = double(T::*)(double);
class TestClass
{
public:
FunctionHandler<TestClass> fooHandler;
__host__ __device__ double AddOne(double x)
{
return x + 1;
}
__host__ __device__ void BuildPointerTable()
{
this->fooHandler = &TestClass::AddOne;
}
};
__global__ void InitializeFooPointer(TestClass* obj)
{
obj->BuildPointerTable();
}
__global__ void Run(TestClass* obj)
{
double x = 1;
FunctionHandler<TestClass> handler = obj->fooHandler;
double result = (obj->*handler)(x);
printf("x + 1: %f \n", result);
}
void main()
{
TestClass* test = new TestClass();
TestClass* dTest = nullptr;
cudaError_t err = cudaMalloc(&dTest, sizeof(TestClass));
if (err != cudaSuccess) { return; }
err = cudaMemcpy(dTest, test, sizeof(TestClass), cudaMemcpyHostToDevice);
if (err != cudaSuccess) { return; }
InitializeFooPointer << < 1, 1 >> > (dTest);
Run << < 1, 1 >> > (dTest);
delete test;
err = cudaFree(dTest);
}