#NEWLINE# def c_code( self , node , nodename , inp , out , sub ) : #NEWLINE# #INDENT# ( x , b , y_idx ) = inp #NEWLINE# ( nll , sm , am ) = out #NEWLINE# classname = self . __class__ . __name__ #NEWLINE# fail = sub [ 'fail' ] #NEWLINE# sio = StringIO ( ) #NEWLINE# print ( ( '\n        if (CudaNdarray_NDIM(%(y_idx)s) != 1)\n        {\n            PyErr_SetString(PyExc_ValueError, "y_idx not 1d tensor");\n            %(fail)s;\n        }\n        if (CudaNdarray_NDIM(%(x)s) != 2)\n        {\n            PyErr_SetString(PyExc_ValueError, "x not 2d tensor");\n            %(fail)s;\n        }\n        if (CudaNdarray_NDIM(%(b)s) != 1)\n        {\n            PyErr_SetString(PyExc_ValueError, "b not 1d tensor");\n            %(fail)s;\n        }\n        if (CudaNdarray_HOST_DIMS(%(x)s)[0] !=\n            CudaNdarray_HOST_DIMS(%(y_idx)s)[0])\n        {\n            PyErr_SetString(PyExc_ValueError,\n                            "dimension mismatch in x,y_idx arguments");\n            %(fail)s;\n        }\n        if (CudaNdarray_HOST_DIMS(%(x)s)[1] != CudaNdarray_HOST_DIMS(%(b)s)[0])\n        {\n            PyErr_SetString(PyExc_ValueError,\n                            "dimension mismatch in x,b arguments");\n            %(fail)s;\n        }\n        if ((NULL == %(nll)s) //initial condition\n            || (CudaNdarray_HOST_DIMS(%(nll)s)[0] !=\n                CudaNdarray_HOST_DIMS(%(y_idx)s)[0]))\n        {\n            Py_XDECREF(%(nll)s);\n            %(nll)s = (CudaNdarray*)CudaNdarray_NewDims(1,\n                CudaNdarray_HOST_DIMS(%(y_idx)s));\n            if(!%(nll)s)\n            {\n                %(fail)s;\n            }\n        }\n        if ((NULL == %(sm)s)\n            || (CudaNdarray_HOST_DIMS(%(sm)s)[0] !=\n                CudaNdarray_HOST_DIMS(%(x)s)[0])\n            || (CudaNdarray_HOST_DIMS(%(sm)s)[1] !=\n                CudaNdarray_HOST_DIMS(%(x)s)[1]))\n        {\n            Py_XDECREF(%(sm)s);\n            %(sm)s = (CudaNdarray*) CudaNdarray_NewDims(2,\n                CudaNdarray_HOST_DIMS(%(x)s));\n            if(!%(sm)s)\n            {\n                PyErr_SetString(PyExc_MemoryError,\n                                "failed to alloc sm output");\n                // no need to decref cnda_nll, the cleanup code should do it up\n                %(fail)s;\n            }\n        }\n        if ((NULL == %(am)s)\n            || (CudaNdarray_HOST_DIMS(%(am)s)[0] !=\n                CudaNdarray_HOST_DIMS(%(y_idx)s)[0]))\n        {\n            Py_XDECREF(%(am)s);\n            %(am)s = (CudaNdarray*) CudaNdarray_NewDims(1,\n                CudaNdarray_HOST_DIMS(%(y_idx)s));\n            if(!%(am)s)\n            {\n                PyErr_SetString(PyExc_MemoryError,\n                                "failed to alloc am output");\n                // no need to decref nll and sm,\n                // the cleanup code should do it up\n                %(fail)s;\n            }\n        }\n        {\n            int n_blocks = std::min(CudaNdarray_HOST_DIMS(%(x)s)[0],\n                                    NUM_VECTOR_OP_BLOCKS);\n            int n_threads = std::min(CudaNdarray_HOST_DIMS(%(x)s)[1],\n                                     NUM_VECTOR_OP_THREADS_PER_BLOCK);\n            int n_shared_bytes = n_threads * sizeof(float);\n            k_xent_sm_1hot_bias<<<n_blocks, n_threads, n_shared_bytes>>>(\n                CudaNdarray_HOST_DIMS(%(x)s)[0],\n                CudaNdarray_HOST_DIMS(%(x)s)[1],\n                CudaNdarray_DEV_DATA(%(x)s),\n                CudaNdarray_HOST_STRIDES(%(x)s)[0],\n                CudaNdarray_HOST_STRIDES(%(x)s)[1],\n                CudaNdarray_DEV_DATA(%(b)s),\n                CudaNdarray_HOST_STRIDES(%(b)s)[0],\n                CudaNdarray_DEV_DATA(%(y_idx)s),\n                CudaNdarray_HOST_STRIDES(%(y_idx)s)[0],\n                CudaNdarray_DEV_DATA(%(nll)s),\n                CudaNdarray_HOST_STRIDES(%(nll)s)[0],\n                CudaNdarray_DEV_DATA(%(sm)s),\n                CudaNdarray_HOST_STRIDES(%(sm)s)[0],\n                CudaNdarray_HOST_STRIDES(%(sm)s)[1],\n                CudaNdarray_DEV_DATA(%(am)s),\n                CudaNdarray_HOST_STRIDES(%(am)s)[0]);\n            CNDA_THREAD_SYNC;\n            cudaError_t err = cudaGetLastError();\n            if (cudaSuccess != err)\n            {\n                PyErr_Format(PyExc_RuntimeError,\n                             "Cuda error: %(classname)s %(nodename)s: %%s.\\n"\n                             "The kernel was launched with %%d threads,"\n                             " %%d blocks and %%d shared memory\\n",\n                             cudaGetErrorString(err),\n                             n_threads, n_blocks, n_shared_bytes);\n                // no need to decref output vars the cleanup code will do it\n                %(fail)s;\n            }\n        }\n        ' % locals ( ) ) , file = x ) #NEWLINE# return sio . getvalue ( )