void TestSseAstro::testDegsToDms() { double decimalDeg(-49.227750000); // -49 13 39.90 char sign; int degs, mins; double secs; SseAstro::decimalDegreesToDms(decimalDeg, &sign, °s, &mins, &secs); cout << "decimalDeg: " << decimalDeg << " d:m:s: " << sign << degs << ":" << mins << ":" << secs << endl; double valueTol(1e-6); cu_assert(sign == '-'); assertLongsEqual(49, degs); assertLongsEqual(13, mins); assertDoublesEqual(39.9, secs, valueTol); // test negative zero degrees decimalDeg = -0.5; // -00 30 00 SseAstro::decimalDegreesToDms(decimalDeg, &sign, °s, &mins, &secs); cout << "decimalDeg: " << decimalDeg << " d:m:s: " << sign << degs << ":" << mins << ":" << secs << endl; cu_assert(sign == '-'); assertLongsEqual(00, degs); assertLongsEqual(30, mins); assertDoublesEqual(00, secs, valueTol); decimalDeg = 45.5; // +45 30 00 SseAstro::decimalDegreesToDms(decimalDeg, &sign, °s, &mins, &secs); cout << "decimalDeg: " << decimalDeg << " d:m:s: " << sign << degs << ":" << mins << ":" << secs << endl; cu_assert(sign == '+'); assertLongsEqual(45, degs); assertLongsEqual(30, mins); assertDoublesEqual(00, secs, valueTol); // handle roundoff of seconds to next minute decimalDeg = 41.9; // +41 54 00 (not: +41 53 60) SseAstro::decimalDegreesToDms(decimalDeg, &sign, °s, &mins, &secs); cout << "decimalDeg: " << decimalDeg << " d:m:s: " << sign << degs << ":" << mins << ":" << secs << endl; cu_assert(sign == '+'); assertLongsEqual(41, degs); assertLongsEqual(54, mins); assertDoublesEqual(00.0, secs, valueTol); }
int main(){ init_test(); const std::string source = ".version 4.2\n" ".target sm_20\n" ".address_size 64\n" ".visible .entry kernel(.param .u64 kernel_param_0) {\n" ".reg .s32 %r<2>;\n" ".reg .s64 %rd<3>;\n" "bra BB1_2;\n" "ld.param.u64 %rd1, [kernel_param_0];\n" "cvta.to.global.u64 %rd2, %rd1;\n" "mov.u32 %r1, 5;\n" "st.global.u32 [%rd2], %r1;\n" "BB1_2: ret;\n" "}\n"; CUmodule modId = 0; CUfunction funcHandle = 0; cu_assert(cuModuleLoadData(&modId, source.c_str())); cu_assert(cuModuleGetFunction(&funcHandle, modId, "kernel")); CUdeviceptr devValue; int hostValue = 10; cu_assert(cuMemAlloc(&devValue, sizeof(int))); cu_assert(cuMemcpyHtoD(devValue, &hostValue, sizeof(hostValue))); void * params[] = {&devValue}; cu_assert(cuLaunchKernel(funcHandle, 1,1,1, 1,1,1, 0,0, params, nullptr)); cu_assert(cuMemcpyDtoH(&hostValue, devValue, sizeof(hostValue))); assert(hostValue == 10); std::cout << hostValue << "\n"; cu_assert(cuMemFree(devValue)); cu_assert(cuModuleUnload(modId)); return 0; }
void TestSite::setUpDxs() { // Put some dxs on the list. // Their names should be the same as the // ones in the expected components configuration info. for (unsigned int i=0; i<DxCount; i++) { DxProxy *proxy; // create proxies proxy = new DxProxy(site_->dxManager()); // give the dx a name stringstream name; name << "dxtest" << i; proxy->setName(name.str()); cu_assert(proxy->getName() == name.str()); // register with the site proxy->notifyInputConnected(); dxList_.push_back(proxy); } }
void TestSite::setUpTscopes() { // Put some tscopes on the list. // These names should match the ones in the // expected components info. for (unsigned int i=0; i<TscopeCount; i++) { TscopeProxy *proxy; // create proxies proxy = new TscopeProxy(site_->tscopeManager()); // give the tscope a name stringstream name; name << "tscopetest" << i; proxy->setName(name.str()); cu_assert(proxy->getName() == name.str()); // register with the site proxy->notifyInputConnected(); tscopeList_.push_back(proxy); } }
void TestSite::setUpIfcs() { // Put some ifcs on the list. // These names should match the ones in the // expected components info, plus one more. for (unsigned int i=0; i<IfcCount; i++) { IfcProxy *proxy; // create proxies proxy = new IfcProxy(site_->ifcManager()); // give the ifc a name stringstream name; name << "ifctest" << i; proxy->setName(name.str()); cu_assert(proxy->getName() == name.str()); // register with the site proxy->notifyInputConnected(); ifcList_.push_back(proxy); } }
void TestSseAstro::testSiderealTimeConvert() { double siderealAngleRads(SseAstro::hoursToRadians(1.0)); int expectedSolarTimeSecs(static_cast<int>(SseAstro::SecsPerHour * SseAstro::SolarDaysPerSideralDay)); int solarTimeSecs = SseAstro::siderealRadsToSolarTimeSecs(siderealAngleRads); //cout << "solarTimeSecs=" << solarTimeSecs << endl; cu_assert(expectedSolarTimeSecs == solarTimeSecs); }
int main(){ init_test(); const std::string test_source = ".version 4.2\n" ".target sm_20\n" ".address_size 64\n" ".visible .entry _Z6kernelPfi(\n" ".param .u64 _Z6kernelPfi_param_0,\n" ".param .u32 _Z6kernelPfi_param_1){\n" ".reg .pred %p<2>;\n" ".reg .f32 %f<3>;\n" ".reg .s32 %r<3>;\n" ".reg .s64 %rd<5>;\n" "ld.param.u64 %rd1, [_Z6kernelPfi_param_0];\n" "ld.param.u32 %r2, [_Z6kernelPfi_param_1];\n" "mov.u32 %r1, %tid.x;\n" "setp.ge.u32 %p1, %r1, %r2;\n" "@%p1 bra BB0_2;\n" "cvta.to.global.u64 %rd2, %rd1;\n" "cvt.rn.f32.u32 %f1, %r1;\n" "mul.f32 %f2, %f1, 0f3FC00000;\n" "mul.wide.u32 %rd3, %r1, 4;\n" "add.s64 %rd4, %rd2, %rd3;\n" "st.global.f32 [%rd4], %f2;\n" "BB0_2:\n" "ret;\n" "}"; CUmodule modId = 0; CUfunction funcHandle = 0; cu_assert(cuModuleLoadData(&modId, test_source.c_str())); cu_assert(cuModuleGetFunction(&funcHandle, modId, "_Z6kernelPfi")); CUdeviceptr devArray; int size = 10; float hostArray[size]; memset(hostArray, 0, size * sizeof(hostArray[0])); cu_assert(cuMemAlloc(&devArray, sizeof(float) * size)); void * params[] = {&devArray, &size}; auto result = cuLaunchKernel(funcHandle, 1,1,1, size*2,1,1, 0,0, params, nullptr); cu_assert(result); cu_assert(cuMemcpyDtoH(&hostArray, devArray, sizeof(hostArray[0])*size)); cu_assert(cuMemFree(devArray)); cu_assert(cuModuleUnload(modId)); for (int i=0 ; i<size ; ++i) std::cout << hostArray[i] << '\n'; return 0; }
int main(){ init_test(); const std::string source = ".version 4.2\n" ".target sm_20\n" ".address_size 64\n" ".visible .entry kernel_4(\n" ".param .u32 kernel_4_param_0,\n" ".param .u64 kernel_4_param_1\n" ")\n" "{\n" ".reg .s32 %r<3>;\n" ".reg .s64 %rd<3>;\n" "ld.param.u32 %r1, [kernel_4_param_0];\n" "ld.param.u64 %rd1, [kernel_4_param_1];\n" "cvta.to.global.u64 %rd2, %rd1;\n" "add.s32 %r2, %r1, 7;\n" "st.global.u32 [%rd2], %r2;\n" "ret;\n" "}"; CUmodule modId = 0; CUfunction funcHandle = 0; cu_assert(cuModuleLoadData(&modId, source.c_str())); cu_assert(cuModuleGetFunction(&funcHandle, modId, "kernel_4")); CUdeviceptr devValue; int hostValue = 10; cu_assert(cuMemAlloc(&devValue, sizeof(int))); void * params[] = {&hostValue, &devValue}; cu_assert(cuLaunchKernel(funcHandle, 1,1,1, 1,1,1, 0,0, params, nullptr)); int result = 0; cu_assert(cuMemcpyDtoH(&result, devValue, sizeof(result))); assert(result == hostValue + 7); std::cout << result << "\n"; cu_assert(cuMemFree(devValue)); cu_assert(cuModuleUnload(modId)); return 0; }
void TestDx::testDxCmdLineArgs() { string progName("TestDxCmdLineArgs"); int defaultMainPort(9999); int defaultRemotePort(5555); string defaultHost("thisHost"); string defaultSciDataDir("fooDir"); string defaultSciDataPrefix("nss.p10"); bool defaultBroadcast(false); bool defaultNoUi(false); bool defaultVaryOutputData(false); string defaultDxName("dxsim33"); bool defaultRemoteMode(false); DxCmdLineArgs cmdArgs(progName, defaultMainPort, defaultRemotePort, defaultHost, defaultSciDataDir, defaultSciDataPrefix, defaultBroadcast, defaultNoUi, defaultVaryOutputData, defaultDxName, defaultRemoteMode); // test usage message output cmdArgs.usage(); cerr << endl; // check the defaults cu_assert (cmdArgs.getMainPort() == 9999); cu_assert (cmdArgs.getRemotePort() == 5555); cu_assert (cmdArgs.getHostName() == "thisHost"); cu_assert (cmdArgs.getSciDataDir() == "fooDir"); cu_assert (cmdArgs.getSciDataPrefix() == "nss.p10"); cu_assert (! cmdArgs.getBroadcast()); cu_assert (! cmdArgs.getNoUi()); cu_assert (! cmdArgs.getVaryOutputData()); cu_assert (cmdArgs.getDxName() == "dxsim33"); cu_assert (! cmdArgs.getRemoteMode()); // try setting good parameters const char *argv[] = { "ProgName", "-host", "matrix", "-mainport", "8888", "-remport", "9999", "-sddir", "../scienceData", "-sdprefix", "nss.p6", "-broadcast", "-name", "dxsim127", "-remote", "-noui", "-vary" }; const int argc = ARRAY_LENGTH(argv); // verify that everything parses cu_assert(cmdArgs.parseArgs(argc, const_cast<char **>(argv))); // check the values cu_assert (cmdArgs.getHostName() == "matrix"); cu_assert (cmdArgs.getMainPort() == 8888); cu_assert (cmdArgs.getRemotePort() == 9999); cu_assert (cmdArgs.getSciDataDir() == "../scienceData"); cu_assert (cmdArgs.getSciDataPrefix() == "nss.p6"); cu_assert (cmdArgs.getBroadcast()); cu_assert (cmdArgs.getDxName() == "dxsim127"); cu_assert (cmdArgs.getRemoteMode()); cu_assert (cmdArgs.getNoUi()); cu_assert (cmdArgs.getVaryOutputData()); cout << "Test a bad port number:" << endl; const char *argvBadPort[] = { "ProgName", "-host", "matrix", "-mainport", "badportnumber", }; const int argcBadPort = ARRAY_LENGTH(argvBadPort); cu_assert(cmdArgs.parseArgs(argcBadPort, const_cast<char **>(argvBadPort)) == false); }
void TestDxArchiver::testDxArchiverCmdLineArgs() { string progName("TestDxArchiverCmdLineArgs"); int defaultSsePort(2222); int defaultDxPort(3333); string defaultSseHostname("fred-main"); string defaultName("archiver7"); bool defaultNoUi(false); DxArchiverCmdLineArgs cmdArgs(progName, defaultSseHostname, defaultSsePort, defaultDxPort, defaultNoUi, defaultName); // test usage message output cmdArgs.usage(); cerr << endl; // check the defaults cu_assert (cmdArgs.getSsePort() == defaultSsePort); cu_assert (cmdArgs.getDxPort() == defaultDxPort); cu_assert (cmdArgs.getSseHostname() == defaultSseHostname); cu_assert (cmdArgs.getNoUi() == defaultNoUi); cu_assert (cmdArgs.getName() == defaultName); // try setting alternate parameters const char *argv[] = { "ProgName", "-host", "barney", "-sse-port", "8888", "-dx-port", "8877", "-name", "archiver8" }; const int argc = ARRAY_LENGTH(argv); // verify that everything parses cu_assert(cmdArgs.parseArgs(argc, const_cast<char **>(argv))); // check the values cu_assert (cmdArgs.getSseHostname() == "barney"); cu_assert (cmdArgs.getSsePort() == 8888); cu_assert (cmdArgs.getDxPort() == 8877); cu_assert (cmdArgs.getName() == "archiver8"); cout << "test a bad bort number" << endl; const char *argvBadPort[] = { "ProgName", "-host", "matrix", "-sse-port", "badportnumber", }; const int argcBadPort = ARRAY_LENGTH(argvBadPort); cu_assert(cmdArgs.parseArgs(argcBadPort, const_cast<char **>(argvBadPort)) == false); }
void TestSite::testGetDxsForIfc() { cout << "testGetDxsForIfc" << endl; // Test that the NssComponentTree correctly // uses the expectedNssComponentsTree information // to return the dx proxies that // are attached to each of the IFs. // expected ifc/dx configuration: // ifctest0: dxtest0 dxtest1 // ifctest1: dxtest2 // Note: // a proxy for ifctest2 exists, but it is not listed // in the configuration info, and so should have no // dxs attached. NssComponentTree * tree = site_->getAllComponents(); IfcList &ifcList = tree->getIfcs(); cu_assert(ifcList.size() == IfcCount); // Get the dx proxies for each of the ifcs // and make sure they are the right ones. unsigned int nIfcsPassed = 0; for(IfcList::iterator it = ifcList.begin(); it != ifcList.end(); ++it) { IfcProxy *ifcProxy = *it; if (ifcProxy->getName() == "ifctest0") { DxList & dxList = tree->getDxsForIfc(ifcProxy); cu_assert(dxList.size() == 2); // make sure the names match DxList::iterator it; for (it = dxList.begin(); it != dxList.end(); ++it) { DxProxy *proxy = *it; cu_assert(proxy->getName() == "dxtest0" || proxy->getName() == "dxtest1"); } nIfcsPassed++; } else if (ifcProxy->getName() == "ifctest1") { DxList & dxList = tree->getDxsForIfc(ifcProxy); cu_assert(dxList.size() == 1); nIfcsPassed++; } else if (ifcProxy->getName() == "ifctest2") { // this one should not have any dxs attached to it DxList & dxList = tree->getDxsForIfc(ifcProxy); cu_assert(dxList.size() == 0); nIfcsPassed++; } else { // unexpected ifc is on the list cout << "error: unexpected ifc on the list: " << ifcProxy->getName() << endl; cu_assert(0); } } cu_assert(nIfcsPassed == IfcCount); delete tree; }
void TestSite::testGetNssComponentTree() { cout << "testGetNssComponentTree" << endl; NssComponentTree * tree = site_->getAllComponents(); DxList &dxList = tree->getDxs(); cu_assert(dxList.size() == DxCount); // make sure all the dxs in the component tree are // on the original list used to create it for(DxList::iterator it = dxList.begin(); it != dxList.end(); ++it) { DxProxy *proxy = *it; cout << proxy->getName() << endl; cu_assert(find(dxList_.begin(), dxList_.end(), proxy) != dxList_.end()); } IfcList &ifcList = tree->getIfcs(); cu_assert(ifcList.size() == IfcCount); // make sure all the ifcs in the component tree are // on the original list used to create it for(IfcList::iterator it = ifcList.begin(); it != ifcList.end(); ++it) { IfcProxy *proxy = *it; cout << proxy->getName() << endl; cu_assert(find(ifcList_.begin(), ifcList_.end(), proxy) != ifcList_.end()); } vector<string> beamNames; beamNames.push_back("beam0"); beamNames.push_back("beam1"); IfcList ifcsForBeams(tree->getIfcsForBeams(beamNames)); cu_assert(ifcsForBeams.size() == 2); for (IfcList::iterator it = ifcsForBeams.begin(); it != ifcsForBeams.end(); ++it) { IfcProxy *proxy = *it; cu_assert(proxy->getName() == "ifctest0" || proxy->getName() == "ifctest1"); } DxList dxsForBeams(tree->getDxsForBeams(beamNames)); cu_assert(dxsForBeams.size() == 3); for (DxList::iterator it = dxsForBeams.begin(); it != dxsForBeams.end(); ++it) { DxProxy *proxy = *it; cu_assert(proxy->getName() == "dxtest0" || proxy->getName() == "dxtest1" || proxy->getName() == "dxtest2"); } TscopeList &tscopeList = tree->getTscopes(); cu_assert(tscopeList.size() == TscopeCount); TestSigList &testSigList = tree->getTestSigs(); cu_assert(testSigList.size() == 0); delete tree; }
T run_function(const std::string& name, const T input, const int shiftValue) { const std::string test_source = "//\n" "// Generated by NVIDIA NVVM Compiler\n" "//\n" "// Compiler Build ID: CL-19856038\n" "// Cuda compilation tools, release 7.5, V7.5.17\n" "// Based on LLVM 3.4svn\n" "//\n" "\n" ".version 4.3\n" ".target sm_20\n" ".address_size 64\n" "\n" " // .globl _Z10kernel_s32Piii\n" "\n" ".visible .entry _Z10kernel_s32Piii(\n" " .param .u64 _Z10kernel_s32Piii_param_0,\n" " .param .u32 _Z10kernel_s32Piii_param_1,\n" " .param .u32 _Z10kernel_s32Piii_param_2\n" ")\n" "{\n" " .reg .b32 %r<4>;\n" " .reg .b64 %rd<3>;\n" "\n" "\n" " ld.param.u64 %rd1, [_Z10kernel_s32Piii_param_0];\n" " ld.param.u32 %r1, [_Z10kernel_s32Piii_param_1];\n" " ld.param.u32 %r2, [_Z10kernel_s32Piii_param_2];\n" " cvta.to.global.u64 %rd2, %rd1;\n" " shr.s32 %r3, %r1, %r2;\n" " st.global.u32 [%rd2], %r3;\n" " ret;\n" "}\n" "\n" " // .globl _Z10kernel_s64Pxxi\n" ".visible .entry _Z10kernel_s64Pxxi(\n" " .param .u64 _Z10kernel_s64Pxxi_param_0,\n" " .param .u64 _Z10kernel_s64Pxxi_param_1,\n" " .param .u32 _Z10kernel_s64Pxxi_param_2\n" ")\n" "{\n" " .reg .b32 %r<2>;\n" " .reg .b64 %rd<5>;\n" "\n" "\n" " ld.param.u64 %rd1, [_Z10kernel_s64Pxxi_param_0];\n" " ld.param.u64 %rd2, [_Z10kernel_s64Pxxi_param_1];\n" " ld.param.u32 %r1, [_Z10kernel_s64Pxxi_param_2];\n" " cvta.to.global.u64 %rd3, %rd1;\n" " shr.s64 %rd4, %rd2, %r1;\n" " st.global.u64 [%rd3], %rd4;\n" " ret;\n" "}\n" "\n" " // .globl _Z10kernel_u32Pjji\n" ".visible .entry _Z10kernel_u32Pjji(\n" " .param .u64 _Z10kernel_u32Pjji_param_0,\n" " .param .u32 _Z10kernel_u32Pjji_param_1,\n" " .param .u32 _Z10kernel_u32Pjji_param_2\n" ")\n" "{\n" " .reg .b32 %r<4>;\n" " .reg .b64 %rd<3>;\n" "\n" "\n" " ld.param.u64 %rd1, [_Z10kernel_u32Pjji_param_0];\n" " ld.param.u32 %r1, [_Z10kernel_u32Pjji_param_1];\n" " ld.param.u32 %r2, [_Z10kernel_u32Pjji_param_2];\n" " cvta.to.global.u64 %rd2, %rd1;\n" " shr.u32 %r3, %r1, %r2;\n" " st.global.u32 [%rd2], %r3;\n" " ret;\n" "}\n" "\n" " // .globl _Z10kernel_u64Pyyi\n" ".visible .entry _Z10kernel_u64Pyyi(\n" " .param .u64 _Z10kernel_u64Pyyi_param_0,\n" " .param .u64 _Z10kernel_u64Pyyi_param_1,\n" " .param .u32 _Z10kernel_u64Pyyi_param_2\n" ")\n" "{\n" " .reg .b32 %r<2>;\n" " .reg .b64 %rd<5>;\n" "\n" "\n" " ld.param.u64 %rd1, [_Z10kernel_u64Pyyi_param_0];\n" " ld.param.u64 %rd2, [_Z10kernel_u64Pyyi_param_1];\n" " ld.param.u32 %r1, [_Z10kernel_u64Pyyi_param_2];\n" " cvta.to.global.u64 %rd3, %rd1;\n" " shr.u64 %rd4, %rd2, %r1;\n" " st.global.u64 [%rd3], %rd4;\n" " ret;\n" "}\n" "\n" "\n" ; CUmodule modId = 0; CUfunction funcHandle = 0; cu_assert(cuModuleLoadData(&modId, test_source.c_str())); cu_assert(cuModuleGetFunction(&funcHandle, modId, name.c_str())); T output; CUdeviceptr devOutput; cu_assert(cuMemAlloc(&devOutput, sizeof(output))); void * params[] = {&devOutput, (void*)&input, (void*)&shiftValue}; auto result = cuLaunchKernel(funcHandle, 1,1,1, 1,1,1, 0,0, params, nullptr); cu_assert(result); cu_assert(cuMemcpyDtoH(&output, devOutput, sizeof(output))); cu_assert(cuMemFree(devOutput)); cu_assert(cuModuleUnload(modId)); return output; }