Hello,
I’m trying to build a list of device function pointers in order to retrieve them when requested.
Moreover, it seems that cudaMemcpyFromSymbol is actually not copying anything when the functions are implemented in another file. Here’s the code:
list.cu
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include <stdlib.h>
#include "foo.h"
#define gpuSafeCall(ans) { gpuAssert((ans), #ans, __FILE__, __LINE__); }
#define gpuSafeMalloc(ans){ gpuMallocSharedAssert((void*)(ans), #ans, __FILE__, __LINE__);}
#define gpuCheckErrors() { gpuLastErrorCheck(__FILE__, __LINE__); }
/*
Check allocation success
*/
void gpuMallocSharedAssert(void *memory, const char *ans, const char* file, int line) {
if (memory == NULL) {
printf("\"%s\": allocation failed in file %s:%d\n", ans, file, line);
abort();
}
}
/*
Check CUDA runtime call success
*/
void gpuAssert(cudaError_t code, const char *ans, const char *file, int line)
{
if (code != cudaSuccess){
printf("\"%s\": \"%s\" in file %s:%d\n", ans, cudaGetErrorString(code), file, line);
abort();
}
}
/*
Check for in-kernel errors
*/
void gpuLastErrorCheck(const char *file, int line)
{
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
printf("Error: \"%s\" in file %s:%d\n", cudaGetErrorString(err), file, line);
abort();
}
}
typedef struct s_node {
component symbol;
const char *name;
struct s_node *next;
} componentNode;
typedef componentNode *componentsList;
componentsList components = NULL;
/*
* Functions
*/
typedef int(*component)();
__device__ component foo1 = Function1;
__device__ component foo2 = Function2;
/*
* Adds a node to the list
*/
void addNodeToList(componentsList *ls, componentNode *add) {
componentNode* current = *ls;
if (*ls == NULL)
*ls = add;
else {
while (current->next != NULL)
current = current->next;
current->next = add;
}
}
/*
* Creates a new node and adds to the list
*/
void addToList(componentsList *ls, const char *name, component symbol) {
componentNode *newElement;
gpuSafeMalloc(newElement = (componentNode*)malloc(sizeof(componentNode)));
newElement->name = name;
newElement->symbol = symbol;
newElement->next = NULL;
addNodeToList(ls, newElement);
}
/*
* Initializes the builtin components
*/
void initializeBuiltinComponents() {
component buffer;
gpuSafeCall(cudaMemcpyFromSymbol(&buffer, foo1, sizeof(component)));
addToList(&components, "Function1", buffer);
gpuSafeCall(cudaMemcpyFromSymbol(&buffer, foo2, sizeof(component)));
addToList(&components, "Function2", buffer);
}
/*
* Get a component from name string
*/
component getComponent(const char* name) {
componentNode *node;
if (strcmp("None", name) == 0)
return NULL;
for (node = components; node; node = node->next)
if (strcmp(node->name, name) == 0)
return node->symbol;
printf("[FATAL ERROR] Cannot find component '%s'\n");
abort();
return NULL;
}
foo.cu
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
__device__ int Function1(){
return 1;
}
__device__ int Function2(){
return 2;
}
foo.h
#ifndef FOO_H
#define FOO_H
#ifdef __cplusplus
extern "C" {
#endif
__device__ int Function1();
__device__ int Function2();
#ifdef __cplusplus
}
#endif
#endif
Inspecting the “initializeBuiltinComponents” with the Visual Studio Profiler, it seems that the “buffer” variable is NULL even if “cudaMemcpyFromSymbol” returns “cudaSuccess”.
Should I define some linking parameters?
Thanks,
Francesco
Working on CUDA 9.1, Windows