aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--libcuda/cuda_api_object.h195
-rw-r--r--libcuda/cuda_runtime_api.cc203
-rw-r--r--libcuda/cuobjdump.h1
-rw-r--r--libcuda/gpgpu_context.h53
-rw-r--r--src/Makefile2
-rw-r--r--src/abstract_hardware_model.h40
-rw-r--r--src/cuda-sim/Makefile8
-rw-r--r--src/cuda-sim/cuda-math.h16
-rw-r--r--src/cuda-sim/instructions.cc10
-rw-r--r--src/cuda-sim/ptx.y23
-rw-r--r--src/cuda-sim/ptx_parser.cc26
-rw-r--r--src/cuda-sim/ptx_parser.h38
-rw-r--r--src/gpgpu-sim/Makefile2
-rw-r--r--src/intersim2/Makefile1
14 files changed, 298 insertions, 320 deletions
diff --git a/libcuda/cuda_api_object.h b/libcuda/cuda_api_object.h
index 2001f91..d931fd5 100644
--- a/libcuda/cuda_api_object.h
+++ b/libcuda/cuda_api_object.h
@@ -1,14 +1,205 @@
#ifndef __cuda_api_object_h__
#define __cuda_api_object_h__
-class cuobjdumpSection;
-class kernel_config;
+#include <list>
+#include <map>
+#include <set>
+#include <string>
+
+#include "builtin_types.h"
+
+#include "../src/gpgpu-sim/gpu-sim.h"
+#include "../src/cuda-sim/ptx_ir.h"
+#include "../src/abstract_hardware_model.h"
+#include "cuobjdump.h"
+
+typedef std::list<gpgpu_ptx_sim_arg> gpgpu_ptx_sim_arg_list_t;
+
+#ifndef OPENGL_SUPPORT
+typedef unsigned long GLuint;
+#endif
+
+struct glbmap_entry {
+ GLuint m_bufferObj;
+ void *m_devPtr;
+ size_t m_size;
+ struct glbmap_entry *m_next;
+};
+
+typedef struct glbmap_entry glbmap_entry_t;
+
+struct _cuda_device_id {
+ _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
+ struct _cuda_device_id *next() { return m_next; }
+ unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
+ int num_devices() const {
+ if( m_next == NULL ) return 1;
+ else return 1 + m_next->num_devices();
+ }
+ struct _cuda_device_id *get_device( unsigned n )
+ {
+ assert( n < (unsigned)num_devices() );
+ struct _cuda_device_id *p=this;
+ for(unsigned i=0; i<n; i++)
+ p = p->m_next;
+ return p;
+ }
+ const struct cudaDeviceProp *get_prop() const
+ {
+ return m_gpgpu->get_prop();
+ }
+ unsigned get_id() const { return m_id; }
+
+ gpgpu_sim *get_gpgpu() { return m_gpgpu; }
+private:
+ unsigned m_id;
+ class gpgpu_sim *m_gpgpu;
+ struct _cuda_device_id *m_next;
+};
+
+struct CUctx_st {
+ CUctx_st( _cuda_device_id *gpu )
+ {
+ m_gpu = gpu;
+ m_binary_info.cmem = 0;
+ m_binary_info.gmem = 0;
+ no_of_ptx=0;
+ }
+
+ _cuda_device_id *get_device() { return m_gpu; }
+
+ void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
+ {
+ m_code[fat_cubin_handle] = symtab;
+ m_last_fat_cubin_handle = fat_cubin_handle;
+ }
+
+ void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_info &info )
+ {
+ symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
+ assert( s != NULL );
+ function_info *f = s->get_pc();
+ assert( f != NULL );
+ f->set_kernel_info(info);
+ }
+
+ void add_ptxinfo( const struct gpgpu_ptx_sim_info &info )
+ {
+ m_binary_info = info;
+ }
+
+ void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
+ {
+ if( m_code.find(fat_cubin_handle) != m_code.end() ) {
+ symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
+ if(s != NULL) {
+ function_info *f = s->get_pc();
+ assert( f != NULL );
+ m_kernel_lookup[hostFun] = f;
+ }
+ else {
+ printf("Warning: cannot find deviceFun %s\n", deviceFun);
+ m_kernel_lookup[hostFun] = NULL;
+ }
+ // assert( s != NULL );
+ // function_info *f = s->get_pc();
+ // assert( f != NULL );
+ // m_kernel_lookup[hostFun] = f;
+ } else {
+ m_kernel_lookup[hostFun] = NULL;
+ }
+ }
+
+ void register_hostFun_function( const char*hostFun, function_info* f){
+ m_kernel_lookup[hostFun] = f;
+ }
+
+ function_info *get_kernel(const char *hostFun)
+ {
+ std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
+ assert( i != m_kernel_lookup.end() );
+ return i->second;
+ }
+
+ int no_of_ptx;
+
+private:
+ _cuda_device_id *m_gpu; // selected gpu
+ std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
+ unsigned m_last_fat_cubin_handle;
+ std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
+ struct gpgpu_ptx_sim_info m_binary_info;
+
+};
+
+class kernel_config {
+public:
+ kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
+ {
+ m_GridDim=GridDim;
+ m_BlockDim=BlockDim;
+ m_sharedMem=sharedMem;
+ m_stream = stream;
+ }
+ kernel_config()
+ {
+ m_GridDim=dim3(-1,-1,-1);
+ m_BlockDim=dim3(-1,-1,-1);
+ m_sharedMem=0;
+ m_stream =NULL;
+ }
+ void set_arg( const void *arg, size_t size, size_t offset )
+ {
+ m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
+ }
+ dim3 grid_dim() const { return m_GridDim; }
+ dim3 block_dim() const { return m_BlockDim; }
+ void set_grid_dim(dim3 *d) { m_GridDim = *d; }
+ void set_block_dim(dim3 *d) { m_BlockDim = *d; }
+ gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
+ struct CUstream_st *get_stream() { return m_stream; }
+
+private:
+ dim3 m_GridDim;
+ dim3 m_BlockDim;
+ size_t m_sharedMem;
+ struct CUstream_st *m_stream;
+ gpgpu_ptx_sim_arg_list_t m_args;
+};
class cuda_runtime_api {
public:
+ cuda_runtime_api() {
+ g_glbmap = NULL;
+ }
// global list
+ std::list<cuobjdumpSection*> cuobjdumpSectionList;
std::list<cuobjdumpSection*> libSectionList;
std::list<kernel_config> g_cuda_launch_stack;
+ std::map<int, bool>fatbin_registered;
+ std::map<int, std::string> fatbinmap;
+ std::map<std::string, symbol_table*> name_symtab;
+ std::map<unsigned long long, size_t> g_mallocPtr_Size;
+ //maps sm version number to set of filenames
+ std::map<unsigned, std::set<std::string> > version_filename;
+ std::map<void *,void **> pinned_memory; //support for pinned memories added
+ std::map<void *, size_t> pinned_memory_size;
+ glbmap_entry_t* g_glbmap;
// member function list
+ void cuobjdumpInit();
+ void extract_code_using_cuobjdump();
+ void extract_ptx_files_using_cuobjdump(CUctx_st *context);
+ void cuobjdumpParseBinary(unsigned int handle);
+ std::list<cuobjdumpSection*> pruneSectionList(CUctx_st *context);
+ std::list<cuobjdumpSection*> mergeMatchingSections(std::string identifier);
+ std::list<cuobjdumpSection*> mergeSections();
+ cuobjdumpELFSection* findELFSection(const std::string identifier);
+ cuobjdumpPTXSection* findPTXSection(const std::string identifier);
+ void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context);
+ kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
+ gpgpu_ptx_sim_arg_list_t args,
+ struct dim3 gridDim,
+ struct dim3 blockDim,
+ struct CUctx_st* context );
};
#endif /* __cuda_api_object_h__ */
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 500b5f2..2d9c4f7 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -141,8 +141,6 @@
#include "../src/gpgpusim_entrypoint.h"
#include "../src/stream_manager.h"
#include "../src/abstract_hardware_model.h"
-typedef void * yyscan_t;
-#include "cuobjdump.h"
#include <pthread.h>
#include <semaphore.h>
@@ -208,145 +206,6 @@ void register_ptx_function( const char *name, function_info *impl )
# endif
#endif
-struct _cuda_device_id {
- _cuda_device_id(gpgpu_sim* gpu) {m_id = 0; m_next = NULL; m_gpgpu=gpu;}
- struct _cuda_device_id *next() { return m_next; }
- unsigned num_shader() const { return m_gpgpu->get_config().num_shader(); }
- int num_devices() const {
- if( m_next == NULL ) return 1;
- else return 1 + m_next->num_devices();
- }
- struct _cuda_device_id *get_device( unsigned n )
- {
- assert( n < (unsigned)num_devices() );
- struct _cuda_device_id *p=this;
- for(unsigned i=0; i<n; i++)
- p = p->m_next;
- return p;
- }
- const struct cudaDeviceProp *get_prop() const
- {
- return m_gpgpu->get_prop();
- }
- unsigned get_id() const { return m_id; }
-
- gpgpu_sim *get_gpgpu() { return m_gpgpu; }
-private:
- unsigned m_id;
- class gpgpu_sim *m_gpgpu;
- struct _cuda_device_id *m_next;
-};
-
-struct CUctx_st {
- CUctx_st( _cuda_device_id *gpu )
- {
- m_gpu = gpu;
- m_binary_info.cmem = 0;
- m_binary_info.gmem = 0;
- no_of_ptx=0;
- }
-
- _cuda_device_id *get_device() { return m_gpu; }
-
- void add_binary( symbol_table *symtab, unsigned fat_cubin_handle )
- {
- m_code[fat_cubin_handle] = symtab;
- m_last_fat_cubin_handle = fat_cubin_handle;
- }
-
- void add_ptxinfo( const char *deviceFun, const struct gpgpu_ptx_sim_info &info )
- {
- symbol *s = m_code[m_last_fat_cubin_handle]->lookup(deviceFun);
- assert( s != NULL );
- function_info *f = s->get_pc();
- assert( f != NULL );
- f->set_kernel_info(info);
- }
-
- void add_ptxinfo( const struct gpgpu_ptx_sim_info &info )
- {
- m_binary_info = info;
- }
-
- void register_function( unsigned fat_cubin_handle, const char *hostFun, const char *deviceFun )
- {
- if( m_code.find(fat_cubin_handle) != m_code.end() ) {
- symbol *s = m_code[fat_cubin_handle]->lookup(deviceFun);
- if(s != NULL) {
- function_info *f = s->get_pc();
- assert( f != NULL );
- m_kernel_lookup[hostFun] = f;
- }
- else {
- printf("Warning: cannot find deviceFun %s\n", deviceFun);
- m_kernel_lookup[hostFun] = NULL;
- }
- // assert( s != NULL );
- // function_info *f = s->get_pc();
- // assert( f != NULL );
- // m_kernel_lookup[hostFun] = f;
- } else {
- m_kernel_lookup[hostFun] = NULL;
- }
- }
-
- void register_hostFun_function( const char*hostFun, function_info* f){
- m_kernel_lookup[hostFun] = f;
- }
-
- function_info *get_kernel(const char *hostFun)
- {
- std::map<const void*,function_info*>::iterator i=m_kernel_lookup.find(hostFun);
- assert( i != m_kernel_lookup.end() );
- return i->second;
- }
-
- int no_of_ptx;
-
-private:
- _cuda_device_id *m_gpu; // selected gpu
- std::map<unsigned,symbol_table*> m_code; // fat binary handle => global symbol table
- unsigned m_last_fat_cubin_handle;
- std::map<const void*,function_info*> m_kernel_lookup; // unique id (CUDA app function address) => kernel entry point
- struct gpgpu_ptx_sim_info m_binary_info;
-
-};
-
-class kernel_config {
-public:
- kernel_config( dim3 GridDim, dim3 BlockDim, size_t sharedMem, struct CUstream_st *stream )
- {
- m_GridDim=GridDim;
- m_BlockDim=BlockDim;
- m_sharedMem=sharedMem;
- m_stream = stream;
- }
- kernel_config()
- {
- m_GridDim=dim3(-1,-1,-1);
- m_BlockDim=dim3(-1,-1,-1);
- m_sharedMem=0;
- m_stream =NULL;
- }
- void set_arg( const void *arg, size_t size, size_t offset )
- {
- m_args.push_front( gpgpu_ptx_sim_arg(arg,size,offset) );
- }
- dim3 grid_dim() const { return m_GridDim; }
- dim3 block_dim() const { return m_BlockDim; }
- void set_grid_dim(dim3 *d) { m_GridDim = *d; }
- void set_block_dim(dim3 *d) { m_BlockDim = *d; }
- gpgpu_ptx_sim_arg_list_t get_args() { return m_args; }
- struct CUstream_st *get_stream() { return m_stream; }
-
-private:
- dim3 m_GridDim;
- dim3 m_BlockDim;
- size_t m_sharedMem;
- struct CUstream_st *m_stream;
- gpgpu_ptx_sim_arg_list_t m_args;
-};
-
struct _cuda_device_id *GPGPUSim_Init()
{
//static _cuda_device_id *the_device = NULL;
@@ -626,7 +485,7 @@ static int get_app_cuda_version() {
}
//! Keep track of the association between filename and cubin handle
-void gpgpu_context::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){
+void cuda_runtime_api::cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context){
fatbinmap[handle] = filename;
}
@@ -705,8 +564,8 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx =
* then for next calls, only returns the appropriate number
*/
assert(fat_cubin_handle >= 1);
- if (fat_cubin_handle==1) ctx->cuobjdumpInit();
- ctx->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context);
+ if (fat_cubin_handle==1) ctx->api->cuobjdumpInit();
+ ctx->api->cuobjdumpRegisterFatBinary(fat_cubin_handle, filename, context);
return (void**)fat_cubin_handle;
}
@@ -803,7 +662,7 @@ void cudaRegisterFunctionInternal(
printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
deviceFun, hostFun, fat_cubin_handle);
if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump())
- ctx->cuobjdumpParseBinary(fat_cubin_handle);
+ ctx->api->cuobjdumpParseBinary(fat_cubin_handle);
context->register_function( fat_cubin_handle, hostFun, deviceFun );
}
@@ -830,7 +689,7 @@ void cudaRegisterVarInternal(
printf("GPGPU-Sim PTX: __cudaRegisterVar: hostVar = %p; deviceAddress = %s; deviceName = %s\n", hostVar, deviceAddress, deviceName);
printf("GPGPU-Sim PTX: __cudaRegisterVar: Registering const memory space of %d bytes\n", size);
if(GPGPUSim_Context()->get_device()->get_gpgpu()->get_config().use_cuobjdump())
- ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
+ ctx->api->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
fflush(stdout);
if ( constant && !global && !ext ) {
gpgpu_ptx_sim_register_const_variable(hostVar,deviceName,size);
@@ -894,7 +753,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
struct CUstream_st *stream = config.get_stream();
printf("\nGPGPU-Sim PTX: cudaLaunch for 0x%p (mode=%s) on stream %u\n", hostFun,
g_ptx_sim_mode?"functional simulation":"performance simulation", stream?stream->get_uid():0 );
- kernel_info_t *grid = ctx->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
+ kernel_info_t *grid = ctx->api->gpgpu_cuda_ptx_sim_init_grid(hostFun,config.get_args(),config.grid_dim(),config.block_dim(),context);
//do dynamic PDOM analysis for performance simulation scenario
std::string kname = grid->name();
function_info *kernel_func_info = grid->entry();
@@ -965,7 +824,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_
*devPtr = context->get_device()->get_gpgpu()->gpu_malloc(size);
if(g_debug_execution >= 3){
printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *devPtr);
- ctx->g_mallocPtr_Size[(unsigned long long)*devPtr] = size;
+ ctx->api->g_mallocPtr_Size[(unsigned long long)*devPtr] = size;
}
if ( *devPtr ) {
return g_last_cudaError = cudaSuccess;
@@ -988,7 +847,7 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu
*ptr = malloc(size);
if ( *ptr ) {
//track pinned memory size allocated in the host so that same amount of memory is also allocated in GPU.
- ctx->pinned_memory_size[*ptr]=size;
+ ctx->api->pinned_memory_size[*ptr]=size;
return g_last_cudaError = cudaSuccess;
} else {
return g_last_cudaError = cudaErrorMemoryAllocation;
@@ -1014,16 +873,16 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
flags=0;
CUctx_st* context = GPGPUSim_Context();
gpgpu_t *gpu = context->get_device()->get_gpgpu();
- std::map<void *, size_t>::const_iterator i = ctx->pinned_memory_size.find(pHost);
- assert(i != ctx->pinned_memory_size.end());
+ std::map<void *, size_t>::const_iterator i = ctx->api->pinned_memory_size.find(pHost);
+ assert(i != ctx->api->pinned_memory_size.end());
size_t size = i->second;
*pDevice = gpu->gpu_malloc(size);
if(g_debug_execution >= 3){
printf("GPGPU-Sim PTX: cudaMallocing %zu bytes starting at 0x%llx..\n",size, (unsigned long long) *pDevice);
- ctx->g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
+ ctx->api->g_mallocPtr_Size[(unsigned long long)*pDevice] = size;
}
if ( *pDevice ) {
- ctx->pinned_memory[pHost]=pDevice;
+ ctx->api->pinned_memory[pHost]=pDevice;
//Copy contents in cpu to gpu
gpu->memcpy_to_gpu((size_t)*pDevice,pHost,size);
return g_last_cudaError = cudaSuccess;
@@ -1050,7 +909,7 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
GLint buffer_size=0;
CUctx_st* context = GPGPUSim_Context();
- glbmap_entry_t *p = ctx->g_glbmap;
+ glbmap_entry_t *p = ctx->api->g_glbmap;
while ( p && p->m_bufferObj != bufferObj )
p = p->m_next;
if ( p == NULL ) {
@@ -1061,8 +920,8 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
// create entry and insert to front of list
glbmap_entry_t *n = (glbmap_entry_t *) calloc(1,sizeof(glbmap_entry_t));
- n->m_next = ctx->g_glbmap;
- ctx->g_glbmap = n;
+ n->m_next = ctx->api->g_glbmap;
+ ctx->api->g_glbmap = n;
// initialize entry
n->m_bufferObj = bufferObj;
@@ -1129,7 +988,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
strcat(file,path);
symbol_table *symtab = gpgpu_ptx_sim_load_ptx_from_filename( file );
std::string fname(path);
- ctx->name_symtab[fname] = symtab;
+ ctx->api->name_symtab[fname] = symtab;
context->add_binary(symtab, 1);
load_static_globals(symtab,STATIC_ALLOC_LIMIT,0xFFFFFFFF,context->get_device()->get_gpgpu());
load_constants(symtab,STATIC_ALLOC_LIMIT,context->get_device()->get_gpgpu());
@@ -1154,7 +1013,7 @@ cudaError_t cudaHostAllocInternal(void **pHost, size_t bytes, unsigned int flag
*pHost = malloc(bytes);
//need to track the size allocated so that cudaHostGetDevicePointer() can function properly.
//TODO: vary this function behavior based on flags value (following nvidia documentation)
- ctx->pinned_memory_size[*pHost]=bytes;
+ ctx->api->pinned_memory_size[*pHost]=bytes;
if( *pHost )
return g_last_cudaError = cudaSuccess;
else
@@ -2431,7 +2290,7 @@ __host__ cudaError_t CUDARTAPI cudaGetExportTable(const void **ppExportTable, co
//#include "../../cuobjdump_to_ptxplus/cuobjdump_parser.h"
//extracts all ptx files from binary and dumps into prog_name.unique_no.sm_<>.ptx files
-void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){
+void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
extern bool g_cdp_enabled;
char command[1000];
char *pytorch_bin = getenv("PYTORCH_BIN");
@@ -2509,7 +2368,7 @@ void gpgpu_context::extract_ptx_files_using_cuobjdump(CUctx_st *context){
* It is also responsible for extracting the libraries linked to the binary if the option is
* enabled
* */
-void gpgpu_context::extract_code_using_cuobjdump(){
+void cuda_runtime_api::extract_code_using_cuobjdump(){
CUctx_st *context = GPGPUSim_Context();
unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
@@ -2633,7 +2492,7 @@ void gpgpu_context::extract_code_using_cuobjdump(){
fclose(cuobjdump_in);
std::getline(libsf, line);
}
- api->libSectionList = cuobjdumpSectionList;
+ libSectionList = cuobjdumpSectionList;
//Restore the original section list
cuobjdumpSectionList = tmpsl;
@@ -2679,7 +2538,7 @@ void printSectionList(std::list<cuobjdumpSection*> sl) {
}
//! Remove unecessary sm versions from the section list
-std::list<cuobjdumpSection*> gpgpu_context::pruneSectionList(CUctx_st *context) {
+std::list<cuobjdumpSection*> cuda_runtime_api::pruneSectionList(CUctx_st *context) {
unsigned forced_max_capability = context->get_device()->get_gpgpu()->get_config().get_forced_max_capability();
//For ptxplus, force the max capability to 19 if it's higher or unspecified(0)
@@ -2732,7 +2591,7 @@ std::list<cuobjdumpSection*> gpgpu_context::pruneSectionList(CUctx_st *context)
}
//! Merge all PTX sections that have a specific identifier into one file
-std::list<cuobjdumpSection*> gpgpu_context::mergeMatchingSections(std::string identifier){
+std::list<cuobjdumpSection*> cuda_runtime_api::mergeMatchingSections(std::string identifier){
const char *ptxcode = "";
std::list<cuobjdumpSection*>::iterator old_iter;
cuobjdumpPTXSection* old_ptxsection = NULL;
@@ -2775,7 +2634,7 @@ std::list<cuobjdumpSection*> gpgpu_context::mergeMatchingSections(std::string id
}
//! Merge any PTX sections with matching identifiers
-std::list<cuobjdumpSection*> gpgpu_context::mergeSections(){
+std::list<cuobjdumpSection*> cuda_runtime_api::mergeSections(){
std::vector<std::string> identifier;
cuobjdumpPTXSection* ptxsection;
@@ -2822,10 +2681,10 @@ cuobjdumpELFSection* findELFSectionInList(std::list<cuobjdumpSection*> sectionli
}
//! Find an ELF section in all the known lists
-cuobjdumpELFSection* gpgpu_context::findELFSection(const std::string identifier){
+cuobjdumpELFSection* cuda_runtime_api::findELFSection(const std::string identifier){
cuobjdumpELFSection* sec = findELFSectionInList(cuobjdumpSectionList, identifier);
if (sec!=NULL)return sec;
- sec = findELFSectionInList(api->libSectionList, identifier);
+ sec = findELFSectionInList(libSectionList, identifier);
if (sec!=NULL)return sec;
std::cout << "Could not find " << identifier << std::endl;
assert(0 && "Could not find the required ELF section");
@@ -2857,10 +2716,10 @@ cuobjdumpPTXSection* findPTXSectionInList(std::list<cuobjdumpSection*> &sectionl
}
//! Find an PTX section in all the known lists
-cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier){
+cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifier){
cuobjdumpPTXSection* sec = findPTXSectionInList(cuobjdumpSectionList, identifier);
if (sec!=NULL)return sec;
- sec = findPTXSectionInList(api->libSectionList, identifier);
+ sec = findPTXSectionInList(libSectionList, identifier);
if (sec!=NULL)return sec;
std::cout << "Could not find " << identifier << std::endl;
assert(0 && "Could not find the required PTX section");
@@ -2870,7 +2729,7 @@ cuobjdumpPTXSection* gpgpu_context::findPTXSection(const std::string identifier)
//! Extract the code using cuobjdump and remove unnecessary sections
-void gpgpu_context::cuobjdumpInit(){
+void cuda_runtime_api::cuobjdumpInit(){
CUctx_st *context = GPGPUSim_Context();
extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.*
const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
@@ -2882,7 +2741,7 @@ void gpgpu_context::cuobjdumpInit(){
//! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
-void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
+void cuda_runtime_api::cuobjdumpParseBinary(unsigned int handle){
CUctx_st *context = GPGPUSim_Context();
if(fatbin_registered[handle]) return;
@@ -3176,7 +3035,7 @@ cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
}
#ifdef OPENGL_SUPPORT
CUctx_st* ctx = GPGPUSim_Context();
- glbmap_entry_t *p = ctx->g_glbmap;
+ glbmap_entry_t *p = ctx->api->g_glbmap;
while ( p && p->m_bufferObj != bufferObj )
p = p->m_next;
if ( p == NULL )
@@ -3531,7 +3390,7 @@ static int load_constants( symbol_table *symtab, addr_t min_gaddr, gpgpu_t *gpu
return nc_bytes;
}
-kernel_info_t * gpgpu_context::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
+kernel_info_t * cuda_runtime_api::gpgpu_cuda_ptx_sim_init_grid( const char *hostFun,
gpgpu_ptx_sim_arg_list_t args,
struct dim3 gridDim,
struct dim3 blockDim,
diff --git a/libcuda/cuobjdump.h b/libcuda/cuobjdump.h
index 49af3e2..6ab6778 100644
--- a/libcuda/cuobjdump.h
+++ b/libcuda/cuobjdump.h
@@ -4,6 +4,7 @@
#include <list>
#include <iostream>
+typedef void * yyscan_t;
struct cuobjdump_parser {
yyscan_t scanner;
int elfserial;
diff --git a/libcuda/gpgpu_context.h b/libcuda/gpgpu_context.h
index 7569ea6..4622f00 100644
--- a/libcuda/gpgpu_context.h
+++ b/libcuda/gpgpu_context.h
@@ -1,68 +1,15 @@
#ifndef __gpgpu_context_h__
#define __gpgpu_context_h__
-#include <list>
-#include <map>
-#include <set>
-#include <string>
#include "cuda_api_object.h"
-class cuobjdumpSection;
-class cuobjdumpELFSection;
-class cuobjdumpPTXSection;
-class symbol_table;
-class gpgpu_ptx_sim_arg;
-class kernel_info_t;
-
-typedef std::list<gpgpu_ptx_sim_arg> gpgpu_ptx_sim_arg_list_t;
-
-#ifndef OPENGL_SUPPORT
-typedef unsigned long GLuint;
-#endif
-
-struct glbmap_entry {
- GLuint m_bufferObj;
- void *m_devPtr;
- size_t m_size;
- struct glbmap_entry *m_next;
-};
-
-typedef struct glbmap_entry glbmap_entry_t;
-
class gpgpu_context {
public:
gpgpu_context() {
api = new cuda_runtime_api();
- g_glbmap = NULL;
}
// global list
- std::list<cuobjdumpSection*> cuobjdumpSectionList;
- std::map<int, bool>fatbin_registered;
- std::map<int, std::string> fatbinmap;
- std::map<std::string, symbol_table*> name_symtab;
- std::map<unsigned long long, size_t> g_mallocPtr_Size;
- //maps sm version number to set of filenames
- std::map<unsigned, std::set<std::string> > version_filename;
- std::map<void *,void **> pinned_memory; //support for pinned memories added
- std::map<void *, size_t> pinned_memory_size;
- glbmap_entry_t* g_glbmap;
// objects pointers for each file
cuda_runtime_api* api;
// member function list
- void cuobjdumpInit();
- void cuobjdumpParseBinary(unsigned int handle);
- void extract_code_using_cuobjdump();
- std::list<cuobjdumpSection*> pruneSectionList(CUctx_st *context);
- std::list<cuobjdumpSection*> mergeMatchingSections(std::string identifier);
- std::list<cuobjdumpSection*> mergeSections();
- cuobjdumpELFSection* findELFSection(const std::string identifier);
- cuobjdumpPTXSection* findPTXSection(const std::string identifier);
- void extract_ptx_files_using_cuobjdump(CUctx_st *context);
- void cuobjdumpRegisterFatBinary(unsigned int handle, const char* filename, CUctx_st *context);
- kernel_info_t *gpgpu_cuda_ptx_sim_init_grid( const char *kernel_key,
- gpgpu_ptx_sim_arg_list_t args,
- struct dim3 gridDim,
- struct dim3 blockDim,
- struct CUctx_st* context );
-
};
#endif /* __gpgpu_context_h__ */
diff --git a/src/Makefile b/src/Makefile
index 6001669..3ad511e 100644
--- a/src/Makefile
+++ b/src/Makefile
@@ -51,6 +51,8 @@ else
CXXFLAGS +=
endif
+CXXFLAGS += -I$(CUDA_INSTALL_PATH)/include
+
OPTFLAGS += -g3 -fPIC
CPP = g++ $(SNOW)
diff --git a/src/abstract_hardware_model.h b/src/abstract_hardware_model.h
index 77d5f58..68cb693 100644
--- a/src/abstract_hardware_model.h
+++ b/src/abstract_hardware_model.h
@@ -173,9 +173,7 @@ enum _memory_op_t {
#include <algorithm>
#if !defined(__VECTOR_TYPES_H__)
-struct dim3 {
- unsigned int x, y, z;
-};
+#include "vector_types.h"
#endif
struct dim3comp {
bool operator() (const dim3 & a, const dim3 & b) const
@@ -454,19 +452,7 @@ protected:
#if !defined(__CUDA_RUNTIME_API_H__)
-enum cudaChannelFormatKind {
- cudaChannelFormatKindSigned,
- cudaChannelFormatKindUnsigned,
- cudaChannelFormatKindFloat
-};
-
-struct cudaChannelFormatDesc {
- int x;
- int y;
- int z;
- int w;
- enum cudaChannelFormatKind f;
-};
+#include "builtin_types.h"
struct cudaArray {
void *devPtr;
@@ -478,28 +464,6 @@ struct cudaArray {
unsigned dimensions;
};
-enum cudaTextureAddressMode {
- cudaAddressModeWrap,
- cudaAddressModeClamp
-};
-
-enum cudaTextureFilterMode {
- cudaFilterModePoint,
- cudaFilterModeLinear
-};
-
-enum cudaTextureReadMode {
- cudaReadModeElementType,
- cudaReadModeNormalizedFloat
-};
-
-struct textureReference {
- int normalized;
- enum cudaTextureFilterMode filterMode;
- enum cudaTextureAddressMode addressMode[3];
- struct cudaChannelFormatDesc channelDesc;
-};
-
#endif
// Struct that record other attributes in the textureReference declaration
diff --git a/src/cuda-sim/Makefile b/src/cuda-sim/Makefile
index 999dad7..85d1c8c 100644
--- a/src/cuda-sim/Makefile
+++ b/src/cuda-sim/Makefile
@@ -79,16 +79,16 @@ libgpgpu_ptx_sim.a: $(OBJS)
ar rcs $(OUTPUT_DIR)/libgpgpu_ptx_sim.a $(OUTPUT_DIR)/ptx.tab.o $(OUTPUT_DIR)/lex.ptx_.o $(OUTPUT_DIR)/ptxinfo.tab.o $(OUTPUT_DIR)/lex.ptxinfo_.o $(OBJS)
$(OUTPUT_DIR)/ptx.tab.o: $(OUTPUT_DIR)/ptx.tab.c
- $(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
+ $(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptx.tab.c -o $(OUTPUT_DIR)/ptx.tab.o
$(OUTPUT_DIR)/lex.ptx_.o: $(OUTPUT_DIR)/lex.ptx_.c
- $(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
+ $(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptx_.c -o $(OUTPUT_DIR)/lex.ptx_.o
$(OUTPUT_DIR)/ptxinfo.tab.o: $(OUTPUT_DIR)/ptxinfo.tab.c
- $(CPP) -c $(OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
+ $(CPP) -c $(CXX_OPT) -DYYDEBUG $(OUTPUT_DIR)/ptxinfo.tab.c -o $(OUTPUT_DIR)/ptxinfo.tab.o
$(OUTPUT_DIR)/lex.ptxinfo_.o: $(OUTPUT_DIR)/lex.ptxinfo_.c $(OUTPUT_DIR)/ptxinfo.tab.c
- $(CPP) -c $(OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
+ $(CPP) -c $(CXX_OPT) $(OUTPUT_DIR)/lex.ptxinfo_.c -o $(OUTPUT_DIR)/lex.ptxinfo_.o
$(OUTPUT_DIR)/ptx.tab.c: ptx.y
bison --name-prefix=ptx_ -v -d ptx.y --file-prefix=$(OUTPUT_DIR)/ptx
diff --git a/src/cuda-sim/cuda-math.h b/src/cuda-sim/cuda-math.h
index a5db337..9a5468c 100644
--- a/src/cuda-sim/cuda-math.h
+++ b/src/cuda-sim/cuda-math.h
@@ -277,10 +277,10 @@ int float2int(float a, enum cudaRoundMode mode)
{
int tmp;
switch (mode) {
- case cuda_math::cudaRoundZero: tmp = truncf(a); break;
- case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break;
- case cuda_math::cudaRoundMinInf: tmp = floorf(a); break;
- case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break;
+ case cudaRoundZero: tmp = truncf(a); break;
+ case cudaRoundNearest: tmp = nearbyintf(a); break;
+ case cudaRoundMinInf: tmp = floorf(a); break;
+ case cudaRoundPosInf: tmp = ceilf(a); break;
default: abort();
}
return tmp;
@@ -296,10 +296,10 @@ unsigned int float2uint(float a, enum cudaRoundMode mode)
{
unsigned int tmp;
switch (mode) {
- case cuda_math::cudaRoundZero: tmp = truncf(a); break;
- case cuda_math::cudaRoundNearest: tmp = nearbyintf(a); break;
- case cuda_math::cudaRoundMinInf: tmp = floorf(a); break;
- case cuda_math::cudaRoundPosInf: tmp = ceilf(a); break;
+ case cudaRoundZero: tmp = truncf(a); break;
+ case cudaRoundNearest: tmp = nearbyintf(a); break;
+ case cudaRoundMinInf: tmp = floorf(a); break;
+ case cudaRoundPosInf: tmp = ceilf(a); break;
default: abort();
}
return tmp;
diff --git a/src/cuda-sim/instructions.cc b/src/cuda-sim/instructions.cc
index 69a97b6..186bb1e 100644
--- a/src/cuda-sim/instructions.cc
+++ b/src/cuda-sim/instructions.cc
@@ -2302,12 +2302,12 @@ ptx_reg_t f2x( ptx_reg_t x, unsigned from_width, unsigned to_width, int to_sign,
half_float::half tmp_h;
//assert( from_width == 32);
- enum cuda_math::cudaRoundMode mode = cuda_math::cudaRoundZero;
+ enum cudaRoundMode mode = cudaRoundZero;
switch (rounding_mode) {
- case RZI_OPTION: mode = cuda_math::cudaRoundZero; break;
- case RNI_OPTION: mode = cuda_math::cudaRoundNearest; break;
- case RMI_OPTION: mode = cuda_math::cudaRoundMinInf; break;
- case RPI_OPTION: mode = cuda_math::cudaRoundPosInf; break;
+ case RZI_OPTION: mode = cudaRoundZero; break;
+ case RNI_OPTION: mode = cudaRoundNearest; break;
+ case RMI_OPTION: mode = cudaRoundMinInf; break;
+ case RPI_OPTION: mode = cudaRoundPosInf; break;
default: break;
}
diff --git a/src/cuda-sim/ptx.y b/src/cuda-sim/ptx.y
index 8447693..837fbe9 100644
--- a/src/cuda-sim/ptx.y
+++ b/src/cuda-sim/ptx.y
@@ -229,7 +229,6 @@ class ptx_recognizer;
#include <string.h>
#include <math.h>
void syntax_not_implemented(yyscan_t yyscanner, ptx_recognizer* recognizer);
- extern int g_func_decl;
int ptx_lex(YYSTYPE * yylval_param, yyscan_t yyscanner, ptx_recognizer* recognizer);
int ptx_error( yyscan_t yyscanner, ptx_recognizer* recognizer, const char *s );
%}
@@ -260,21 +259,21 @@ block_spec_list: block_spec
function_decl: function_decl_header LEFT_PAREN { recognizer->start_function($1); recognizer->func_header_info("(");} param_entry RIGHT_PAREN {recognizer->func_header_info(")");} function_ident_param { $$ = recognizer->reset_symtab(); }
| function_decl_header { recognizer->start_function($1); } function_ident_param { $$ = recognizer->reset_symtab(); }
- | function_decl_header { recognizer->start_function($1); recognizer->add_function_name(""); g_func_decl=0; $$ = recognizer->reset_symtab(); }
+ | function_decl_header { recognizer->start_function($1); recognizer->add_function_name(""); recognizer->g_func_decl=0; $$ = recognizer->reset_symtab(); }
;
-function_ident_param: IDENTIFIER { recognizer->add_function_name($1); } LEFT_PAREN {recognizer->func_header_info("(");} param_list RIGHT_PAREN { g_func_decl=0; recognizer->func_header_info(")"); }
- | IDENTIFIER { recognizer->add_function_name($1); g_func_decl=0; }
+function_ident_param: IDENTIFIER { recognizer->add_function_name($1); } LEFT_PAREN {recognizer->func_header_info("(");} param_list RIGHT_PAREN { recognizer->g_func_decl=0; recognizer->func_header_info(")"); }
+ | IDENTIFIER { recognizer->add_function_name($1); recognizer->g_func_decl=0; }
;
-function_decl_header: ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
- | VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
- | WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; g_func_decl=1; recognizer->func_header(".entry"); }
- | FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
- | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
- | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
- | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; g_func_decl=1; recognizer->func_header(".func"); }
- | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; g_func_decl=1; recognizer->func_header(".func"); }
+function_decl_header: ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+ | VISIBLE_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+ | WEAK_DIRECTIVE ENTRY_DIRECTIVE { $$ = 1; recognizer->g_func_decl=1; recognizer->func_header(".entry"); }
+ | FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+ | VISIBLE_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+ | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+ | EXTERN_DIRECTIVE FUNC_DIRECTIVE { $$ = 2; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
+ | WEAK_DIRECTIVE FUNC_DIRECTIVE { $$ = 0; recognizer->g_func_decl=1; recognizer->func_header(".func"); }
;
param_list: /*empty*/
diff --git a/src/cuda-sim/ptx_parser.cc b/src/cuda-sim/ptx_parser.cc
index a61dc16..897beaf 100644
--- a/src/cuda-sim/ptx_parser.cc
+++ b/src/cuda-sim/ptx_parser.cc
@@ -49,7 +49,6 @@ void set_ptx_warp_size(const struct core_config * warp_size)
static bool g_debug_ir_generation=false;
const char *g_filename;
-unsigned g_max_regs_per_thread = 0;
// the program intermediate representation...
static symbol_table *g_global_allfiles_symbol_table = NULL;
@@ -61,27 +60,6 @@ static symbol *g_last_symbol = NULL;
int g_error_detected = 0;
-// type specifier stuff:
-memory_space_t g_space_spec = undefined_space;
-memory_space_t g_ptr_spec = undefined_space;
-int g_scalar_type_spec = -1;
-int g_vector_spec = -1;
-int g_alignment_spec = -1;
-int g_extern_spec = 0;
-
-// variable declaration stuff:
-type_info *g_var_type = NULL;
-
-// instruction definition stuff:
-const symbol *g_pred;
-int g_neg_pred;
-int g_pred_mod;
-symbol *g_label;
-int g_opcode = -1;
-std::list<operand_info> g_operands;
-std::list<int> g_options;
-std::list<int> g_wmma_options;
-std::list<int> g_scalar_type;
#define PTX_PARSE_DPRINTF(...) \
if( g_debug_ir_generation ) { \
@@ -362,10 +340,6 @@ bool check_for_duplicates( const char *identifier )
extern std::set<std::string> g_globals;
extern std::set<std::string> g_constants;
-int g_func_decl = 0;
-int g_ident_add_uid = 0;
-unsigned g_const_alloc = 1;
-
// Returns padding that needs to be inserted ahead of address to make it aligned to min(size, maxalign)
/*
* @param address the address in bytes
diff --git a/src/cuda-sim/ptx_parser.h b/src/cuda-sim/ptx_parser.h
index fb1b7f2..c52fe20 100644
--- a/src/cuda-sim/ptx_parser.h
+++ b/src/cuda-sim/ptx_parser.h
@@ -29,6 +29,7 @@
#define ptx_parser_INCLUDED
#include "../abstract_hardware_model.h"
+#include "ptx_ir.h"
extern const char *g_filename;
extern int g_error_detected;
@@ -44,7 +45,20 @@ class ptx_recognizer {
scanner = NULL;
g_size = -1;
g_add_identifier_cached__identifier = NULL;
+ g_alignment_spec = -1;
+ g_var_type = NULL;
+ g_opcode = -1;
+ g_space_spec = undefined_space;
+ g_ptr_spec = undefined_space;
+ g_scalar_type_spec = -1;
+ g_vector_spec = -1;
+ g_extern_spec = 0;
+ g_func_decl = 0;
+ g_ident_add_uid = 0;
+ g_const_alloc = 1;
+ g_max_regs_per_thread = 0;
}
+ // global list
yyscan_t scanner;
#define LINEBUF_SIZE (4*1024)
char linebuf[LINEBUF_SIZE];
@@ -53,7 +67,31 @@ class ptx_recognizer {
char *g_add_identifier_cached__identifier;
int g_add_identifier_cached__array_dim;
int g_add_identifier_cached__array_ident;
+ int g_alignment_spec;
+ // variable declaration stuff:
+ type_info *g_var_type;
+ // instruction definition stuff:
+ const symbol *g_pred;
+ int g_neg_pred;
+ int g_pred_mod;
+ symbol *g_label;
+ int g_opcode;
+ std::list<operand_info> g_operands;
+ std::list<int> g_options;
+ std::list<int> g_wmma_options;
+ std::list<int> g_scalar_type;
+ // type specifier stuff:
+ memory_space_t g_space_spec;
+ memory_space_t g_ptr_spec;
+ int g_scalar_type_spec;
+ int g_vector_spec;
+ int g_extern_spec;
+ int g_func_decl;
+ int g_ident_add_uid;
+ unsigned g_const_alloc;
+ unsigned g_max_regs_per_thread;
+ // member function list
void init_directive_state();
void init_instruction_state();
void start_function( int entry_point );
diff --git a/src/gpgpu-sim/Makefile b/src/gpgpu-sim/Makefile
index f10a8a4..4994577 100644
--- a/src/gpgpu-sim/Makefile
+++ b/src/gpgpu-sim/Makefile
@@ -53,6 +53,8 @@ else
CXXFLAGS +=
endif
+CXXFLAGS += -I$(CUDA_INSTALL_PATH)/include
+
POWER_FLAGS=
ifneq ($(GPGPUSIM_POWER_MODEL),)
POWER_FLAGS = -I$(GPGPUSIM_POWER_MODEL) -DGPGPUSIM_POWER_MODEL
diff --git a/src/intersim2/Makefile b/src/intersim2/Makefile
index 7d10b3f..3eeeb70 100644
--- a/src/intersim2/Makefile
+++ b/src/intersim2/Makefile
@@ -48,6 +48,7 @@ CPPFLAGS += -O3
endif
CPPFLAGS += -g
CPPFLAGS += -fPIC
+CPPFLAGS += -I$(CUDA_INSTALL_PATH)/include
LFLAGS +=
generated by cgit v1.3 (git 2.53.0) at 2026-07-21 04:50:30 +0000