{"action":"created","issue":{"url":"https://api.github.com/repos/hughperkins/clnn/issues/38","repository_url":"https://api.github.com/repos/hughperkins/clnn","labels_url":"https://api.github.com/repos/hughperkins/clnn/issues/38/labels{/name}","comments_url":"https://api.github.com/repos/hughperkins/clnn/issues/38/comments","events_url":"https://api.github.com/repos/hughperkins/clnn/issues/38/events","html_url":"https://github.com/hughperkins/clnn/issues/38","id":143946624,"number":38,"title":"clnn for neuralconvo support. LogSoftMax.lua:23","user":{"login":"tigerneil","id":5799436,"avatar_url":"https://avatars.githubusercontent.com/u/5799436?v=3","gravatar_id":"","url":"https://api.github.com/users/tigerneil","html_url":"https://github.com/tigerneil","followers_url":"https://api.github.com/users/tigerneil/followers","following_url":"https://api.github.com/users/tigerneil/following{/other_user}","gists_url":"https://api.github.com/users/tigerneil/gists{/gist_id}","starred_url":"https://api.github.com/users/tigerneil/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/tigerneil/subscriptions","organizations_url":"https://api.github.com/users/tigerneil/orgs","repos_url":"https://api.github.com/users/tigerneil/repos","events_url":"https://api.github.com/users/tigerneil/events{/privacy}","received_events_url":"https://api.github.com/users/tigerneil/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"milestone":null,"comments":1,"created_at":"2016-03-28T10:29:56Z","updated_at":"2016-03-28T10:32:31Z","closed_at":null,"body":"Hey I am trying to add cltorch for [neuralconvo](https://github.com/macournoyer/neuralconvo) . and my code is here [cltorch support neural conversation model](https://github.com/tigerneil/neuralconvo-cl). \r\n\r\n\r\n```\r\nbash-3.2$ th train.lua --dataset 1000 --hiddenSize 100\r\nlibthclnn_searchpath\t/Users/zhuxiaohu/torch/install/lib/lua/5.1/libTHCLNN.so\r\n-- Loading dataset\r\ndata/vocab.t7 not found\r\n-- Parsing Cornell movie dialogs data set ...\r\n [=============================================================== 387810/387810 =======>] Tot: 2s514ms | Step: 0ms\r\n-- Pre-processing data\r\n [============================================================= 1000/1000 =============>] Tot: 242ms | Step: 0ms\r\n-- Removing low frequency words\r\n [============================================================ 1569/1569 ==============>] Tot: 99ms | Step: 0ms\r\nWriting data/examples.t7 ...\r\n [============================================================= 1569/1569 =============>] Tot: 323ms | Step: 0ms\r\nWriting data/vocab.t7 ...\r\n\r\nDataset stats:\r\n  Vocabulary size: 2536\r\n         Examples: 1569\r\nUsing Apple , OpenCL platform: Apple\r\nUsing OpenCL device: HD Graphics 4000\r\n```\r\n\r\nIt seems all right until the following error occurs: \r\n\r\n\r\n```\r\nkernel source:\r\n1: // Threads per thread block\r\n2: #define THCL_NONCONTIG_REDUCE_BLOCK_SIZE 32 * 16\r\n3:\r\n4: inline float modifyOp(float _in1) {\r\n5:   float _out;\r\n6:   float *in1 = &_in1;\r\n7:   float *out = &_out;\r\n8:   *out = *in1;\r\n9:   return _out;\r\n10: }\r\n11:\r\n12: inline float reduceOp(float _in1, float _in2) {\r\n13:   // I guess the compiler can sort this stuff out :-P\r\n14:   float _out;\r\n15:   float *in1 = &_in1;\r\n16:   float *in2 = &_in2;\r\n17:   float *out = &_out;\r\n18:   *out = *in1 + *in2;\r\n19:   return _out;\r\n20: }\r\n21:\r\n22: // kernel argument that defines tensor layout\r\n23: typedef struct TensorInfoCl {\r\n24:   // Extracts size/stride information for the kernel.\r\n25:   // Successive dimensions can be collapsed if the size/strides match\r\n26:   // up and thus there are no holes between the dimensions. This is used\r\n27:   // to reduce the complexity of the problem.\r\n28:   // The optional `reduceDim` indicates a reduction dimension for the\r\n29:   // given tensor, so that the output size for this dimension will be 1.\r\n30:\r\n31:   int sizes[25];\r\n32:   int strides[25];\r\n33:   int offset;\r\n34:   int dims;\r\n35: } TensorInfoCl;\r\n36: // Contiguous tensors of more than one dimension are collapsed down\r\n37: // to one tensor\r\n38:\r\n39:\r\n40: // Translate a linear index for the apply to a float* offset;\r\n41: // specialized on `Dims` to reduce nvcc compilation time\r\n42:\r\n43:\r\n44: inline int IndexToOffset_998_get(int linearId, global const TensorInfoCl *info) {\r\n45:     return linearId + info->offset;\r\n46: }\r\n47:\r\n48: inline int IndexToOffset_999_get(int linearId, global const TensorInfoCl *info) {\r\n49:   int offset = info->offset;\r\n50:\r\n51:   // Use dynamic dims\r\n52:   for (int i = info->dims - 1; i >= 0; --i) {\r\n53:     int curDimIndex = linearId % info->sizes[i];\r\n54:     int curDimOffset = curDimIndex * info->strides[i];\r\n55:     offset += curDimOffset;\r\n56:\r\n57:     linearId /= info->sizes[i];\r\n58:   }\r\n59:\r\n60:   return offset;\r\n61: }\r\n62:\r\n63: inline int getLinearBlockId() {\r\n64:   return get_group_id(2) * get_num_groups(1) * get_num_groups(0) +\r\n65:     get_group_id(1) * get_num_groups(0) +\r\n66:     get_group_id(0);\r\n67: }\r\n68:\r\n69: // Block-wide reduction in shared memory helper; only /*threadIdx.x*/ get_local_id(0) == 0 will\r\n70: // return the reduced value\r\n71:\r\n72: inline float reduceBlock( local float* smem,\r\n73:                    int numVals,\r\n74:                    float threadVal,\r\n75:                    float init) {\r\n76:   if (numVals == 0) {\r\n77:     return init;\r\n78:   }\r\n79:\r\n80:   if ((int)get_local_id(0) < numVals) {\r\n81:     smem[ get_local_id(0)] = threadVal;\r\n82:   }\r\n83:\r\n84:   // First warp will perform reductions across warps\r\n85:   barrier(CLK_LOCAL_MEM_FENCE);\r\n86:   if ((get_local_id(0) / 32) == 0) {\r\n87:     float r = (int)get_local_id(0) < numVals ? smem[get_local_id(0)] : init;\r\n88:\r\n89:     for (int i = 32 + get_local_id(0); i < numVals; i += 32) {\r\n90:       r = reduceOp(r, smem[i]);\r\n91:     }\r\n92:\r\n93:     smem[get_local_id(0)] = r;\r\n94:   }\r\n95:\r\n96:   // First thread will perform reductions across the block\r\n97:   barrier(CLK_LOCAL_MEM_FENCE);\r\n98:\r\n99:   float r = init;\r\n100:   if (get_local_id(0) == 0) {\r\n101:     r = smem[0];\r\n102:\r\n103:     int numLanesParticipating = min(numVals, 32);\r\n104:\r\n105:     if (numLanesParticipating == 32) {\r\n106:       // Unroll for 32 == 32 and numVals >= 32\r\n107:       // #pragma unroll\r\n108:       // unrolling by hand, so compiler-independent\r\n109:\r\n110:         r = reduceOp(r, smem[1]);\r\n111:\r\n112:         r = reduceOp(r, smem[2]);\r\n113:\r\n114:         r = reduceOp(r, smem[3]);\r\n115:\r\n116:         r = reduceOp(r, smem[4]);\r\n117:\r\n118:         r = reduceOp(r, smem[5]);\r\n119:\r\n120:         r = reduceOp(r, smem[6]);\r\n121:\r\n122:         r = reduceOp(r, smem[7]);\r\n123:\r\n124:         r = reduceOp(r, smem[8]);\r\n125:\r\n126:         r = reduceOp(r, smem[9]);\r\n127:\r\n128:         r = reduceOp(r, smem[10]);\r\n129:\r\n130:         r = reduceOp(r, smem[11]);\r\n131:\r\n132:         r = reduceOp(r, smem[12]);\r\n133:\r\n134:         r = reduceOp(r, smem[13]);\r\n135:\r\n136:         r = reduceOp(r, smem[14]);\r\n137:\r\n138:         r = reduceOp(r, smem[15]);\r\n139:\r\n140:         r = reduceOp(r, smem[16]);\r\n141:\r\n142:         r = reduceOp(r, smem[17]);\r\n143:\r\n144:         r = reduceOp(r, smem[18]);\r\n145:\r\n146:         r = reduceOp(r, smem[19]);\r\n147:\r\n148:         r = reduceOp(r, smem[20]);\r\n149:\r\n150:         r = reduceOp(r, smem[21]);\r\n151:\r\n152:         r = reduceOp(r, smem[22]);\r\n153:\r\n154:         r = reduceOp(r, smem[23]);\r\n155:\r\n156:         r = reduceOp(r, smem[24]);\r\n157:\r\n158:         r = reduceOp(r, smem[25]);\r\n159:\r\n160:         r = reduceOp(r, smem[26]);\r\n161:\r\n162:         r = reduceOp(r, smem[27]);\r\n163:\r\n164:         r = reduceOp(r, smem[28]);\r\n165:\r\n166:         r = reduceOp(r, smem[29]);\r\n167:\r\n168:         r = reduceOp(r, smem[30]);\r\n169:\r\n170:         r = reduceOp(r, smem[31]);\r\n171:\r\n172:     } else {\r\n173:       for (int i = 1; i < numLanesParticipating; ++i) {\r\n174:         r = reduceOp(r, smem[i]);\r\n175:       }\r\n176:     }\r\n177:   }\r\n178:\r\n179:   return r;\r\n180: }\r\n181:\r\n182:\r\n183:\r\n184:\r\n185: inline int getReduceNoncontigDimSliceIndex() {\r\n186:   // Each thread handles one slice\r\n187:   return getLinearBlockId() * THCL_NONCONTIG_REDUCE_BLOCK_SIZE + /*threadIdx.x*/ get_local_id(0);\r\n188: }\r\n189:\r\n190: // Kernel that handles an entire reduction of a slice of a tensor per each thread\r\n191: kernel void\r\n192: THClTensor_reduceNoncontigDim(global TensorInfoCl *out_info,\r\n193:                               global float *out_data,\r\n194:                               global TensorInfoCl *in_info,\r\n195:                               global float *in_data,\r\n196:                               int reductionStride,\r\n197:                               int reductionSize,\r\n198:                               int totalSlices,\r\n199:                               float init) {\r\n200:   const int sliceIndex = getReduceNoncontigDimSliceIndex();\r\n201:\r\n202:   if ((int)sliceIndex >= totalSlices) {\r\n203:     return;\r\n204:   }\r\n205:\r\n206:   // Each thread picks a point in `out` and `in` for which it is\r\n207:   // producing the reduction\r\n208:   const int outOffset =\r\n209:     IndexToOffset_998_get(sliceIndex, &out_info[0]);\r\n210:   const int inBaseOffset =\r\n211:     IndexToOffset_998_get(sliceIndex, &in_info[0]);\r\n212:\r\n213:   // For each point in reductionSize, reduce into `r`\r\n214:   int inOffset = inBaseOffset;\r\n215:   float r = init;\r\n216:\r\n217:   for (int i = 0; (int)i < reductionSize; ++i) {\r\n218:     r = reduceOp(r, modifyOp(in_data[inOffset]));\r\n219:     inOffset += reductionStride;\r\n220:   }\r\n221:\r\n222:   // Write out reduced value\r\n223:   out_data[outOffset] = r;\r\n224: }\r\n225:\r\n226: inline int getReduceContigDimSliceIndex() {\r\n227:   // Each block handles one slice\r\n228:   return getLinearBlockId();\r\n229: }\r\n230:\r\n231: // Kernel that handles an entire reduction of a slice of a tensor per\r\n232: // each block\r\n233: kernel void\r\n234: THClTensor_reduceContigDim(global TensorInfoCl *out_info,\r\n235:                            global float *out_data,\r\n236:                            global TensorInfoCl *in_info,\r\n237:                            global float *in_data,\r\n238:                            int reductionSize,\r\n239:                            int totalSlices,\r\n240:                            float init,\r\n241:                            local float *smem) {\r\n242:   const int sliceIndex = getReduceContigDimSliceIndex();\r\n243:\r\n244:   if ((int)sliceIndex >= totalSlices) {\r\n245:     return;\r\n246:   }\r\n247:\r\n248:   // Get the offset in `out` for the reduction\r\n249:   const int outOffset =\r\n250:     IndexToOffset_998_get(sliceIndex, &out_info[0]);\r\n251:\r\n252:   // Get the base offset in `in` for this block's reduction\r\n253:   const int inBaseOffset =\r\n254:     IndexToOffset_998_get(sliceIndex, &in_info[0]);\r\n255:\r\n256:   // Each thread in the block will reduce some subset of elements in\r\n257:   // the slice. The elements are guaranteed contiguous starting at\r\n258:   // `inBaseOffset`.\r\n259:   float r = init;\r\n260:   for (int i = /*threadIdx.x*/ get_local_id(0); (int)i < reductionSize; i += /*blockDim.x*/ get_local_size(0)) {\r\n261:     r = reduceOp(r, modifyOp(in_data[inBaseOffset + i]));\r\n262:   }\r\n263:\r\n264:   // Reduce within the block\r\n265: //  extern __shared__ float smem[];\r\n266:   r = reduceBlock(smem, /*blockDim.x*/ get_local_size(0), r, init);\r\n267:\r\n268:   if (/*threadIdx.x*/ get_local_id(0) == 0) {\r\n269:     // Write out reduced value\r\n270:     out_data[outOffset] = r;\r\n271:   }\r\n272: }\r\n273:\r\n274:\r\n\r\nInvalid work group size, code -54\r\n/Users/zhuxiaohu/torch/install/bin/luajit: ...huxiaohu/torch/install/share/lua/5.1/clnn/LogSoftMax.lua:23:\r\nkernel source:\r\n1: // Threads per thread block\r\n2: #define THCL_NONCONTIG_REDUCE_BLOCK_SIZE 32 * 16\r\n3:\r\n4: inline float modifyOp(float _in1) {\r\n5:   float _out;\r\n6:   float *in1 = &_in1;\r\n7:   float *out = &_out;\r\n8:   *out = *in1;\r\n9:   return _out;\r\n10: }\r\n11:\r\n12: inline float reduceOp(float _in1, float _in2) {\r\n13:   // I guess the compiler can sort this stuff out :-P\r\n14:   float _out;\r\n15:   float *in1 = &_in1;\r\n16:   float *in2 = &_in2;\r\n17:   float *out = &_out;\r\n18:   *out = *in1 + *in2;\r\n19:   return _out;\r\n20: }\r\n21:\r\n22: // kernel argument that defines tensor layout\r\n23: typedef struct TensorInfoCl {\r\n24:   // Extracts size/stride information for the kernel.\r\n25:   // Successive dimensions can be collapsed if the size/strides match\r\n26:   // up and thus there are no holes between the dimensions. This is used\r\n27:   // to reduce the complexity of the problem.\r\n28:   // The optional `reduceDim` indicates a reduction dimension for the\r\n29:   // given tensor, so that the output size for this dimension will be 1.\r\n30:\r\n31:   int sizes[25];\r\n32:   int strides[25];\r\n33:   int offset;\r\n34:   int dims;\r\n35: } TensorInfoCl;\r\n36: // Contiguous tensors of more than one dimension are collapsed down\r\n37: // to one tensor\r\n38:\r\n39:\r\n40: // Translate a linear index for the apply to a float* offset;\r\n41: // specialized on `Dims` to reduce nvcc compilation time\r\n42:\r\n43:\r\n44: inline int IndexToOffset_998_get(int linearId, global const TensorInfoCl *info) {\r\n45:     return linearId + info->offset;\r\n46: }\r\n47:\r\n48: inline int IndexToOffset_999_get(int linearId, global const TensorInfoCl *info) {\r\n49:   int offset = info->offset;\r\n50:\r\n51:   // Use dynamic dims\r\n52:   for (int i = info->dims - 1; i >= 0; --i) {\r\n53:     int curDimIndex = linearId -1180192744nfo->sizes[i];\r\n54:     int curDimOffset = curDimIndex * info->strides[i];\r\n55:     offset += curDimOffset;\r\n56:\r\n57:     linearId /= info->sizes[i];\r\n58:   }\r\n59:\r\n60:   return offset;\r\n61: }\r\n62:\r\n63: inline int getLinearBlockId() {\r\n64:   return get_group_id(2) * get_num_groups(1) * get_num_groups(0) +\r\n65:     ge\r\nstack traceback:\r\n\t[C]: in function 'sum'\r\n\t...huxiaohu/torch/install/share/lua/5.1/clnn/LogSoftMax.lua:23: in function 'updateOutput'\r\n\t...s/zhuxiaohu/torch/install/share/lua/5.1/rnn/Recursor.lua:24: in function 'updateOutput'\r\n\t.../zhuxiaohu/torch/install/share/lua/5.1/rnn/Sequencer.lua:47: in function 'updateOutput'\r\n\t.../zhuxiaohu/torch/install/share/lua/5.1/nn/Sequential.lua:44: in function 'forward'\r\n\t./seq2seq.lua:79: in function 'train'\r\n\ttrain.lua:82: in main chunk\r\n\t[C]: in function 'dofile'\r\n\t...aohu/torch/install/lib/luarocks/rocks/trepl/scm-1/bin/th:143: in main chunk\r\n\t[C]: at 0x0109d1cbc0\r\n```\r\n\r\nSeems like I should implement some methods in `clnn` LogSoftMax.lua. "},"comment":{"url":"https://api.github.com/repos/hughperkins/clnn/issues/comments/202335334","html_url":"https://github.com/hughperkins/clnn/issues/38#issuecomment-202335334","issue_url":"https://api.github.com/repos/hughperkins/clnn/issues/38","id":202335334,"user":{"login":"hughperkins","id":123560,"avatar_url":"https://avatars.githubusercontent.com/u/123560?v=3","gravatar_id":"","url":"https://api.github.com/users/hughperkins","html_url":"https://github.com/hughperkins","followers_url":"https://api.github.com/users/hughperkins/followers","following_url":"https://api.github.com/users/hughperkins/following{/other_user}","gists_url":"https://api.github.com/users/hughperkins/gists{/gist_id}","starred_url":"https://api.github.com/users/hughperkins/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/hughperkins/subscriptions","organizations_url":"https://api.github.com/users/hughperkins/orgs","repos_url":"https://api.github.com/users/hughperkins/repos","events_url":"https://api.github.com/users/hughperkins/events{/privacy}","received_events_url":"https://api.github.com/users/hughperkins/received_events","type":"User","site_admin":false},"created_at":"2016-03-28T10:32:31Z","updated_at":"2016-03-28T10:32:31Z","body":"In theory, this was fixed yesterday, in https://github.com/hughperkins/cltorch/commit/2df6d9928d16613ef1cbc51e1324ae929db95979 . Do you mind installing the latest cltorch, `luarocks install cltorch`, and confirm the problem does/doesnt persist?"}}