summaryrefslogtreecommitdiff
path: root/libcuda
diff options
context:
space:
mode:
Diffstat (limited to 'libcuda')
-rw-r--r--libcuda/cuda_runtime_api.cc1177
1 files changed, 689 insertions, 488 deletions
diff --git a/libcuda/cuda_runtime_api.cc b/libcuda/cuda_runtime_api.cc
index 716e297..175cbc5 100644
--- a/libcuda/cuda_runtime_api.cc
+++ b/libcuda/cuda_runtime_api.cc
@@ -199,8 +199,7 @@ void register_ptx_function( const char *name, function_info *impl )
struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
{
- //static _cuda_device_id *the_device = NULL;
- _cuda_device_id *the_device = GPGPUsim_ctx_ptr()->the_cude_device;
+ _cuda_device_id *the_device = the_gpgpusim->the_cude_device;
if( !the_device ) {
gpgpu_sim *the_gpu = gpgpu_ptx_sim_init_perf();
@@ -252,15 +251,14 @@ struct _cuda_device_id *gpgpu_context::GPGPUSim_Init()
return the_device;
}
-static CUctx_st* GPGPUSim_Context()
+CUctx_st* GPGPUSim_Context(gpgpu_context * ctx)
{
//static CUctx_st *the_context = NULL;
- gpgpu_context *cur_ctx = GPGPU_Context();
- CUctx_st *the_context = GPGPUsim_ctx_ptr()->the_context;
+ CUctx_st *the_context = ctx->the_gpgpusim->the_context;
if( the_context == NULL ) {
- _cuda_device_id *the_gpu = cur_ctx->GPGPUSim_Init();
- GPGPUsim_ctx_ptr()->the_context = new CUctx_st(the_gpu);
- the_context = GPGPUsim_ctx_ptr()->the_context;
+ _cuda_device_id *the_gpu = ctx->GPGPUSim_Init();
+ ctx->the_gpgpusim->the_context = new CUctx_st(the_gpu);
+ the_context = ctx->the_gpgpusim->the_context;
}
return the_context;
}
@@ -276,9 +274,9 @@ gpgpu_context* GPGPU_Context()
void ptxinfo_data::ptxinfo_addinfo()
{
+ CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
if(!get_ptxinfo_kname()){
/* This info is not per kernel (since CUDA 5.0 some info (e.g. gmem, and cmem) is added at the beginning for the whole binary ) */
- CUctx_st *context = GPGPUSim_Context();
print_ptxinfo();
context->add_ptxinfo(get_ptxinfo());
clear_ptxinfo();
@@ -289,7 +287,6 @@ gpgpu_context* GPGPU_Context()
clear_ptxinfo();
return;
}
- CUctx_st *context = GPGPUSim_Context();
print_ptxinfo();
context->add_ptxinfo( get_ptxinfo_kname(), get_ptxinfo() );
clear_ptxinfo();
@@ -585,7 +582,7 @@ void** cudaRegisterFatBinaryInternal( void *fatCubin, gpgpu_context* gpgpu_ctx =
printf("GPGPU-Sim PTX: ERROR ** this version of GPGPU-Sim requires CUDA 2.1 or higher\n");
exit(1);
#endif
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(ctx);
static unsigned next_fat_bin_handle = 1;
if(context->get_device()->get_gpgpu()->get_config().use_cuobjdump()) {
// The following workaround has only been verified on 64-bit systems.
@@ -732,7 +729,7 @@ void cudaRegisterFunctionInternal(
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(ctx);
unsigned fat_cubin_handle = (unsigned)(unsigned long long)fatCubinHandle;
printf("GPGPU-Sim PTX: __cudaRegisterFunction %s : hostFun 0x%p, fat_cubin_handle = %u\n",
deviceFun, hostFun, fat_cubin_handle);
@@ -763,7 +760,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())
+ if(GPGPUSim_Context(ctx)->get_device()->get_gpgpu()->get_config().use_cuobjdump())
ctx->cuobjdumpParseBinary((unsigned)(unsigned long long)fatCubinHandle);
fflush(stdout);
if ( constant && !global && !ext ) {
@@ -872,7 +869,7 @@ cudaError_t cudaLaunchInternal( const char *hostFun, gpgpu_context* gpgpu_ctx =
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- CUctx_st* context = GPGPUSim_Context();
+ CUctx_st* context = GPGPUSim_Context(ctx);
char *mode = getenv("PTX_SIM_MODE_FUNC");
if( mode )
sscanf(mode,"%u", &(ctx->func_sim->g_ptx_sim_mode));
@@ -960,7 +957,7 @@ cudaError_t cudaMallocInternal(void **devPtr, size_t size, gpgpu_context* gpgpu_
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- CUctx_st* context = GPGPUSim_Context();
+ CUctx_st* context = GPGPUSim_Context(ctx);
*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);
@@ -994,6 +991,29 @@ cudaError_t cudaMallocHostInternal(void **ptr, size_t size, gpgpu_context* gpgpu
}
}
+__host__ cudaError_t CUDARTAPI cudaMallocPitchInternal(void **devPtr, size_t *pitch, size_t width, size_t height, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ unsigned malloc_width_inbytes = width;
+ printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
+ CUctx_st* context = GPGPUSim_Context(ctx);
+ *devPtr = context->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height);
+ pitch[0] = malloc_width_inbytes;
+ if ( *devPtr ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
+}
+
cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
{
gpgpu_context *ctx;
@@ -1011,7 +1031,7 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
//only cpu memory allocation happens in cudaHostAlloc. Linking with device pointer to pinned memory happens here.
//TODO: once kernel is executed, the contents in global pointer of GPU must be copied back to CPU host pointer!
flags=0;
- CUctx_st* context = GPGPUSim_Context();
+ CUctx_st* context = GPGPUSim_Context(ctx);
gpgpu_t *gpu = context->get_device()->get_gpgpu();
std::map<void *, size_t>::const_iterator i = ctx->api->pinned_memory_size.find(pHost);
assert(i != ctx->api->pinned_memory_size.end());
@@ -1031,7 +1051,164 @@ cudaError_t cudaHostGetDevicePointerInternal(void **pDevice, void *pHost, unsign
}
}
-cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+__host__ cudaError_t CUDARTAPI cudaMallocArrayInternal(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1), gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
+ CUctx_st* context = GPGPUSim_Context(ctx);
+ (*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
+ (*array)->desc = *desc;
+ (*array)->width = width;
+ (*array)->height = height;
+ (*array)->size = size;
+ (*array)->dimensions = 2;
+ ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
+ printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
+ ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
+ if ( ((*array)->devPtr) ) {
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorMemoryAllocation;
+ }
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpyToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ size_t size = count;
+ printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
+ if( kind == cudaMemcpyHostToDevice )
+ gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+ return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DInternal(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ size_t size = spitch*height;
+ gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
+ if( kind == cudaMemcpyHostToDevice )
+ gpu->memcpy_to_gpu( (size_t)dst, src, size );
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst, (size_t)src, size );
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArrayInternal(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ size_t size = spitch*height;
+ size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z;
+ gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
+ unsigned elem_size = channel_size/8;
+ gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
+ gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
+ gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
+ gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
+ gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
+ gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
+ if( kind == cudaMemcpyHostToDevice )
+ gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
+ else if( kind == cudaMemcpyDeviceToHost )
+ gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
+ else if( kind == cudaMemcpyDeviceToDevice )
+ gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
+ else {
+ printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
+ abort();
+ }
+ dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
+ return g_last_cudaError = cudaSuccess;
+}
+
+#if (CUDART_VERSION >= 8000)
+cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ function_info *entry = context->get_kernel(hostFunc);
+ printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
+ if (flags == cudaOccupancyDefault) {
+ //create kernel_info based on entry
+ dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
+ * context->get_device()->get_gpgpu()->get_config().num_shader());
+ dim3 blockDim(blockSize);
+ kernel_info_t result(gridDim, blockDim, entry);
+ //if(entry == NULL){
+ // *numBlocks = 1;
+ // return g_last_cudaError = cudaErrorUnknown;
+ //}
+ *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
+ printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ cuda_not_implemented(__my_func__,__LINE__);
+ return g_last_cudaError = cudaErrorUnknown;
+ }
+}
+
+#endif
+
+__host__ cudaError_t CUDARTAPI cudaMemsetInternal(void *mem, int c, size_t count, gpgpu_context* gpgpu_ctx = NULL)
{
gpgpu_context *ctx;
if (gpgpu_ctx){
@@ -1042,12 +1219,45 @@ cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ gpu->gpu_memset((size_t)mem, c, count);
+ return g_last_cudaError = cudaSuccess;
+}
+
+//memset operation is done but i think its not async?
+__host__ cudaError_t CUDARTAPI cudaMemsetAsyncInternal(void *mem, int c, size_t count, cudaStream_t stream=0, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ gpu->gpu_memset((size_t)mem, c, count);
+ return g_last_cudaError = cudaSuccess;
+}
+
+cudaError_t cudaGLMapBufferObjectInternal(void** devPtr, GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+{
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
#ifdef OPENGL_SUPPORT
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
GLint buffer_size=0;
- CUctx_st* context = GPGPUSim_Context();
+ CUctx_st* context = GPGPUSim_Context(ctx);
glbmap_entry_t *p = ctx->api->g_glbmap;
while ( p && p->m_bufferObj != bufferObj )
@@ -1118,7 +1328,7 @@ cuLinkAddFileInternal(CUlinkState state, CUjitInputType type, const char *path,
//blocking
assert(type==CU_JIT_INPUT_PTX);
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(ctx);
char *file = getenv("PTX_JIT_PATH");
if(file==NULL){
printf("GPGPU-Sim PTX: ERROR: PTX_JIT_PATH has not been set\n");
@@ -1190,7 +1400,7 @@ cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *a
if(g_debug_execution >= 3){
announce_call(__my_func__);
}
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(ctx);
function_info *entry = context->get_kernel(hostFun);
if( entry ) {
const struct gpgpu_ptx_sim_info *kinfo = entry->get_kernel_info();
@@ -1210,6 +1420,443 @@ cudaError_t CUDARTAPI cudaFuncGetAttributesInternal(struct cudaFuncAttributes *a
return g_last_cudaError = cudaSuccess;
}
+#if (CUDART_VERSION > 5000)
+__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+
+ const struct cudaDeviceProp *prop;
+ _cuda_device_id *dev = ctx->GPGPUSim_Init();
+
+ if (device <= dev->num_devices() ) {
+ prop = dev->get_prop();
+ switch (attr) {
+ case 1:
+ *value= prop->maxThreadsPerBlock;
+ break;
+ case 2:
+ *value= prop->maxThreadsDim[0];
+ break;
+ case 3:
+ *value= prop->maxThreadsDim[1];
+ break;
+ case 4:
+ *value= prop->maxThreadsDim[2];
+ break;
+ case 5:
+ *value= prop->maxGridSize[0];
+ break;
+ case 6:
+ *value= prop->maxGridSize[1];
+ break;
+ case 7:
+ *value= prop->maxGridSize[2];
+ break;
+ case 8:
+ *value= prop->sharedMemPerBlock;
+ break;
+ case 9:
+ *value= prop->totalConstMem;
+ break;
+ case 10:
+ *value= prop->warpSize;
+ break;
+ case 11:
+ *value= 16;//dummy value
+ break;
+ case 12:
+ *value= prop->regsPerBlock;
+ break;
+ case 13:
+ *value= 1480000;//for 1080ti
+ break;
+ case 14:
+ *value= prop->textureAlignment ;
+ break;
+ case 15:
+ *value = 0;
+ break;
+ case 16:
+ *value= prop->multiProcessorCount ;
+ break;
+ case 17:
+ case 18:
+ case 19:
+ *value = 0;
+ break;
+ case 21:
+ case 22:
+ case 23:
+ case 24:
+ case 25:
+ case 26:
+ case 27:
+ case 28:
+ case 42:
+ case 45:
+ case 46:
+ case 47:
+ case 48:
+ case 49:
+ case 52:
+ case 53:
+ case 55:
+ case 56:
+ case 57:
+ case 58:
+ case 59:
+ case 60:
+ case 61:
+ case 62:
+ case 63:
+ case 64:
+ case 66:
+ case 67:
+ case 69:
+ case 70:
+ case 71:
+ case 73:
+ case 74:
+ case 77:
+ *value = 1000;//dummy value
+ break;
+ case 29:
+ case 43:
+ case 54:
+ case 65:
+ case 68:
+ case 72:
+ *value = 10;//dummy value
+ break;
+ case 30:
+ case 51:
+ *value = 128;//dummy value
+ break;
+ case 31:
+ *value = 1;
+ break;
+ case 32:
+ *value = 0;
+ break;
+ case 33:
+ case 50:
+ *value = 0;//dummy value
+ break;
+ case 34:
+ *value= 0;
+ break;
+ case 35:
+ *value = 0;
+ break;
+ case 36:
+ *value = 1250000;//CK value for 1080ti
+ break;
+ case 37:
+ *value = 352;//value for 1080ti
+ break;
+ case 38:
+ *value = 3000000;//value for 1080ti
+ break;
+ case 39:
+ *value= dev->get_gpgpu()->threads_per_core();
+ break;
+ case 40:
+ *value= 0;
+ break;
+ case 41:
+ *value= 0;
+ break;
+ case 75://cudaDevAttrComputeCapabilityMajor
+ *value= prop->major ;
+ break;
+ case 76://cudaDevAttrComputeCapabilityMinor
+ *value= prop->minor ;
+ break;
+ case 78:
+ *value= 0 ; //TODO: as of now, we dont support stream priorities.
+ break;
+ case 79:
+ *value= 0;
+ break;
+ case 80:
+ *value= 0;
+ break;
+ #if (CUDART_VERSION > 5050)
+ case 81:
+ *value= prop->sharedMemPerMultiprocessor;
+ break;
+ case 82:
+ *value= prop->regsPerMultiprocessor;
+ break;
+ #endif
+ case 83:
+ case 84:
+ case 85:
+ case 86:
+ *value= 0;
+ break;
+ case 87:
+ *value= 4;//dummy value
+ break;
+ case 88:
+ case 89:
+ *value= 0;
+ break;
+ default:
+ printf("ERROR: Attribute number %d unimplemented \n",attr);
+ abort();
+ }
+ return g_last_cudaError = cudaSuccess;
+ } else {
+ return g_last_cudaError = cudaErrorInvalidDevice;
+ }
+}
+#endif
+
+__host__ cudaError_t CUDARTAPI cudaBindTextureInternal(size_t *offset,
+ const struct textureReference *texref,
+ const void *devPtr,
+ const struct cudaChannelFormatDesc *desc,
+ size_t size __dv(UINT_MAX),
+ gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
+ struct cudaArray *array;
+ array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
+ array->desc = *desc;
+ array->size = size;
+ array->width = size;
+ array->height = 1;
+ array->dimensions = 1;
+ array->devPtr = (void*)devPtr;
+ array->devPtr32 = (int)(long long)devPtr;
+ offset = 0;
+ printf("GPGPU-Sim PTX: size = %zu\n", size);
+ printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array);
+ printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
+ printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+ printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
+ printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
+ gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
+ devPtr = (void*)(long long)array->devPtr32;
+ printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
+ return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaBindTextureToArrayInternal(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array);
+ printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
+ printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+ printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
+ gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
+ return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaUnbindTextureInternal(const struct textureReference *texref, gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
+ printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
+
+ gpu->gpgpu_ptx_sim_unbindTexture(texref);
+ return g_last_cudaError = cudaSuccess;
+}
+
+__host__ cudaError_t CUDARTAPI cudaLaunchKernelInternal( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream, gpgpu_context* gpgpu_ctx = NULL )
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ function_info *entry = context->get_kernel(hostFun);
+#if CUDART_VERSION < 10000
+ cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream, ctx);
+#endif
+ for(unsigned i = 0; i < entry->num_args(); i++){
+ std::pair<size_t, unsigned> p = entry->get_param_config(i);
+ cudaSetupArgumentInternal(args[i], p.first, p.second);
+ }
+
+ cudaLaunchInternal(hostFun);
+ return g_last_cudaError = cudaSuccess;
+}
+
+void __cudaRegisterTextureInternal(
+ void **fatCubinHandle,
+ const struct textureReference *hostVar,
+ const void **deviceAddress,
+ const char *deviceName,
+ int dim,
+ int norm,
+ int ext,
+ gpgpu_context* gpgpu_ctx = NULL
+) //passes in a newly created textureReference
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ std::string devStr (deviceName);
+ #if (CUDART_VERSION > 4020)
+ if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':')
+ devStr = devStr.replace(0, 2, "");
+ #endif
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ gpgpu_t *gpu = context->get_device()->get_gpgpu();
+ printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n");
+ gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext);
+ printf("GPGPU-Sim PTX: int dim = %d\n", dim);
+ printf("GPGPU-Sim PTX: int norm = %d\n", norm);
+ printf("GPGPU-Sim PTX: int ext = %d\n", ext);
+ printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ );
+}
+
+cudaError_t cudaGLUnmapBufferObjectInternal(GLuint bufferObj, gpgpu_context* gpgpu_ctx = NULL)
+{
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+#ifdef OPENGL_SUPPORT
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ CUctx_st* ctx = GPGPUSim_Context(ctx);
+ glbmap_entry_t *p = ctx->api->g_glbmap;
+ while ( p && p->m_bufferObj != bufferObj )
+ p = p->m_next;
+ if ( p == NULL )
+ return g_last_cudaError = cudaErrorUnknown;
+
+ char *data = (char *) calloc(p->m_size,1);
+ memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size );
+ glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data);
+ free(data);
+
+ return g_last_cudaError = cudaSuccess;
+#else
+ fflush(stdout);
+ fflush(stderr);
+ printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n");
+ fflush(stdout);
+ exit(50);
+#endif
+}
+
+#if CUDART_VERSION >= 3000
+
+__host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfigInternal(const char *func, enum cudaFuncCache cacheConfig, gpgpu_context* gpgpu_ctx = NULL )
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig);
+ return g_last_cudaError = cudaSuccess;
+}
+
+#endif
+
+#if CUDART_VERSION >= 4000
+CUresult CUDAAPI cuLaunchKernelInternal(CUfunction f,
+ unsigned int gridDimX,
+ unsigned int gridDimY,
+ unsigned int gridDimZ,
+ unsigned int blockDimX,
+ unsigned int blockDimY,
+ unsigned int blockDimZ,
+ unsigned int sharedMemBytes,
+ CUstream hStream,
+ void **kernelParams,
+ void **extra,
+ gpgpu_context* gpgpu_ctx = NULL)
+{
+ gpgpu_context *ctx;
+ if (gpgpu_ctx){
+ ctx = gpgpu_ctx;
+ } else {
+ ctx = GPGPU_Context();
+ }
+ if(g_debug_execution >= 3){
+ announce_call(__my_func__);
+ }
+ if (extra!=NULL){
+ printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n");
+ abort();
+ }
+ const char *hostFun = (const char*) f;
+ CUctx_st *context = GPGPUSim_Context(ctx);
+ function_info *entry = context->get_kernel(hostFun);
+ cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream, ctx);
+ for(unsigned i = 0; i < entry->num_args(); i++){
+ std::pair<size_t, unsigned> p = entry->get_param_config(i);
+ cudaSetupArgumentInternal(kernelParams[i], p.first, p.second, ctx);
+ }
+ cudaLaunchInternal(hostFun, ctx);
+ return CUDA_SUCCESS;
+}
+#endif /* CUDART_VERSION >= 4000 */
/*******************************************************************************
* *
@@ -1240,42 +1887,12 @@ __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size)
}
__host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t width, size_t height)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- unsigned malloc_width_inbytes = width;
- printf("GPGPU-Sim PTX: cudaMallocPitch (width = %d)\n", malloc_width_inbytes);
- CUctx_st* ctx = GPGPUSim_Context();
- *devPtr = ctx->get_device()->get_gpgpu()->gpu_malloc(malloc_width_inbytes*height);
- pitch[0] = malloc_width_inbytes;
- if ( *devPtr ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ return cudaMallocPitchInternal(devPtr, pitch, width, height);
}
__host__ cudaError_t CUDARTAPI cudaMallocArray(struct cudaArray **array, const struct cudaChannelFormatDesc *desc, size_t width, size_t height __dv(1))
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- unsigned size = width * height * ((desc->x + desc->y + desc->z + desc->w)/8);
- CUctx_st* context = GPGPUSim_Context();
- (*array) = (struct cudaArray*) malloc(sizeof(struct cudaArray));
- (*array)->desc = *desc;
- (*array)->width = width;
- (*array)->height = height;
- (*array)->size = size;
- (*array)->dimensions = 2;
- ((*array)->devPtr32)= (int) (long long)context->get_device()->get_gpgpu()->gpu_mallocarray(size);
- printf("GPGPU-Sim PTX: cudaMallocArray: devPtr32 = %d\n", ((*array)->devPtr32));
- ((*array)->devPtr) = (void*) (long long) ((*array)->devPtr32);
- if ( ((*array)->devPtr) ) {
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorMemoryAllocation;
- }
+ return cudaMallocArrayInternal(array, desc, width, height __dv(1));
}
__host__ cudaError_t CUDARTAPI cudaFree(void *devPtr)
@@ -1351,25 +1968,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy(void *dst, const void *src, size_t cou
__host__ cudaError_t CUDARTAPI cudaMemcpyToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t count, enum cudaMemcpyKind kind)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- size_t size = count;
- printf("GPGPU-Sim PTX: cudaMemcpyToArray\n");
- if( kind == cudaMemcpyHostToDevice )
- gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)(dst->devPtr), (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpyToArray - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
- return g_last_cudaError = cudaSuccess;
+ return cudaMemcpyToArrayInternal(dst, wOffset, hOffset, src, count, kind);
}
@@ -1395,59 +1994,14 @@ __host__ cudaError_t CUDARTAPI cudaMemcpyArrayToArray(struct cudaArray *dst, siz
__host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- size_t size = spitch*height;
- gpgpusim_ptx_assert( (dpitch==spitch), "different src and dst pitch not supported yet" );
- if( kind == cudaMemcpyHostToDevice )
- gpu->memcpy_to_gpu( (size_t)dst, src, size );
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst, (size_t)src, size );
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)dst, (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- return g_last_cudaError = cudaSuccess;
+ return cudaMemcpy2DInternal(dst, dpitch, src, spitch, width, height, kind);
}
-
__host__ cudaError_t CUDARTAPI cudaMemcpy2DToArray(struct cudaArray *dst, size_t wOffset, size_t hOffset, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- size_t size = spitch*height;
- size_t channel_size = dst->desc.w+dst->desc.x+dst->desc.y+dst->desc.z;
- gpgpusim_ptx_assert( ((channel_size%8) == 0), "none byte multiple destination channel size not supported (sz=%u)", channel_size );
- unsigned elem_size = channel_size/8;
- gpgpusim_ptx_assert( (dst->dimensions==2), "copy to none 2D array not supported" );
- gpgpusim_ptx_assert( (wOffset==0), "non-zero wOffset not yet supported" );
- gpgpusim_ptx_assert( (hOffset==0), "non-zero hOffset not yet supported" );
- gpgpusim_ptx_assert( (dst->height == (int)height), "partial copy not supported" );
- gpgpusim_ptx_assert( (elem_size*dst->width == width), "partial copy not supported" );
- gpgpusim_ptx_assert( (spitch == width), "spitch != width not supported" );
- if( kind == cudaMemcpyHostToDevice )
- gpu->memcpy_to_gpu( (size_t)(dst->devPtr), src, size);
- else if( kind == cudaMemcpyDeviceToHost )
- gpu->memcpy_from_gpu( dst->devPtr, (size_t)src, size);
- else if( kind == cudaMemcpyDeviceToDevice )
- gpu->memcpy_gpu_to_gpu( (size_t)dst->devPtr, (size_t)src, size);
- else {
- printf("GPGPU-Sim PTX: cudaMemcpy2D - ERROR : unsupported cudaMemcpyKind\n");
- abort();
- }
- dst->devPtr32 = (unsigned) (size_t)(dst->devPtr);
- return g_last_cudaError = cudaSuccess;
+ return cudaMemcpy2DToArrayInternal(dst, wOffset, hOffset, src, spitch, width, height, kind);
}
-
__host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, const struct cudaArray *src, size_t wOffset, size_t hOffset, size_t width, size_t height, enum cudaMemcpyKind kind)
{
if(g_debug_execution >= 3){
@@ -1457,7 +2011,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArray(void *dst, size_t dpitch, c
return g_last_cudaError = cudaErrorUnknown;
}
-
__host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, size_t wOffsetDst, size_t hOffsetDst, const struct cudaArray *src, size_t wOffsetSrc, size_t hOffsetSrc, size_t width, size_t height, enum cudaMemcpyKind kind __dv(cudaMemcpyDeviceToDevice))
{
if(g_debug_execution >= 3){
@@ -1467,7 +2020,6 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DArrayToArray(struct cudaArray *dst, s
return g_last_cudaError = cudaErrorUnknown;
}
-
__host__ cudaError_t CUDARTAPI cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset __dv(0), enum cudaMemcpyKind kind __dv(cudaMemcpyHostToDevice))
{
if(g_debug_execution >= 3){
@@ -1582,27 +2134,7 @@ __host__ cudaError_t CUDARTAPI cudaMemcpy2DFromArrayAsync(void *dst, size_t dpit
#if (CUDART_VERSION >= 8000)
cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks, const char *hostFunc, int blockSize, size_t dynamicSMemSize, unsigned int flags)
{
- printf("GPGPU-Sim PTX: cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags %p\n", hostFunc);
- CUctx_st *context = GPGPUSim_Context();
- function_info *entry = context->get_kernel(hostFunc);
- printf("Calculate Maxium Active Block with function ptr=%p, blockSize=%d, SMemSize=%d\n", hostFunc, blockSize, dynamicSMemSize);
- if (flags == cudaOccupancyDefault) {
- //create kernel_info based on entry
- dim3 gridDim(context->get_device()->get_gpgpu()->max_cta_per_core()
- * context->get_device()->get_gpgpu()->get_config().num_shader());
- dim3 blockDim(blockSize);
- kernel_info_t result(gridDim, blockDim, entry);
- //if(entry == NULL){
- // *numBlocks = 1;
- // return g_last_cudaError = cudaErrorUnknown;
- //}
- *numBlocks = context->get_device()->get_gpgpu()->get_max_cta(result);
- printf("Maximum size is %d with gridDim %d and blockDim %d\n", *numBlocks, gridDim.x, blockDim.x);
- return g_last_cudaError = cudaSuccess;
- } else {
- cuda_not_implemented(__my_func__,__LINE__);
- return g_last_cudaError = cudaErrorUnknown;
- }
+ return cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsInternal(numBlocks, hostFunc, blockSize, dynamicSMemSize, flags);
}
#endif
@@ -1616,227 +2148,13 @@ cudaError_t CUDARTAPI cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int
*******************************************************************************/
__host__ cudaError_t CUDARTAPI cudaMemset(void *mem, int c, size_t count)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- gpu->gpu_memset((size_t)mem, c, count);
- return g_last_cudaError = cudaSuccess;
+ return cudaMemsetInternal(mem, c, count);
}
-#if (CUDART_VERSION > 5000)
-__host__ cudaError_t CUDARTAPI cudaDeviceGetAttributeInternal(int *value, enum cudaDeviceAttr attr, int device, gpgpu_context* gpgpu_ctx = NULL)
-{
- gpgpu_context *ctx;
- if (gpgpu_ctx){
- ctx = gpgpu_ctx;
- } else {
- ctx = GPGPU_Context();
- }
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
-
- const struct cudaDeviceProp *prop;
- _cuda_device_id *dev = ctx->GPGPUSim_Init();
-
- if (device <= dev->num_devices() ) {
- prop = dev->get_prop();
- switch (attr) {
- case 1:
- *value= prop->maxThreadsPerBlock;
- break;
- case 2:
- *value= prop->maxThreadsDim[0];
- break;
- case 3:
- *value= prop->maxThreadsDim[1];
- break;
- case 4:
- *value= prop->maxThreadsDim[2];
- break;
- case 5:
- *value= prop->maxGridSize[0];
- break;
- case 6:
- *value= prop->maxGridSize[1];
- break;
- case 7:
- *value= prop->maxGridSize[2];
- break;
- case 8:
- *value= prop->sharedMemPerBlock;
- break;
- case 9:
- *value= prop->totalConstMem;
- break;
- case 10:
- *value= prop->warpSize;
- break;
- case 11:
- *value= 16;//dummy value
- break;
- case 12:
- *value= prop->regsPerBlock;
- break;
- case 13:
- *value= 1480000;//for 1080ti
- break;
- case 14:
- *value= prop->textureAlignment ;
- break;
- case 15:
- *value = 0;
- break;
- case 16:
- *value= prop->multiProcessorCount ;
- break;
- case 17:
- case 18:
- case 19:
- *value = 0;
- break;
- case 21:
- case 22:
- case 23:
- case 24:
- case 25:
- case 26:
- case 27:
- case 28:
- case 42:
- case 45:
- case 46:
- case 47:
- case 48:
- case 49:
- case 52:
- case 53:
- case 55:
- case 56:
- case 57:
- case 58:
- case 59:
- case 60:
- case 61:
- case 62:
- case 63:
- case 64:
- case 66:
- case 67:
- case 69:
- case 70:
- case 71:
- case 73:
- case 74:
- case 77:
- *value = 1000;//dummy value
- break;
- case 29:
- case 43:
- case 54:
- case 65:
- case 68:
- case 72:
- *value = 10;//dummy value
- break;
- case 30:
- case 51:
- *value = 128;//dummy value
- break;
- case 31:
- *value = 1;
- break;
- case 32:
- *value = 0;
- break;
- case 33:
- case 50:
- *value = 0;//dummy value
- break;
- case 34:
- *value= 0;
- break;
- case 35:
- *value = 0;
- break;
- case 36:
- *value = 1250000;//CK value for 1080ti
- break;
- case 37:
- *value = 352;//value for 1080ti
- break;
- case 38:
- *value = 3000000;//value for 1080ti
- break;
- case 39:
- *value= dev->get_gpgpu()->threads_per_core();
- break;
- case 40:
- *value= 0;
- break;
- case 41:
- *value= 0;
- break;
- case 75://cudaDevAttrComputeCapabilityMajor
- *value= prop->major ;
- break;
- case 76://cudaDevAttrComputeCapabilityMinor
- *value= prop->minor ;
- break;
- case 78:
- *value= 0 ; //TODO: as of now, we dont support stream priorities.
- break;
- case 79:
- *value= 0;
- break;
- case 80:
- *value= 0;
- break;
- #if (CUDART_VERSION > 5050)
- case 81:
- *value= prop->sharedMemPerMultiprocessor;
- break;
- case 82:
- *value= prop->regsPerMultiprocessor;
- break;
- #endif
- case 83:
- case 84:
- case 85:
- case 86:
- *value= 0;
- break;
- case 87:
- *value= 4;//dummy value
- break;
- case 88:
- case 89:
- *value= 0;
- break;
- default:
- printf("ERROR: Attribute number %d unimplemented \n",attr);
- abort();
- }
- return g_last_cudaError = cudaSuccess;
- } else {
- return g_last_cudaError = cudaErrorInvalidDevice;
- }
-}
-#endif
-
//memset operation is done but i think its not async?
__host__ cudaError_t CUDARTAPI cudaMemsetAsync(void *mem, int c, size_t count, cudaStream_t stream=0)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- printf("GPGPU-Sim PTX: WARNING: Asynchronous memset not supported (%s)\n", __my_func__);
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- gpu->gpu_memset((size_t)mem, c, count);
- return g_last_cudaError = cudaSuccess;
+ return cudaMemsetAsyncInternal(mem, c, count, stream=0);
}
__host__ cudaError_t CUDARTAPI cudaMemset2D(void *mem, size_t pitch, int c, size_t width, size_t height)
@@ -1986,61 +2304,18 @@ __host__ cudaError_t CUDARTAPI cudaBindTexture(size_t *offset,
const struct cudaChannelFormatDesc *desc,
size_t size __dv(UINT_MAX))
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in cudaBindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
- struct cudaArray *array;
- array = (struct cudaArray*) malloc(sizeof(struct cudaArray));
- array->desc = *desc;
- array->size = size;
- array->width = size;
- array->height = 1;
- array->dimensions = 1;
- array->devPtr = (void*)devPtr;
- array->devPtr32 = (int)(long long)devPtr;
- offset = 0;
- printf("GPGPU-Sim PTX: size = %zu\n", size);
- printf("GPGPU-Sim PTX: texref = %p, array = %p\n", texref, array);
- printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
- printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
- printf("GPGPU-Sim PTX: ChannelFormatDesc: x=%d, y=%d, z=%d, w=%d\n", desc->x, desc->y, desc->z, desc->w);
- printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
- gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
- devPtr = (void*)(long long)array->devPtr32;
- printf("GPGPU-Sim PTX: devPtr = %p\n", devPtr);
- return g_last_cudaError = cudaSuccess;
+ return cudaBindTextureInternal(offset, texref, devPtr, desc, size __dv(UINT_MAX));
}
__host__ cudaError_t CUDARTAPI cudaBindTextureToArray(const struct textureReference *texref, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in cudaBindTextureToArray: %p %p\n", texref, array);
- printf("GPGPU-Sim PTX: devPtr32 = %x\n", array->devPtr32);
- printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
- printf("GPGPU-Sim PTX: Texture Normalized? = %d\n", texref->normalized);
- gpu->gpgpu_ptx_sim_bindTextureToArray(texref, array);
- return g_last_cudaError = cudaSuccess;
+ return cudaBindTextureToArrayInternal(texref, array, desc);
}
-__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref){
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in cudaUnbindTexture: sizeof(struct textureReference) = %zu\n", sizeof(struct textureReference));
- printf("GPGPU-Sim PTX: Name corresponding to textureReference: %s\n", gpu->gpgpu_ptx_sim_findNamefromTexture(texref));
-
- gpu->gpgpu_ptx_sim_unbindTexture(texref);
- return g_last_cudaError = cudaSuccess;
+__host__ cudaError_t CUDARTAPI cudaUnbindTexture(const struct textureReference *texref)
+{
+ return cudaUnbindTextureInternal(texref);
}
__host__ cudaError_t CUDARTAPI cudaGetTextureAlignmentOffset(size_t *offset, const struct textureReference *texref)
@@ -2125,24 +2400,9 @@ __host__ cudaError_t CUDARTAPI cudaLaunch( const char *hostFun )
return cudaLaunchInternal( hostFun );
}
-__host__ cudaError_t CUDARTAPI cudaLaunchKernel ( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
+__host__ cudaError_t CUDARTAPI cudaLaunchKernel( const char* hostFun, dim3 gridDim, dim3 blockDim, const void** args, size_t sharedMem, cudaStream_t stream )
{
-
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- function_info *entry = context->get_kernel(hostFun);
-#if CUDART_VERSION < 10000
- cudaConfigureCallInternal(gridDim, blockDim, sharedMem, stream);
-#endif
- for(unsigned i = 0; i < entry->num_args(); i++){
- std::pair<size_t, unsigned> p = entry->get_param_config(i);
- cudaSetupArgumentInternal(args[i], p.first, p.second);
- }
-
- cudaLaunchInternal(hostFun);
- return g_last_cudaError = cudaSuccess;
+ return cudaLaunchKernelInternal(hostFun, gridDim, blockDim, args, sharedMem, stream);
}
@@ -2528,7 +2788,7 @@ void cuda_runtime_api::extract_ptx_files_using_cuobjdump(CUctx_st *context){
* enabled
* */
void cuda_runtime_api::extract_code_using_cuobjdump(){
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
//prevent the dumping by cuobjdump everytime we execute the code!
const char *override_cuobjdump = getenv("CUOBJDUMP_SIM_FILE");
@@ -2886,7 +3146,7 @@ cuobjdumpPTXSection* cuda_runtime_api::findPTXSection(const std::string identifi
//! Extract the code using cuobjdump and remove unnecessary sections
void cuda_runtime_api::cuobjdumpInit(){
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(gpgpu_ctx);
extract_code_using_cuobjdump(); //extract all the output of cuobjdump to _cuobjdump_*.*
const char* pre_load = getenv("CUOBJDUMP_SIM_FILE");
if (pre_load ==NULL || strlen(pre_load)==0){
@@ -2899,7 +3159,7 @@ void cuda_runtime_api::cuobjdumpInit(){
//! Either submit PTX for simulation or convert SASS to PTXPlus and submit it
void gpgpu_context::cuobjdumpParseBinary(unsigned int handle){
- CUctx_st *context = GPGPUSim_Context();
+ CUctx_st *context = GPGPUSim_Context(this);
if(api->fatbin_registered[handle]) return;
api->fatbin_registered[handle] = true;
std::string fname = api->fatbinmap[handle];
@@ -3140,22 +3400,7 @@ void __cudaRegisterTexture(
int ext
) //passes in a newly created textureReference
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- std::string devStr (deviceName);
- #if (CUDART_VERSION > 4020)
- if (devStr.size() > 2 && devStr.data()[0] == ':' && devStr.data()[1] == ':')
- devStr = devStr.replace(0, 2, "");
- #endif
- CUctx_st *context = GPGPUSim_Context();
- gpgpu_t *gpu = context->get_device()->get_gpgpu();
- printf("GPGPU-Sim PTX: in __cudaRegisterTexture:\n");
- gpu->gpgpu_ptx_sim_bindNameToTexture(devStr.data(), hostVar, dim, norm, ext);
- printf("GPGPU-Sim PTX: int dim = %d\n", dim);
- printf("GPGPU-Sim PTX: int norm = %d\n", norm);
- printf("GPGPU-Sim PTX: int ext = %d\n", ext);
- printf("GPGPU-Sim PTX: Execution warning: Not finished implementing \"%s\"\n", __my_func__ );
+ __cudaRegisterTextureInternal(fatCubinHandle, hostVar, deviceAddress, deviceName, dim, norm, ext);
}
@@ -3187,30 +3432,7 @@ cudaError_t cudaGLMapBufferObject(void** devPtr, GLuint bufferObj)
cudaError_t cudaGLUnmapBufferObject(GLuint bufferObj)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
-#ifdef OPENGL_SUPPORT
- CUctx_st* ctx = GPGPUSim_Context();
- glbmap_entry_t *p = ctx->api->g_glbmap;
- while ( p && p->m_bufferObj != bufferObj )
- p = p->m_next;
- if ( p == NULL )
- return g_last_cudaError = cudaErrorUnknown;
-
- char *data = (char *) calloc(p->m_size,1);
- memcpy_from_gpu( data,(size_t)p->m_devPtr,p->m_size );
- glBufferSubData(GL_ARRAY_BUFFER,0,p->m_size,data);
- free(data);
-
- return g_last_cudaError = cudaSuccess;
-#else
- fflush(stdout);
- fflush(stderr);
- printf("GPGPU-Sim PTX: support for OpenGL integration disabled -- exiting\n");
- fflush(stdout);
- exit(50);
-#endif
+ return cudaGLUnmapBufferObjectInternal(bufferObj);
}
cudaError_t cudaGLUnregisterBufferObject(GLuint bufferObj)
@@ -3322,12 +3544,7 @@ cudaError_t CUDARTAPI cudaRuntimeGetVersion(int *runtimeVersion)
#if CUDART_VERSION >= 3000
__host__ cudaError_t CUDARTAPI cudaFuncSetCacheConfig(const char *func, enum cudaFuncCache cacheConfig )
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- CUctx_st *context = GPGPUSim_Context();
- context->get_device()->get_gpgpu()->set_cache_config(context->get_kernel(func)->get_name(), (FuncCache)cacheConfig);
- return g_last_cudaError = cudaSuccess;
+ return cudaFuncSetCacheConfigInternal(func, cacheConfig);
}
//Jin: hack for cdp
@@ -5033,23 +5250,7 @@ CUresult CUDAAPI cuLaunchKernel(CUfunction f,
void **kernelParams,
void **extra)
{
- if(g_debug_execution >= 3){
- announce_call(__my_func__);
- }
- if (extra!=NULL){
- printf("GPGPU-Sim CUDA DRIVER API: ERROR: Currently do not support void** extra.\n");
- abort();
- }
- const char *hostFun = (const char*) f;
- CUctx_st *context = GPGPUSim_Context();
- function_info *entry = context->get_kernel(hostFun);
- cudaConfigureCallInternal(dim3(gridDimX, gridDimY, gridDimZ), dim3(blockDimX, blockDimY, blockDimZ), sharedMemBytes, (cudaStream_t) hStream);
- for(unsigned i = 0; i < entry->num_args(); i++){
- std::pair<size_t, unsigned> p = entry->get_param_config(i);
- cudaSetupArgument(kernelParams[i], p.first, p.second);
- }
- cudaLaunchInternal(hostFun);
- return CUDA_SUCCESS;
+ return cuLaunchKernelInternal(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra);
}
#endif /* CUDART_VERSION >= 4000 */
generated by cgit v1.3 (git 2.53.0) at 2026-07-10 23:54:50 +0000