{"action":"created","issue":{"url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/70","repository_url":"https://api.github.com/repos/Genoil/cpp-ethereum","labels_url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/70/labels{/name}","comments_url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/70/comments","events_url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/70/events","html_url":"https://github.com/Genoil/cpp-ethereum/issues/70","id":163568389,"number":70,"title":"cleaner copy for dag generation","user":{"login":"nerdralph","id":5561594,"avatar_url":"https://avatars.githubusercontent.com/u/5561594?v=3","gravatar_id":"","url":"https://api.github.com/users/nerdralph","html_url":"https://github.com/nerdralph","followers_url":"https://api.github.com/users/nerdralph/followers","following_url":"https://api.github.com/users/nerdralph/following{/other_user}","gists_url":"https://api.github.com/users/nerdralph/gists{/gist_id}","starred_url":"https://api.github.com/users/nerdralph/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/nerdralph/subscriptions","organizations_url":"https://api.github.com/users/nerdralph/orgs","repos_url":"https://api.github.com/users/nerdralph/repos","events_url":"https://api.github.com/users/nerdralph/events{/privacy}","received_events_url":"https://api.github.com/users/nerdralph/received_events","type":"User","site_admin":false},"labels":[],"state":"open","locked":false,"assignee":null,"assignees":[],"milestone":null,"comments":10,"created_at":"2016-07-03T16:06:10Z","updated_at":"2016-07-03T21:04:04Z","closed_at":null,"body":"<pre>\r\ndiff --git a/libethash-cl/ethash_cl_miner_kernel.cl b/libethash-cl/ethash_cl_min\r\nindex 390a021..54027bd 100644\r\n--- a/libethash-cl/ethash_cl_miner_kernel.cl\r\n+++ b/libethash-cl/ethash_cl_miner_kernel.cl\r\n@@ -227,12 +227,14 @@ typedef struct\r\n } hash32_t;\r\n\r\n typedef union {\r\n+       uint16   vec;\r\n        uint     words[64 / sizeof(uint)];\r\n        uint2    uint2s[64 / sizeof(uint2)];\r\n        uint4    uint4s[64 / sizeof(uint4)];\r\n } hash64_t;\r\n\r\n typedef union {\r\n+       uint16   vec;\r\n        uint     words[200 / sizeof(uint)];\r\n        uint2    uint2s[200 / sizeof(uint2)];\r\n        uint4    uint4s[200 / sizeof(uint4)];\r\n@@ -348,7 +350,7 @@ __kernel void ethash_search(\r\n\r\n static void SHA3_512(uint2* s, uint isolate)\r\n {\r\n-       for (uint i = 8; i != 25; ++i)\r\n+       for (uint i = 9; i != 25; ++i)\r\n        {\r\n                s[i] = (uint2){ 0, 0 };\r\n        }\r\n@@ -363,7 +365,9 @@ __kernel void ethash_calculate_dag_item(uint start, __global\r\n        if (node_index > DAG_SIZE * 2) return;\r\n\r\n        hash200_t dag_node;\r\n-       copy(dag_node.uint4s, g_light[node_index % LIGHT_SIZE].uint4s, 4);\r\n+       //copy(dag_node.uint4s, g_light[node_index % LIGHT_SIZE].uint4s, 4);\r\n+       dag_node.vec = g_light[node_index % LIGHT_SIZE].vec;\r\n+\r\n        dag_node.words[0] ^= node_index;\r\n        SHA3_512(dag_node.uint2s, isolate);\r\n\r\n@@ -375,5 +379,6 @@ __kernel void ethash_calculate_dag_item(uint start, __global\r\n                }\r\n        }\r\n        SHA3_512(dag_node.uint2s, isolate);\r\n-       copy(g_dag[node_index].uint4s, dag_node.uint4s, 4);\r\n+       //copy(g_dag[node_index].uint4s, dag_node.uint4s, 4);\r\n+       g_dag[node_index].vec = dag_node.vec;\r\n }\r\n</pre>\r\n\r\nIt doesn't seem to be any faster, but I think assigning one vector to another is cleaner, and in theory the OpenCL complier should auto-vectorize it (i.e. a simd-16 assignment of int32s).\r\n"},"comment":{"url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/comments/230175252","html_url":"https://github.com/Genoil/cpp-ethereum/issues/70#issuecomment-230175252","issue_url":"https://api.github.com/repos/Genoil/cpp-ethereum/issues/70","id":230175252,"user":{"login":"nerdralph","id":5561594,"avatar_url":"https://avatars.githubusercontent.com/u/5561594?v=3","gravatar_id":"","url":"https://api.github.com/users/nerdralph","html_url":"https://github.com/nerdralph","followers_url":"https://api.github.com/users/nerdralph/followers","following_url":"https://api.github.com/users/nerdralph/following{/other_user}","gists_url":"https://api.github.com/users/nerdralph/gists{/gist_id}","starred_url":"https://api.github.com/users/nerdralph/starred{/owner}{/repo}","subscriptions_url":"https://api.github.com/users/nerdralph/subscriptions","organizations_url":"https://api.github.com/users/nerdralph/orgs","repos_url":"https://api.github.com/users/nerdralph/repos","events_url":"https://api.github.com/users/nerdralph/events{/privacy}","received_events_url":"https://api.github.com/users/nerdralph/received_events","type":"User","site_admin":false},"created_at":"2016-07-03T21:04:03Z","updated_at":"2016-07-03T21:04:03Z","body":"@wolf9466 Thanks. Now I remember reading the 64 work-items/wavefront before,  so 256 VGPRs * 4 bytes each * 64 work-items = 64KB.\r\nMy brain is still learning how to visualize multi-threaded SIMD execution flow..."}}
{"actor":{"display_login":"nerdralph"}}