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;
}
