{"action":"created","comment":{"url":"https://api.github.com/repos/Theano/libgpuarray/pulls/comments/65942723","id":65942723,"diff_hunk":"@@ -0,0 +1,548 @@\n+/* Includes */\n+#ifdef _MSC_VER\n+#define _CRT_SECURE_NO_WARNINGS\n+#endif\n+\n+#include <assert.h>\n+#include <stdarg.h>\n+#include <stddef.h>\n+#if !defined(_MSC_VER) || _MSC_VER < 1600\n+#include <stdint.h>\n+#endif\n+#include <stdlib.h>\n+#include <string.h>\n+#include <errno.h>\n+\n+#include \"private.h\"\n+#include \"gpuarray/array.h\"\n+#include \"gpuarray/error.h\"\n+#include \"gpuarray/kernel.h\"\n+#include \"gpuarray/util.h\"\n+\n+#include \"util/strb.h\"\n+\n+\n+/* Datatypes */\n+struct GEN_KERNEL_CTX{\n+\tint         numIdx;\n+\tconst int*  isReduced;\n+\tint         numRedIdx;\n+\tint         numFreeIdx;\n+\tconst char* dstMaxType;\n+\tconst char* dstArgmaxType;\n+};\n+typedef struct GEN_KERNEL_CTX GEN_KERNEL_CTX;\n+\n+\n+\n+/* Function prototypes */\n+static int   getRdxIdx              (const int       numIdx,\n+                                     const int*      isReduced);\n+static char* genkernel_maxandargmax (const int       numIdx,\n+                                     const int*      isReduced,\n+                                     const char*     dstMaxType,\n+                                     const char*     dstArgmaxType);\n+static void  appendKernel           (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendTypedefs         (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendPrototype        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendIndexDeclarations(strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendIdxes            (strb*           s,\n+                                     const char*     prologue,\n+                                     const char*     prefix,\n+                                     int             startIdx,\n+                                     int             endIdx,\n+                                     const char*     suffix,\n+                                     const char*     epilogue);\n+static void  appendRangeCalculations(strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendLoops            (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendLoopMacroDefs    (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendLoopOuter        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendLoopInner        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  appendLoopMacroUndefs  (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx);\n+static void  scheduleMaxAndArgmax   (size_t*         blockSize,\n+                                     size_t*         gridSize,\n+                                     const GpuArray* src,\n+                                     const int*      isReduced);\n+static void  invokeMaxAndArgmax     (GpuKernel*      kernel,\n+                                     const GpuArray* src,\n+                                     const int*      isReduced);\n+\n+\n+/* Function implementation */\n+\n+/**\n+ * @brief Computes simultaneously the maxima and the arguments of maxima over\n+ * specified axes of the tensor.\n+ *\n+ * Returns two tensors of identical shape. Both tensors' axes are a subset of\n+ * the axes of the original tensor. The axes to be reduced are specified by\n+ * the caller, and the maxima and arguments of maxima are computed over them.\n+ *\n+ * @param [out] dstMax     The resulting tensor of maxima\n+ * @param [out] dstArgmax  the resulting tensor of arguments at maxima\n+ * @param [in]  src        The source tensor.\n+ * @param [in]  isReduced  Either NULL, or an array of booleans of the same\n+ *                         size as the dimensionality of the source tensor.\n+ *                         Axis k is reduced if isReduced[k] is non-zero,\n+ *                         and is preserved otherwise.\n+ * @return GA_NO_ERROR if the operation was successful, or a non-zero error\n+ *         code otherwise.\n+ */\n+\n+GPUARRAY_PUBLIC int GpuArray_maxandargmax(GpuArray*       dstMax,\n+                                          GpuArray*       dstArgmax,\n+                                          const GpuArray* src,\n+                                          const int*      isReduced){\n+\t/**\n+\t * Generate kernel source code\n+\t */\n+\t\n+\tconst char*  dstMaxType      = gpuarray_get_type(src->typecode) -> cluda_name;\n+\tconst char*  dstArgmaxType   = gpuarray_get_type(GA_LONG)       -> cluda_name;\n+\tconst char*  s               = genkernel_maxandargmax(src->nd,\n+\t                                                      isReduced,\n+\t                                                      dstMaxType,\n+\t                                                      dstArgmaxType);\n+\tif(!s){return GA_MEMORY_ERROR;}\n+\t\n+\t/* Compile it */\n+\tconst int    ARG_TYPECODE[8] = {\n+\t\tGA_POINTER, /* src */\n+\t\tGA_POINTER, /* srcSteps */\n+\t\tGA_POINTER, /* srcSize */\n+\t\tGA_POINTER, /* numBlk */\n+\t\tGA_POINTER, /* dstMax */\n+\t\tGA_POINTER, /* dstMaxSteps */\n+\t\tGA_POINTER, /* dstArgmax */\n+\t\tGA_POINTER  /* dstArgmaxSteps */\n+\t};\n+\tconst size_t l = strlen(s);\n+\tGpuKernel kernel;\n+\tGpuKernel_init(&kernel, 0, 1, &s, &l, \"maxandargmax\",\n+\t               8, ARG_TYPECODE, 0, (char**)0);\n+\t\n+\t/* Invoke it. */\n+\tinvokeMaxAndArgmax(&kernel, src, isReduced);\n+\n+\t/* Return error code */\n+\treturn GA_NO_ERROR;\n+}\n+\n+/**\n+ * Count the number of dimensions to be reduced.\n+ */\n+\n+static int  getRdxIdx(const int numIdx, const int* isReduced){\n+\tint i, countReduced;\n+\tfor(i=0, countReduced = 0;i<numIdx;i++){\n+\t\tcountReduced += !!isReduced[i];\n+\t}\n+\treturn countReduced;\n+}\n+\n+/**\n+ * @brief Generate the kernel code for MaxAndArgmax.\n+ * \n+ * @param [in]  numIdx\n+ * @param [in]  isReduced\n+ * @param [in]  dstMaxType\n+ * @param [in]  dstArgmaxType\n+ * @return A free()'able string containing source code implementing the\n+ *         kernel, or else NULL.\n+ */\n+\n+static char* genkernel_maxandargmax(const int       numIdx,\n+                                    const int*      isReduced,\n+                                    const char*     dstMaxType,\n+                                    const char*     dstArgmaxType){\n+\t/* Obtain the parameters of the reduction. */\n+\tGEN_KERNEL_CTX ctx;\n+\tctx.numIdx        = numIdx;\n+\tctx.isReduced     = isReduced;\n+\tctx.numRedIdx     = getRdxIdx(ctx.numIdx, ctx.isReduced);\n+\tctx.numFreeIdx    = ctx.numIdx - ctx.numRedIdx;\n+\tctx.dstMaxType    = dstMaxType;\n+\tctx.dstArgmaxType = dstArgmaxType;\n+\t\n+\t/**\n+\t * Allocate string buffer.\n+\t */\n+\t\n+\tstrb* s = strb_alloc(5*1024);\n+\tif(!s){return NULL;}\n+\t\n+\t/**\n+\t * Generate kernel\n+\t */\n+\t\n+\tappendKernel(s, &ctx);\n+\t\n+\t/**\n+\t * Return the kernel.\n+\t */\n+\t\n+\tchar* kernelSource = strb_cstr(s);\n+\tstrb_free(s);\n+\treturn kernelSource;\n+}\n+\n+static void  appendKernel           (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tappendTypedefs         (s, ctx);\n+\tappendPrototype        (s, ctx);\n+\tstrb_appends           (s, \"{\\n\");\n+\tappendIndexDeclarations(s, ctx);\n+\tappendRangeCalculations(s, ctx);\n+\tappendLoops            (s, ctx);\n+\tstrb_appends           (s, \"}\\n\");\n+}\n+static void  appendTypedefs         (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tstrb_appends(s, \"/* Typedefs */\\n\");\n+\tstrb_appendf(s, \"typedef %s     T;/* The type of the array being processed. */\\n\", ctx->dstMaxType);\n+\tstrb_appendf(s, \"typedef %s     X;/* Index type: signed 32/64-bit. */\\n\",          ctx->dstArgmaxType);\n+\tstrb_appends(s, \"\\n\");\n+\tstrb_appends(s, \"\\n\");\n+\tstrb_appends(s, \"\\n\");\n+}\n+static void  appendPrototype        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tstrb_appends(s, \"__global__ void maxandargmax(const T*        src,\\n\");\n+\tstrb_appends(s, \"                             const X*        srcSteps,\\n\");\n+\tstrb_appends(s, \"                             const X*        srcSize,\\n\");\n+\tstrb_appends(s, \"                             const X*        blkNum,\\n\");\n+\tstrb_appends(s, \"                             T*              dstMax,\\n\");\n+\tstrb_appends(s, \"                             const X*        dstMaxSteps,\\n\");\n+\tstrb_appends(s, \"                             X*              dstArgmax,\\n\");\n+\tstrb_appends(s, \"                             const X*        dstArgmaxSteps)\");\n+}\n+static void  appendIndexDeclarations(strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tstrb_appends(s, \"\\t/* GPU kernel coordinates. Always 3D. */\\n\");\n+\t\n+\tstrb_appends(s, \"\\tX bi0 = blockIdx.x,  bi1 = blockIdx.y,  bi2 = blockIdx.z;\\n\");\n+\tstrb_appends(s, \"\\tX bd0 = blockDim.x,  bd1 = blockDim.y,  bd2 = blockDim.z;\\n\");\n+\tstrb_appends(s, \"\\tX ti0 = threadIdx.x, ti1 = threadIdx.y, ti2 = threadIdx.z;\\n\");\n+\t\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\t/* Free indices & Reduction indices */\\n\");\n+\t\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"Blk\",   \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"\",      \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"Dim\",   \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"Start\", \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"End\",   \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numIdx,     \"SStep\", \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numFreeIdx, \"MStep\", \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", 0,               ctx->numFreeIdx, \"AStep\", \";\\n\");\n+\tappendIdxes (s, \"\\tX \", \"i\", ctx->numFreeIdx, ctx->numIdx,     \"PDim\",  \";\\n\");\n+\t\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+}\n+static void  appendIdxes            (strb*           s,\n+                                     const char*     prologue,\n+                                     const char*     prefix,\n+                                     int             startIdx,\n+                                     int             endIdx,\n+                                     const char*     suffix,\n+                                     const char*     epilogue){\n+\tint i;\n+\t\n+\tprologue = prologue ? prologue : \"\";\n+\tprefix   = prefix   ? prefix   : \"\";\n+\tsuffix   = suffix   ? suffix   : \"\";\n+\tepilogue = epilogue ? epilogue : \"\";\n+\t\n+\tstrb_appends(s, prologue);\n+\tfor(i=startIdx;i<endIdx;i++){\n+\t\tstrb_appendf(s, \"%s%d%s%s\", prefix, i, suffix, &\",\"[i==endIdx-1]);\n+\t}\n+\tstrb_appends(s, epilogue);\n+}\n+static void  appendRangeCalculations(strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tint i;\n+\t\n+\tstrb_appends(s, \"\\t/* Compute ranges for this thread. */\\n\");\n+\t\n+\tfor(i=0;i<ctx->numIdx    ;i++){/* i*Dim   = srcSize[*]; */\n+\t\tstrb_appendf(s, \"\\ti%dDim   = srcSize[%d];\\n\", i, i);\n+\t}\n+\tfor(i=0;i<ctx->numIdx    ;i++){/* i*SStep = srcSteps[*]; */\n+\t\tstrb_appendf(s, \"\\ti%dSStep = srcSteps[%d];\\n\", i, i);\n+\t}\n+\tfor(i=0;i<ctx->numFreeIdx;i++){/* i*MStep = dstMaxSteps[*]; */\n+\t\tstrb_appendf(s, \"\\ti%dMStep = dstMaxSteps[%d];\\n\", i, i);\n+\t}\n+\tfor(i=0;i<ctx->numFreeIdx;i++){/* i*AStep = dstArgmaxSteps[*]; */\n+\t\tstrb_appendf(s, \"\\ti%dMStep = dstArgmaxSteps[%d];\\n\", i, i);\n+\t}\n+\tfor(i=0;i<ctx->numIdx    ;i++){/* i*Blk   = numBlk[*]; */\n+\t\tstrb_appendf(s, \"\\ti%dBlk   = numBlk[%d];\\n\", i, i);\n+\t}\n+\tfor(i=ctx->numIdx-1;i>=ctx->numFreeIdx;i--){/* i*PDim  = ...; */\n+\t\t/**\n+\t\t * If this is the last index, it's the first cumulative dimension\n+\t\t * product we generate, and thus we initialize to 1.\n+\t\t */\n+\t\t\n+\t\tif(i == ctx->numIdx-1){\n+\t\t\tstrb_appendf(s, \"\\ti%dPDim  = 1;\\n\", i);\n+\t\t}else{\n+\t\t\tstrb_appendf(s, \"\\ti%dPDim  = i%dPDim * i%dDim;\\n\", i, i+1, i);\n+\t\t}\n+\t}\n+\tfor(i=0;i<ctx->numIdx    ;i++){/* i*Start = ...; */\n+\t\t/**\n+\t\t * The first 3 dimensions get to rely on hardware loops.\n+\t\t * The others, if any, have to use software looping beginning at 0.\n+\t\t */\n+\t\t\n+\t\tif(i < 3){\n+\t\t\tstrb_appendf(s, \"\\ti%dStart = ((bi%d * bd%d) + ti%d) * i%dBlk;\\n\", i, i, i, i, i);\n+\t\t}else{\n+\t\t\tstrb_appendf(s, \"\\ti%dStart = 0;\\n\", i);\n+\t\t}\n+\t}\n+\tfor(i=0;i<ctx->numIdx    ;i++){/* i*End   = ...; */\n+\t\t/**\n+\t\t * The first 3 dimensions get to rely on hardware loops.\n+\t\t * The others, if any, have to use software looping beginning at 0.\n+\t\t */\n+\t\t\n+\t\tif(i < 3){\n+\t\t\tstrb_appendf(s, \"\\ti%dEnd   = i%dStart + bd%d * i%dBlk;\\n\", i, i, i, i);\n+\t\t}else{\n+\t\t\tstrb_appendf(s, \"\\ti%dEnd   = i%dStart + i%dDim;\\n\", i, i, i);\n+\t\t}\n+\t}\n+\t\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+}\n+static void  appendLoops            (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tstrb_appends(s, \"\\t/**\\n\");\n+\tstrb_appends(s, \"\\t * FREE LOOPS.\\n\");\n+\tstrb_appends(s, \"\\t */\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\t\n+\tappendLoopMacroDefs  (s, ctx);\n+\tappendLoopOuter      (s, ctx);\n+\tappendLoopMacroUndefs(s, ctx);\n+}\n+static void  appendLoopMacroDefs    (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tint i;\n+\t\n+\t/**\n+\t * FOROVER Macro\n+\t */\n+\t\n+\tstrb_appends(s, \"#define FOROVER(idx)    for(i##idx = i##idx##Start; i##idx < i##idx##End; i##idx++)\\n\");\n+\t\n+\t/**\n+\t * ESCAPE Macro\n+\t */\n+\t\n+\tstrb_appends(s, \"#define ESCAPE(idx)     if(i##idx >= i##idx##Dim){continue;}\\n\");\n+\t\n+\t/**\n+\t * SRCINDEXER Macro\n+\t */\n+\t\n+\tappendIdxes (s, \"#define SRCINDEXER(\", \"i\", 0, ctx->numIdx, \"\", \")   src[\");\n+\tfor(i=0;i<ctx->numIdx;i++){\n+\t\tstrb_appendf(s, \"i%d*i%dSStep + \\\\\\n                                            \", i, i);\n+\t}\n+\tstrb_appends(s, \"0]\\n\");\n+\t\n+\t/**\n+\t * RDXINDEXER Macro\n+\t */\n+\t\n+\tappendIdxes (s, \"#define RDXINDEXER(\", \"i\", ctx->numFreeIdx, ctx->numIdx, \"\", \")              (\");\n+\tfor(i=ctx->numFreeIdx;i<ctx->numIdx;i++){\n+\t\tstrb_appendf(s, \"i%d*i%dPDim + \\\\\\n                                        \", i, i);\n+\t}\n+\tstrb_appends(s, \"0)\\n\");\n+\t\n+\t/**\n+\t * DSTMINDEXER Macro\n+\t */\n+\t\n+\tappendIdxes (s, \"#define DSTMINDEXER(\", \"i\", 0, ctx->numFreeIdx, \"\", \")        dstMax[\");\n+\tfor(i=0;i<ctx->numFreeIdx;i++){\n+\t\tstrb_appendf(s, \"i%d*i%dMStep + \\\\\\n                                                  \", i, i);\n+\t}\n+\tstrb_appends(s, \"0]\\n\");\n+\t\n+\t/**\n+\t * DSTAINDEXER Macro\n+\t */\n+\t\n+\tappendIdxes (s, \"#define DSTAINDEXER(\", \"i\", 0, ctx->numFreeIdx, \"\", \")        dstArgmax[\");\n+\tfor(i=0;i<ctx->numFreeIdx;i++){\n+\t\tstrb_appendf(s, \"i%d*i%dAStep + \\\\\\n                                                     \", i, i);\n+\t}\n+\tstrb_appends(s, \"0]\\n\");\n+}\n+static void  appendLoopOuter        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tint i;\n+\t\n+\t/**\n+\t * Outer Loop Header Generation\n+\t */\n+\t\n+\tfor(i=0;i<ctx->numFreeIdx;i++){\n+\t\tstrb_appendf(s, \"\\tFOROVER(%d){ESCAPE(%d)\\n\", i, i);\n+\t}\n+\t\n+\t/**\n+\t * Inner Loop Generation\n+\t */\n+\t\n+\tappendLoopInner(s, ctx);\n+\t\n+\t/**\n+\t * Outer Loop Trailer Generation\n+\t */\n+\t\n+\tfor(i=0;i<ctx->numFreeIdx;i++){\n+\t\tstrb_appends(s, \"\\t}\\n\");\n+\t}\n+}\n+static void  appendLoopInner        (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tint i;\n+\t\n+\t/**\n+\t * Inner Loop Prologue\n+\t */\n+\t\n+\tstrb_appends(s, \"\\t/**\\n\");\n+\tstrb_appends(s, \"\\t * Reduction initialization.\\n\");\n+\tstrb_appends(s, \"\\t */\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\t\n+\tappendIdxes (s, \"\\tT maxV = SRCINDEXER(\", \"i\", 0, ctx->numFreeIdx, \"\", \"\");\n+\tif(ctx->numFreeIdx && ctx->numRedIdx){strb_appends(s, \",\");}\n+\tappendIdxes (s, \"\", \"i\", ctx->numFreeIdx, ctx->numIdx, \"Start\", \");\\n\");\n+\t\n+\tappendIdxes (s, \"\\tX maxI = RDXINDEXER(\", \"i\", ctx->numFreeIdx, ctx->numIdx, \"Start\", \");\\n\");\n+\t\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\t/**\\n\");\n+\tstrb_appends(s, \"\\t * REDUCTION LOOPS.\\n\");\n+\tstrb_appends(s, \"\\t */\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\t\n+\t/**\n+\t * Inner Loop Header Generation\n+\t */\n+\t\n+\tfor(i=ctx->numFreeIdx;i<ctx->numIdx;i++){\n+\t\tstrb_appendf(s, \"\\tFOROVER(%d){ESCAPE(%d)\\n\", i, i);\n+\t}\n+\t\n+\t/**\n+\t * Inner Loop Body Generation\n+\t */\n+\t\n+\tappendIdxes (s, \"\\tT V = SRCINDEXER(\", \"i\", 0, ctx->numIdx, \"\", \");\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\tstrb_appends(s, \"\\tif(V > maxV){\\n\");\n+\tstrb_appends(s, \"\\t\\tmaxV = V;\\n\");\n+\tappendIdxes (s, \"\\t\\tmaxI = RDXINDEXER(\", \"i\", ctx->numFreeIdx, ctx->numIdx, \"\", \");\\n\");\n+\tstrb_appends(s, \"\\t}\\n\");\n+\t\n+\t/**\n+\t * Inner Loop Trailer Generation\n+\t */\n+\t\n+\tfor(i=ctx->numFreeIdx;i<ctx->numIdx;i++){\n+\t\tstrb_appends(s, \"\\t}\\n\");\n+\t}\n+\tstrb_appends(s, \"\\t\\n\");\n+\t\n+\t/**\n+\t * Inner Loop Epilogue Generation\n+\t */\n+\t\n+\tstrb_appends(s, \"\\t/**\\n\");\n+\tstrb_appends(s, \"\\t * Destination writeback.\\n\");\n+\tstrb_appends(s, \"\\t */\\n\");\n+\tstrb_appends(s, \"\\t\\n\");\n+\tappendIdxes (s, \"\\tDSTMINDEXER(\", \"i\", 0, ctx->numFreeIdx, \"\", \") = maxV;\\n\");\n+\tappendIdxes (s, \"\\tDSTAINDEXER(\", \"i\", 0, ctx->numFreeIdx, \"\", \") = maxI;\\n\");\n+}\n+static void  appendLoopMacroUndefs  (strb*           s,\n+                                     GEN_KERNEL_CTX* ctx){\n+\tstrb_appends(s, \"\\t#undef FOROVER\\n\");\n+\tstrb_appends(s, \"\\t#undef ESCAPE\\n\");\n+\tstrb_appends(s, \"\\t#undef SRCINDEXER\\n\");\n+\tstrb_appends(s, \"\\t#undef RDXINDEXER\\n\");\n+\tstrb_appends(s, \"\\t#undef DSTMINDEXER\\n\");\n+\tstrb_appends(s, \"\\t#undef DSTAINDEXER\\n\");\n+}\n+\n+\n+/**\n+ * FIXME: Implement a working scheduler and invoker.\n+ * \n+ * To schedule effectively the work across several dimensions of possibly\n+ * ugly numbers, taking into account the limitations on thread & block\n+ * scheduling, will require some library capable of providing a \"factoring\"\n+ * of the tensor dimensions into \"nice\" prime numbers. Their product may\n+ * be allowed to be larger than the original number, provided it remains\n+ * within bounds.\n+ */\n+\n+/**\n+ * Compute a good thread block size / grid size for Nvidia.\n+ */\n+\n+static void  scheduleMaxAndArgmax   (size_t*         blockSize,\n+                                     size_t*         gridSize,\n+                                     const GpuArray* src,\n+                                     const int*      isReduced){\n+\t//int maxThreadPerBlock = 1024;\n+\t//int numFreeIdx = src->nd - getRdxIdx(src->nd, isReduced);\n+\t\n+\t/* Naive solution. Optimization is a tough problem. */\n+\tblockSize[0] = blockSize[1] = blockSize[2] = 1;\n+\tgridSize [0] = gridSize [1] = gridSize [2] = 1;\n+}\n+\n+/**\n+ * Invoke the kernel.\n+ */\n+\n+static void  invokeMaxAndArgmax     (GpuKernel*      kernel,\n+                                     const GpuArray* src,\n+                                     const int*      isReduced){\n+\tsize_t blockSize[3];\n+\tsize_t gridSize[3];\n+\t\n+\tscheduleMaxAndArgmax(blockSize, gridSize, src, isReduced);\n+\tGpuKernel_call(kernel,\n+\t               getRdxIdx(src->nd, isReduced),\n+\t               blockSize,\n+\t               gridSize,\n+\t               0,\n+\t               NULL);","path":"src/gpuarray_reduction.c","position":546,"original_position":546,"commit_id":"850b5834353bfd703c8e3a8c1208ce7e58fddcd2","original_commit_id":"850b5834353bfd703c8e3a8c1208ce7e58fddcd2","user":{"login":"abergeron","id":1002504,"avatar_url":"https://avatars.githubusercontent.com/u/1002504?v=3","gravatar_id":"","url":"https://api.github.com/users/abergeron","html_url":"https://github.com/abergeron","followers_url":"https://api.github.com/users/abergeron/followers","following_url":"https://api.github.com/users/abergeron/following{/other_user}","gists_url":"https://api.github.com/users/abergeron/gists{/gist_id}","starred_url":"https://api.github.com/users/abergeron/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/abergeron/subscriptions","organizations_url":"https://api.github.com/users/abergeron/orgs","repos_url":"https://api.github.com/users/abergeron/repos","events_url":"https://api.github.com/users/abergeron/events{/privacy}","received_events_url":"https://api.github.com/users/abergeron/received_events","type":"User","site_admin":false},"body":"You don't pass any arguments to the kernel anywhere.  I don't see how this can work.","created_at":"2016-06-06T18:28:34Z","updated_at":"2016-06-06T18:28:34Z","html_url":"https://github.com/Theano/libgpuarray/pull/185#discussion_r65942723","pull_request_url":"https://api.github.com/repos/Theano/libgpuarray/pulls/185","_links":{"self":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/comments/65942723"},"html":{"href":"https://github.com/Theano/libgpuarray/pull/185#discussion_r65942723"},"pull_request":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/185"}}},"pull_request":{"url":"https://api.github.com/repos/Theano/libgpuarray/pulls/185","id":72593326,"html_url":"https://github.com/Theano/libgpuarray/pull/185","diff_url":"https://github.com/Theano/libgpuarray/pull/185.diff","patch_url":"https://github.com/Theano/libgpuarray/pull/185.patch","issue_url":"https://api.github.com/repos/Theano/libgpuarray/issues/185","number":185,"state":"open","locked":false,"title":"Initial support for MaxAndArgmax and GpuArray_maxandargmax().","user":{"login":"obilaniu","id":3268344,"avatar_url":"https://avatars.githubusercontent.com/u/3268344?v=3","gravatar_id":"","url":"https://api.github.com/users/obilaniu","html_url":"https://github.com/obilaniu","followers_url":"https://api.github.com/users/obilaniu/followers","following_url":"https://api.github.com/users/obilaniu/following{/other_user}","gists_url":"https://api.github.com/users/obilaniu/gists{/gist_id}","starred_url":"https://api.github.com/users/obilaniu/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/obilaniu/subscriptions","organizations_url":"https://api.github.com/users/obilaniu/orgs","repos_url":"https://api.github.com/users/obilaniu/repos","events_url":"https://api.github.com/users/obilaniu/events{/privacy}","received_events_url":"https://api.github.com/users/obilaniu/received_events","type":"User","site_admin":false},"body":"What remains incomplete is a scheduler for the kernel, as noted in the\r\nsource code.","created_at":"2016-06-04T05:05:40Z","updated_at":"2016-06-06T18:28:34Z","closed_at":null,"merged_at":null,"merge_commit_sha":"5f4345e7ad33f41a9af1d5b07cfa46ce9c544e91","assignee":null,"milestone":null,"commits_url":"https://api.github.com/repos/Theano/libgpuarray/pulls/185/commits","review_comments_url":"https://api.github.com/repos/Theano/libgpuarray/pulls/185/comments","review_comment_url":"https://api.github.com/repos/Theano/libgpuarray/pulls/comments{/number}","comments_url":"https://api.github.com/repos/Theano/libgpuarray/issues/185/comments","statuses_url":"https://api.github.com/repos/Theano/libgpuarray/statuses/850b5834353bfd703c8e3a8c1208ce7e58fddcd2","head":{"label":"obilaniu:master","ref":"master","sha":"850b5834353bfd703c8e3a8c1208ce7e58fddcd2","user":{"login":"obilaniu","id":3268344,"avatar_url":"https://avatars.githubusercontent.com/u/3268344?v=3","gravatar_id":"","url":"https://api.github.com/users/obilaniu","html_url":"https://github.com/obilaniu","followers_url":"https://api.github.com/users/obilaniu/followers","following_url":"https://api.github.com/users/obilaniu/following{/other_user}","gists_url":"https://api.github.com/users/obilaniu/gists{/gist_id}","starred_url":"https://api.github.com/users/obilaniu/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/obilaniu/subscriptions","organizations_url":"https://api.github.com/users/obilaniu/orgs","repos_url":"https://api.github.com/users/obilaniu/repos","events_url":"https://api.github.com/users/obilaniu/events{/privacy}","received_events_url":"https://api.github.com/users/obilaniu/received_events","type":"User","site_admin":false},"repo":{"id":60393896,"name":"libgpuarray","full_name":"obilaniu/libgpuarray","owner":{"login":"obilaniu","id":3268344,"avatar_url":"https://avatars.githubusercontent.com/u/3268344?v=3","gravatar_id":"","url":"https://api.github.com/users/obilaniu","html_url":"https://github.com/obilaniu","followers_url":"https://api.github.com/users/obilaniu/followers","following_url":"https://api.github.com/users/obilaniu/following{/other_user}","gists_url":"https://api.github.com/users/obilaniu/gists{/gist_id}","starred_url":"https://api.github.com/users/obilaniu/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/obilaniu/subscriptions","organizations_url":"https://api.github.com/users/obilaniu/orgs","repos_url":"https://api.github.com/users/obilaniu/repos","events_url":"https://api.github.com/users/obilaniu/events{/privacy}","received_events_url":"https://api.github.com/users/obilaniu/received_events","type":"User","site_admin":false},"private":false,"html_url":"https://github.com/obilaniu/libgpuarray","description":"Library to manipulate tensors on the GPU.","fork":true,"url":"https://api.github.com/repos/obilaniu/libgpuarray","forks_url":"https://api.github.com/repos/obilaniu/libgpuarray/forks","keys_url":"https://api.github.com/repos/obilaniu/libgpuarray/keys{/key_id}","collaborators_url":"https://api.github.com/repos/obilaniu/libgpuarray/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/obilaniu/libgpuarray/teams","hooks_url":"https://api.github.com/repos/obilaniu/libgpuarray/hooks","issue_events_url":"https://api.github.com/repos/obilaniu/libgpuarray/issues/events{/number}","events_url":"https://api.github.com/repos/obilaniu/libgpuarray/events","assignees_url":"https://api.github.com/repos/obilaniu/libgpuarray/assignees{/user}","branches_url":"https://api.github.com/repos/obilaniu/libgpuarray/branches{/branch}","tags_url":"https://api.github.com/repos/obilaniu/libgpuarray/tags","blobs_url":"https://api.github.com/repos/obilaniu/libgpuarray/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/obilaniu/libgpuarray/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/obilaniu/libgpuarray/git/refs{/sha}","trees_url":"https://api.github.com/repos/obilaniu/libgpuarray/git/trees{/sha}","statuses_url":"https://api.github.com/repos/obilaniu/libgpuarray/statuses/{sha}","languages_url":"https://api.github.com/repos/obilaniu/libgpuarray/languages","stargazers_url":"https://api.github.com/repos/obilaniu/libgpuarray/stargazers","contributors_url":"https://api.github.com/repos/obilaniu/libgpuarray/contributors","subscribers_url":"https://api.github.com/repos/obilaniu/libgpuarray/subscribers","subscription_url":"https://api.github.com/repos/obilaniu/libgpuarray/subscription","commits_url":"https://api.github.com/repos/obilaniu/libgpuarray/commits{/sha}","git_commits_url":"https://api.github.com/repos/obilaniu/libgpuarray/git/commits{/sha}","comments_url":"https://api.github.com/repos/obilaniu/libgpuarray/comments{/number}","issue_comment_url":"https://api.github.com/repos/obilaniu/libgpuarray/issues/comments{/number}","contents_url":"https://api.github.com/repos/obilaniu/libgpuarray/contents/{+path}","compare_url":"https://api.github.com/repos/obilaniu/libgpuarray/compare/{base}...{head}","merges_url":"https://api.github.com/repos/obilaniu/libgpuarray/merges","archive_url":"https://api.github.com/repos/obilaniu/libgpuarray/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/obilaniu/libgpuarray/downloads","issues_url":"https://api.github.com/repos/obilaniu/libgpuarray/issues{/number}","pulls_url":"https://api.github.com/repos/obilaniu/libgpuarray/pulls{/number}","milestones_url":"https://api.github.com/repos/obilaniu/libgpuarray/milestones{/number}","notifications_url":"https://api.github.com/repos/obilaniu/libgpuarray/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/obilaniu/libgpuarray/labels{/name}","releases_url":"https://api.github.com/repos/obilaniu/libgpuarray/releases{/id}","deployments_url":"https://api.github.com/repos/obilaniu/libgpuarray/deployments","created_at":"2016-06-04T04:54:54Z","updated_at":"2016-06-04T04:54:55Z","pushed_at":"2016-06-04T04:58:28Z","git_url":"git://github.com/obilaniu/libgpuarray.git","ssh_url":"git@github.com:obilaniu/libgpuarray.git","clone_url":"https://github.com/obilaniu/libgpuarray.git","svn_url":"https://github.com/obilaniu/libgpuarray","homepage":null,"size":1559,"stargazers_count":0,"watchers_count":0,"language":"C","has_issues":false,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":0,"mirror_url":null,"open_issues_count":0,"forks":0,"open_issues":0,"watchers":0,"default_branch":"master"}},"base":{"label":"Theano:master","ref":"master","sha":"1bc15f489a521d51980d105e2c1e5367c8959440","user":{"login":"Theano","id":970517,"avatar_url":"https://avatars.githubusercontent.com/u/970517?v=3","gravatar_id":"","url":"https://api.github.com/users/Theano","html_url":"https://github.com/Theano","followers_url":"https://api.github.com/users/Theano/followers","following_url":"https://api.github.com/users/Theano/following{/other_user}","gists_url":"https://api.github.com/users/Theano/gists{/gist_id}","starred_url":"https://api.github.com/users/Theano/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/Theano/subscriptions","organizations_url":"https://api.github.com/users/Theano/orgs","repos_url":"https://api.github.com/users/Theano/repos","events_url":"https://api.github.com/users/Theano/events{/privacy}","received_events_url":"https://api.github.com/users/Theano/received_events","type":"Organization","site_admin":false},"repo":{"id":19312996,"name":"libgpuarray","full_name":"Theano/libgpuarray","owner":{"login":"Theano","id":970517,"avatar_url":"https://avatars.githubusercontent.com/u/970517?v=3","gravatar_id":"","url":"https://api.github.com/users/Theano","html_url":"https://github.com/Theano","followers_url":"https://api.github.com/users/Theano/followers","following_url":"https://api.github.com/users/Theano/following{/other_user}","gists_url":"https://api.github.com/users/Theano/gists{/gist_id}","starred_url":"https://api.github.com/users/Theano/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/Theano/subscriptions","organizations_url":"https://api.github.com/users/Theano/orgs","repos_url":"https://api.github.com/users/Theano/repos","events_url":"https://api.github.com/users/Theano/events{/privacy}","received_events_url":"https://api.github.com/users/Theano/received_events","type":"Organization","site_admin":false},"private":false,"html_url":"https://github.com/Theano/libgpuarray","description":"Library to manipulate tensors on the GPU.","fork":false,"url":"https://api.github.com/repos/Theano/libgpuarray","forks_url":"https://api.github.com/repos/Theano/libgpuarray/forks","keys_url":"https://api.github.com/repos/Theano/libgpuarray/keys{/key_id}","collaborators_url":"https://api.github.com/repos/Theano/libgpuarray/collaborators{/collaborator}","teams_url":"https://api.github.com/repos/Theano/libgpuarray/teams","hooks_url":"https://api.github.com/repos/Theano/libgpuarray/hooks","issue_events_url":"https://api.github.com/repos/Theano/libgpuarray/issues/events{/number}","events_url":"https://api.github.com/repos/Theano/libgpuarray/events","assignees_url":"https://api.github.com/repos/Theano/libgpuarray/assignees{/user}","branches_url":"https://api.github.com/repos/Theano/libgpuarray/branches{/branch}","tags_url":"https://api.github.com/repos/Theano/libgpuarray/tags","blobs_url":"https://api.github.com/repos/Theano/libgpuarray/git/blobs{/sha}","git_tags_url":"https://api.github.com/repos/Theano/libgpuarray/git/tags{/sha}","git_refs_url":"https://api.github.com/repos/Theano/libgpuarray/git/refs{/sha}","trees_url":"https://api.github.com/repos/Theano/libgpuarray/git/trees{/sha}","statuses_url":"https://api.github.com/repos/Theano/libgpuarray/statuses/{sha}","languages_url":"https://api.github.com/repos/Theano/libgpuarray/languages","stargazers_url":"https://api.github.com/repos/Theano/libgpuarray/stargazers","contributors_url":"https://api.github.com/repos/Theano/libgpuarray/contributors","subscribers_url":"https://api.github.com/repos/Theano/libgpuarray/subscribers","subscription_url":"https://api.github.com/repos/Theano/libgpuarray/subscription","commits_url":"https://api.github.com/repos/Theano/libgpuarray/commits{/sha}","git_commits_url":"https://api.github.com/repos/Theano/libgpuarray/git/commits{/sha}","comments_url":"https://api.github.com/repos/Theano/libgpuarray/comments{/number}","issue_comment_url":"https://api.github.com/repos/Theano/libgpuarray/issues/comments{/number}","contents_url":"https://api.github.com/repos/Theano/libgpuarray/contents/{+path}","compare_url":"https://api.github.com/repos/Theano/libgpuarray/compare/{base}...{head}","merges_url":"https://api.github.com/repos/Theano/libgpuarray/merges","archive_url":"https://api.github.com/repos/Theano/libgpuarray/{archive_format}{/ref}","downloads_url":"https://api.github.com/repos/Theano/libgpuarray/downloads","issues_url":"https://api.github.com/repos/Theano/libgpuarray/issues{/number}","pulls_url":"https://api.github.com/repos/Theano/libgpuarray/pulls{/number}","milestones_url":"https://api.github.com/repos/Theano/libgpuarray/milestones{/number}","notifications_url":"https://api.github.com/repos/Theano/libgpuarray/notifications{?since,all,participating}","labels_url":"https://api.github.com/repos/Theano/libgpuarray/labels{/name}","releases_url":"https://api.github.com/repos/Theano/libgpuarray/releases{/id}","deployments_url":"https://api.github.com/repos/Theano/libgpuarray/deployments","created_at":"2014-04-30T12:33:14Z","updated_at":"2016-05-26T04:39:17Z","pushed_at":"2016-06-06T16:38:21Z","git_url":"git://github.com/Theano/libgpuarray.git","ssh_url":"git@github.com:Theano/libgpuarray.git","clone_url":"https://github.com/Theano/libgpuarray.git","svn_url":"https://github.com/Theano/libgpuarray","homepage":null,"size":2075,"stargazers_count":70,"watchers_count":70,"language":"C","has_issues":true,"has_downloads":true,"has_wiki":true,"has_pages":false,"forks_count":42,"mirror_url":null,"open_issues_count":31,"forks":42,"open_issues":31,"watchers":70,"default_branch":"master"}},"_links":{"self":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/185"},"html":{"href":"https://github.com/Theano/libgpuarray/pull/185"},"issue":{"href":"https://api.github.com/repos/Theano/libgpuarray/issues/185"},"comments":{"href":"https://api.github.com/repos/Theano/libgpuarray/issues/185/comments"},"review_comments":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/185/comments"},"review_comment":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/comments{/number}"},"commits":{"href":"https://api.github.com/repos/Theano/libgpuarray/pulls/185/commits"},"statuses":{"href":"https://api.github.com/repos/Theano/libgpuarray/statuses/850b5834353bfd703c8e3a8c1208ce7e58fddcd2"}}}}
{"actor":{"display_login":"abergeron"}}